
    ukiF                   1   d Z ddlmZ ddlmZmZmZmZ ddlZddl	Z	ddl
Z
ddlZddlZddlmZmZmZmZ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l0m1Z1 ddl2m3Z3 ddl2m4Z5 ddl6m7Z7 ddl6mZ8 ddl9m:Z: dd l;m<Z< dd!l=m>Z> dd"l?m@Z@ dd#l?mAZA dd$l?mBZB dd%l?mCZC dd&l?mDZD dd'l?mEZE dd(l?mFZF ddlGmZH dd)lGmIZJ dd*lGmKZK dd+lGmLZM ddlNmZO dd,lNmPZP dd*lNmKZQ dd-lNmRZS dd.lTmUZU dd*lTmKZV dd/lWmXZXmYZY dd0lZm[Z[m\Z\ dd1l]m^Z^ dd2l]m_Z_ dd3l]m`Z` dd4l]maZa dd5lbmcZc ddldmeZf ddleZgeUj                  ZheOj                  ZjeHj                  ejz  Zkejj                  Zlejj                  ZmeHj                  j                  Zn egj                  d6      Zpd7Zq egj                  egj                        j                  Zud8Zve
j                  Zwe_excZxZye`ezcZzZ{eHj                  hZ}dd9Z~d: Z G d; d<      Zeeej                  d=f   geed=f   f   Ze	j
                   G d> d?             Z G d@ dAe      Ze	j
                   G dB dC             Z	 	 	 	 ddDZ	 	 ddEZ	 d	 	 	 ddGZ	 	 	 	 d	 	 	 	 	 ddIZdddJZeOj                  D  ci c]  } | i  c} Z e       Z edK      ZeOj                  j&                  fdHdL	 	 	 	 	 	 	 ddMZdN Z	 d	 	 	 	 	 	 	 ddOZ	 	 	 	 ddPZ e	j
                  dFQ       G dR dS             Z	 	 	 	 	 	 	 	 ddTZddFdU	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d dVZddW	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddXZddFdY	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddZZdd[Z G d\ d]e      Z	 	 	 	 	 	 dd^Z	 	 	 	 	 	 	 	 dd_Zdd`Z eeVjF                  dFa      	 	 ddb       Z eeVjJ                  dFa      	 	 ddc       Zdd Zde Z	 	 	 	 	 	 ddfZ	 	 	 	 	 	 	 	 d	dgZd
dhZ	 	 	 	 	 	 	 	 	 	 ddiZ	 	 	 	 	 	 	 	 	 	 ddjZ	 	 	 	 	 	 	 	 	 	 ddkZdl Z e	j
                  dHm       G dn do             Zdp Z eeKjd                  dFa      ddq       ZddrZ	 	 	 	 	 	 ddsZ	 	 	 	 	 	 ddtZ eeKjJ                  dFa      	 	 ddu       Z eeKjp                  g eOj                  v      ddw       Zdx Zefjv                  eFjx                  jz                  efj|                  eFjx                  j~                  efj                  eFjx                  j                  iZefjv                   edy      efj|                   egj                  egj                        j                  iZ eefj                  eeū      Z  eej                        eƫ       efjv                  eFjx                  j                  efj|                  eFjx                  j                  efj                  eFjx                  j                  iZefjv                   edz      efj|                   egj                  egj                        j                  iZ eefj                  ee̫      Z  eej                        eͫ       efjv                  eFjx                  j                  efj|                  eFjx                  j                  efj                  eFjx                  j                  iZefjv                  d{efj|                  diZ eefj                  eeѫ      Z  eej                        eӫ        eej                        dd|       Z eej                        dd}       Z eeVj                        	 	 	 	 dd~       Z eej                  g eOj                  v      	 	 dd       Zd Z eej                        	 	 dd       ZߐddZ eej                  g eOj                  v      	 	 dd       Z eej                  g eOj                  v      dd       Z eej                  g eOj                  v      dd       Z eej                  g eOj                  v      dd       Z eej                  g eOj                  v      	 	 dd       Z eej                        dd       Z eej                        	 	 dd       Z eej                        dd       Z	 	 	 	 	 	 	 	 	 	 	 	 ddZ eej                  g eOj                  dFL       eej                  dFa      dd              Z G d de      Zd Z eej                        dd       Z eej                  dFg eOj                        dd       Z eej                  dFg eOj                        dd       Z	 	 ddZ eej                  dFa      dd       Z eej                   dFa      dd       Z eej                  g eOj                  dFL      dd       Z eej                  g eOj                  dFL      dd       Z eej                  g eOj                  dFL      dd       Z eej                  g eOj                  dFL      dd       Z	 eej                  g eOj                  v      dd       Z eej                  g eOj                  dFL      dd       Z eej                  g eOj                  v      dd       Z eej                         dd       Z eej$                        dd       Z eej(                        dd       Z eej,                        dd       Z eej0                  g eOj                  v      dd       Z eej4                  dFa      dd       Z eej8                        dd       Z eej<                  dFa      dd       Z eej@                        dd       Z! eejD                        dd       Z# eejH                        dd       Z% eejL                        dd       Z' eejP                        dd       Z) eejT                        dd       Z+ eejX                        dd       Z- eej\                        dd       Z/ eej`                        dd       Z1 eejd                        dd       Z3 eejh                        dd       Z5 eejl                        dd       Z7ejp                  e@jr                  jt                  ejv                  e@jr                  jx                  ejz                  e@jr                  j|                  ej~                  e@jr                  j                  ej                  e@jr                  j                  ej                  e@jr                  j                  iZEejp                  e@jr                  jt                  ejv                  e@jr                  jx                  ejz                  e@jr                  j                  ej~                  e@jr                  j                  ej                  e@jr                  j                  ej                  e@jr                  j                  iZJejp                  e@j                  j                  ejv                  e@j                  j                  ejz                  e@j                  j                  ej~                  e@j                  j                  ej                  e@j                  j                  ej                  e@j                  j                  iZRddZSddZTejp                  ejv                  ejz                  ej~                  ej                  ej                  fD ]4  ZU  eeUg eOj                  v       e
j                  eTeU             6  eej                  g eOj                  dFL      dd       ZW eej                        dd       ZY eej                  g eOj                  dFL      dd       Z[ eej                  g eOj                  v      dd       Z] eej                  g eOj                  v      dd       Z_d Z` eej                        dd       Zb	 	 	 	 	 	 	 	 	 	 ddZc eej                  g eOj                  dFL      	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dd       Ze	 	 ddZf eej                  g eOj                  v      	 	 dd       Zh eej                  g eOj                  v      dd       Zj ee&j                  g eOj                  v      dd       Zl ee&j                        dd       Zn eej                        	 	 	 	 	 	 	 	 	 	 dd       Zp eej                        	 	 ddÄ       Zr eej                        ddĄ       Zt eeKj                  g eOj                  v      ddń       Zv eeKj                  g eOj                  v      ddƄ       Zx eeQj                        ddǄ       Zz eeQj                        	 	 ddȄ       Z| eej                  g eOj                  v      	 	 ddɄ       Z~ eej                  g eOj                  dFL      ddʄ       Z eej                  g eOj                  dFL      dd˄       Z eej                  g eOj                  dFL      dd̄       Z eej
                  g eOj                  dFL      dd̈́       Z eej                        dd΄       Z eeKj                        ddτ       Z eeQj                        	 	 ddЄ       Zdф Z eeQj                        	 	 dd҄       Z eeQj                         	 	 ddӄ       Z eeQj$                        ddԄ       Z eej(                  g eOj                  v      	 	 ddՄ       Z	 	 	 	 	 	 ddքZ eeKj.                        edל	 dd؄       Z	 	 	 	 ddلZ eeKj4                  g eOj                  v      	 	 ddڄ       Z eeKj8                  g eOj                  v      	 	 	 	 ddۄ       Z eeKj<                  g eOj                  v      dd܄       Z eeQj@                        	 	 	 	 	 	 	 	 dd݄       Z eeQjD                        	 	 ddބ       Z eejH                  g eOj                  v      dd߄       Z eeQjL                  g eOj                  v      dd       Z eeKjP                        dd       Z eejT                        	 	 	 	 dd       Z eeQjX                        dd       Z eeQj\                        dd       Z ee'j`                        dd       Z ee'jd                        dd       Z ee'jh                        dd       Z ee'jl                        dd       Z ee'jp                        dd       Z eeQjt                        	 	 	 	 d d       Z eeQjx                        dd       Z eej|                  g eOj                  v      	 	 dd       Z ee'j                        dd       Z ee'j                        dd       Z eej                        dd       Z ee7j                        	 	 	 	 d!d       Z ee/j                        dd       Z eeQj                        d"d       Zyc c} w (#  z;Module for lowering JAX to Mosaic-compatible MLIR dialects.    )annotations)Callable
CollectionHashableSequenceN)AnyLiteralProtocolSelfTypeVarcast)api_util)lax)	tree_util)ad_util)checkify)config)core)custom_derivatives)	debugging)dtypes)linear_util)literals)mesh)pjit)prng)source_info_util)state)traceback_util)
xla_bridge)is_cloud_tpu_older_than)
shape_poly)export)mlir)partial_eval)control_flow)BranchesPlatforms)
xla_client)ir)arith)cf)func)math)memref)scf)vector)helpers)
primitives)utils)error_handling)random)indexing)RefBitcasterRefReshaper)Array	DTypeLike)foreach)safe_map)safe_zip)
split_list)tpuint32l         ic                     t        j                   t        j                        xr t	         fdt
        D               S )z=Returns whether a dtype should be lowered to a physical type.c              3  J   K   | ]  }t        j                  |        y wNjnp
issubdtype).0tdtypes     Z/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/_src/pallas/mosaic/lowering.py	<genexpr>z+should_physicalize_dtype.<locals>.<genexpr>s   s     I1cnnUA&Is    #)rE   rF   r   extendedanyPHYSICAL_EXTENDED_DTYPESrI   s   `rJ   should_physicalize_dtyperP   o   s4     
nnUFOO, J
I0HI
II    c                    t        | j                        r.t        j                  | j                        }||j                  z  }|S rC   )rP   rI   jax_corephysical_element_avalshape)avalblock_shaperT   s      rJ   _maybe_physicalize_block_shaperX   w   s:    djj)$::4::F(...K	rQ   c                  2    e Zd ZU i Zded<   i Zded<   ddZy)LoweringDynamicShapeEnvzdict[shape_poly._DimExpr, int]dim_expr_to_placeholderzdict[int, shape_poly._DimExpr]placeholder_to_dim_exprc                   t        j                  |      r|S || j                  vrNt        t	        | j                        z
  }|t
        k  rt        d      || j                  |<   || j                  |<   | j                  |   S )NzhToo many dynamic shapes in the input. Mosaic currently only supports up to 128 dynamic dimension values.)rS   is_constant_dimr[   DIM_UPPER_BOUNDlenDIM_LOWER_BOUND
ValueErrorr\   )selfdim_exprnext_vals      rJ   to_placeholderz&LoweringDynamicShapeEnv.to_placeholder   s    )ot333 3t'C'C#DDh	O	# <
 	
 08d""8, 08d""8,''11rQ   N)rd   r   returnir.Value)__name__
__module____qualname__r[   __annotations__r\   rf    rQ   rJ   rZ   rZ      s    <>9><>9>2rQ   rZ   .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<   ded<   ded<   ded<   ded<   ddZed        Zej                  d        Z	y)LoweringContexttuple[int, ...]
grid_sizestuple[Hashable, ...] | None
grid_namesvmapped_dimszSequence[ir.Value] | Noneuser_grid_indicesz,list[tuple[int | pallas_core.Squeezed, ...]]block_shapeszsource_info_util.NameStack
name_stackpallas_utils.MeshInfo | Nonemesh_contexttpu_core.KernelTypekernel_typezmlir.TracebackCachestraceback_cachesboolforward_compatiblezxla_client.Client | NonebackendDynamicShapeReplacementFndynamic_shape_replacement_fnc                .    t        j                  | fi |S rC   dataclassesreplacerc   changess     rJ   r   zLoweringContext.replace       t/w//rQ   c                ,    t        | j                        S rC   )r`   rq   rc   s    rJ   	grid_rankzLoweringContext.grid_rank   s    trQ   c              #  
   K    j                   sd  y  j                   }t         fdt         j                        D              }t	        ||      }t        j                  |      5  d  d d d        y # 1 sw Y   y xY ww)Nc              3  F   K   | ]  \  }}|j                   vs|  y wrC   rt   )rG   idrc   s      rJ   rK   z4LoweringContext.grid_name_context.<locals>.<genexpr>   s'      aAT=N=N4N   !!)rs   tuple	enumeraterq   ziprS   extend_axis_env_nd)rc   rs   valid_grid_sizesgrid_envs   `   rJ   grid_name_contextz!LoweringContext.grid_name_context   st      ??J 0  :/0H		$	$X	.   s   A&B)A7.	B7B <BN)r   r   rg   ro   )
ri   rj   rk   rl   r   propertyr   
contextlibcontextmanagerr   rm   rQ   rJ   ro   ro      s}    ))..<<((,,""((## 990      rQ   ro   c                  4    e Zd ZU ded<   ded<   ded<   d	dZy)
ShapedAbstractValueztuple[jax_core.DimSize, ...]rU   	jnp.dtyperI   r}   	weak_typec                    t         rC   )NotImplementedError)rc   kwargss     rJ   updatezShapedAbstractValue.update   s    
rQ   N)r   r   rg   r   )ri   rj   rk   rl   r   rm   rQ   rJ   r   r      s    	%%	/rQ   r   c                  V    e Zd ZU ded<   ded<   ded<   ded<   ddZed	        Zdd
Zy)LoweringRuleContextro   lowering_contextzSequence[ShapedAbstractValue]avals_in	avals_outz7Sequence[tuple[int | pallas_core.Squeezed, ...] | None]rv   c                .    t        j                  | fi |S rC   r   r   s     rJ   r   zLoweringRuleContext.replace   r   rQ   c                .    | j                   j                  S rC   )r   r~   r   s    rJ   r~   z&LoweringRuleContext.forward_compatible   s      333rQ   c                x    | j                   j                  y| j                   j                  }t        ||||      S )NT)r   r   r!   )rc   yearmonthdayr   s        rJ   r!   z+LoweringRuleContext.is_cloud_tpu_older_than   s9    $$,##++G"4W==rQ   N)r   r   rg   r   )r   intr   r   r   r   )ri   rj   rk   rl   r   r   r~   r!   rm   rQ   rJ   r   r      s8    ##))**GG0 4 4>rQ   r   c                ,   | t         j                  j                  k(  rt        S | x t        j
                  S xt        j                  j                  k(  r t        S xt        j                  j                  k(  r t        j                  S xxt        j                  j                  k(  rn>xt        j                  j                  k(  rn xt        j                  j                  k(  rn n  t        j                  S t        d x\   | S  	 t        d|        )Nrm   zInvalid memory space: )rS   MemorySpaceDeviceANYTPUMemorySpaceVMEMpallas_coreHOSTERRORINDEXKEYSMEMrb   memory_spaces    rJ   !_memory_space_to_tpu_memory_spacer      s     X))000J	    	$	 	 	$	$j	%	 	 	%	%   	&%%%
'+
!
!
'
'
%+
!
!
%
%	&    		 
 
/~>??rQ   c                ^    t        |       }t        j                  j                  d| d      S )Nz#tpu.memory_space<>)r   r)   	Attributeparse)r   tpu_memory_spaces     rJ   !_memory_space_to_mosaic_attributer     s.    6|D			01A0B!D	EErQ   Fc                   t        j                  | t        j                        rt        j                  | t        j
                        rt        j                  j                  d      S t        j                  | t        j                        rt        j                  j                  d      S t        j                  | t        j                        rt        j                  j                  d      S t        t        j                  | t        j                        rt        d|  d      |r*t        j                  | t         j                        rt        } t!        j"                  t        j$                  |             }t'        |t        j(                        r)t        j(                  j+                  |j,                        S |S )N!tpu.dma_semaphore!tpu.semaphorezExtended dtype z is unsupported.)rE   rF   r   semaphore_dtypetpu_coredma_semaphorer)   Typer   	semaphorebarrier_semaphorer   r   rL   r}   BOOL_MEMREF_TYPEr$   dtype_to_ir_typerI   
isinstanceIntegerTypeget_signlesswidth)rI   is_kernel_boundarytypes      rJ   _dtype_to_ir_typer     s   ^^E;667
~~eX334WW]]/00	{44	5WW]]+,,	{<<	=WW]]+,,^^E6??+
w6FG
HHCNN5#((;E 

		% 0
1$bnn%>>&&tzz22KrQ   Tc                   |rt        |j                        rt        |t        j                        r2t        j                  |j                        }|j                  |      }nt        j                  |      }| t        j                  ||j                        }t        | ||||d      S t        |t        j                        r&|j                  t        j                  j                  u r t         j"                  j%                  d      }n|j                  t        j                  j&                  u r t         j"                  j%                  d      }n_|j                  t        j                  j(                  u r t         j"                  j%                  d      }nt+        d|j                   d      t-        t.        j0                        }	t         j2                  j5                  d||		      S t        |t        j                        rf||j6                  }||j8                  }t-        |      }	 | |      }t         j2                  j5                  |t;        |j                  d
      |		      S t        |t
        j<                        rd||j6                  }|st;        |j                  |      S  | |      }t         j>                  j5                  |t;        |j                  |            S tA        |      )N)
inner_avalF)rV   rU   r   r   allow_extended_typesr   r   Cannot allocate .rm   r   Tr   )!rP   rI   r   r   AbstractRefrS   physical_avalr   r   physical_shapeaval_to_ir_typer   AbstractSemaphoresem_typeSemaphoreTypeDMAr)   r   r   REGULARBARRIERrb   r   r   	SEMAPHORE
MemRefTypegetrU   r   r   ShapedArray
VectorTyper   )
r   rV   rU   r   r   r   r   r   r   memspaces
             rJ   r   r   6  sW    6tzzB$))*))$//:jkkZk8m,,T2m%%eTZZ8e7 -!&\.@05	7 7
 h001}}..22234h	(0088	8/0h	(0088	8/0h)$--:;;01I1IJH==RAAe''(}jje&&l0>H(/E==U

t<    h**+}jje
**);= =(/E==$**9KLN N 	D!!rQ   c           	        t        | d      sit        | t              r%t        j                  | t        j
                        } n4t        | t              r$t        j                  | t        j                        } |st        | j                        }t        | t              s.t        j                  | j                  t        j                        r=t        j                  |t        j                   j#                  |t        |                   S t        | t              s.t        j                  | j                  t        j$                        r=t        j                  |t        j&                  j#                  |t        |                   S | j                  t        j(                  k(  r<t        j                  |t        j*                  j#                  t-        |                   S t/        | j                        )NrI   )hasattrr   r   nparrayr@   floatfloat32r   rI   rE   rF   integerr*   constantr)   IntegerAttrr   floating	FloatAttrbool_BoolAttrr}   r   )x	mlir_types     rJ   ir_constantr   m  s$   	G	!S
((1bhh
a	Au	
((1bjj
!a	!!''*I33>>!''2::>>>)R^^%7%7	3q6%JKK!Us~~aggs||D>>)R\\%5%5iq%JKKww#))>>)R[[__T!W%=>>AGG$$rQ   T)kernel_typesensure_mlir_valuesc                    d fd}|S )Nc                b    D ](  }| t         |   <   rt        j                  |f       * | S rC   )lowering_rulesskip_mlir_conversionsadd)ruler{   r   r   prims     rJ   	decoratorz)register_lowering_rule.<locals>.decorator  s>    # 7*.n[!$'!!4"567 KrQ   )r  r   rg   r   rm   )r  r   r   r  s   ``` rJ   register_lowering_ruler    s     
rQ   c                    t        | j                        r7t        j                  |       }|j                  t        | j                        d  S y)Nrm   )rP   rI   rS   r   rU   r`   )rV   r   s     rJ   _get_aval_physical_dtype_shaper
    s=    djj)**40Ms4::011rQ   c                    d }t        |t        j                        rt        |j                        }t        | |||      S )N)rU   r   )r   r   r   r   r   r   )r   rV   rU   r   s       rJ   _get_arg_typer    s?    
 ,e''(4T5F5FGL	"DL
 rQ   c                R    t        | t        j                        r| j                  S | S rC   )r   r   GridDimensionSemanticsvalue)dimension_semantics    rJ    _canonicalize_dimension_semanticr    s'     "H$C$CD###	rQ   )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
y)#MosaicGridMappingz"pallas_core.GridMappingGrid | Nonegridrr   rs   jax_core.Jaxprjaxpr$tuple[pallas_core.BlockMapping, ...]block_mappingsrp   rt   ztuple[ir.Type, ...]scalar_prefetch_typesoperand_typesscratch_types
grid_typesztuple[tuple[int, ...], ...]scalar_prefetch_block_shapesz2tuple[tuple[int | pallas_core.Squeezed, ...], ...]operand_block_shapesscratch_block_shapesrx   	mesh_infozCallable[..., Any]get_grid_indicesc                    |j                    _         |j                   _        | _        |j                   _        |j                   _        t         fdt         j                         D              }|dt        |      z  }t        d |D              }t        |      t        |      k7  r#t        dt        |       dt        |            t         j                        t        |      z   t         j                         k(  s$J d j                  d|d j                          t        |      t         fd	t        t         j                               D               _         j                  j                  D cg c]  }t        t        |j                         }}||j                      }	||j"                     }
||j$                     }t        fd
|	D               _        t        d |	D              }t        t)        t*        |	|             _        g }g }t/        |
 j                        D ]z  \  }}t1        j2                  |j4                        }t        d |j4                  D              }t+        ||      }|j7                  t9        ||             |j7                  |       | t        |       _        t        |       _        t        fd|D               _        t        d |D               _         t9        t0        jB                        ft         j                         z   _"         jG                  |       |jH                   j                  dfd}| _$        y |jH                   _$        y c c}w )Nc              3  F   K   | ]  \  }}|j                   vs|  y wrC   r   )rG   r   grc   s      rJ   rK   z-MosaicGridMapping.__init__.<locals>.<genexpr>  s'      aat7H7H.Hr   )	arbitraryc              3  2   K   | ]  }t        |        y wrC   )r  rG   ss     rJ   rK   z-MosaicGridMapping.__init__.<locals>.<genexpr>  s       01(+    zGLength of grid does not match length of dimension semantics. len(grid)=z, len(dimension_semantics)=z&Misconfigured grid: self.vmapped_dims=z, dimension_semantics=z, self.grid=c              3  T   K   | ]  }|j                   vrt              nd  ! yw)parallelN)rt   next)rG   r   rc   semantics_iters     rJ   rK   z-MosaicGridMapping.__init__.<locals>.<genexpr>  s0      & !"):): :^
J&s   %(c              3  6   K   | ]  }t        |        y wrC   r  rG   rV   r   s     rJ   rK   z-MosaicGridMapping.__init__.<locals>.<genexpr>  s!      ' 	2D9'   c              3  4   K   | ]  }|j                     y wrC   rU   rG   rV   s     rJ   rK   z-MosaicGridMapping.__init__.<locals>.<genexpr>  s      )6

)6s   c              3     K   | ]C  }t        |t        j                        rt        j                  nt        j                  |       E y wrC   )r   r   Squeezedsqueezed_get_block_dim_size)rG   bs     rJ   rK   z-MosaicGridMapping.__init__.<locals>.<genexpr>  sE        ;//0 

..q12   A	Ar4  c              3  6   K   | ]  }t        |        y wrC   r0  r1  s     rJ   rK   z-MosaicGridMapping.__init__.<locals>.<genexpr>  s!       	2D9r2  c              3  l   K   | ],  }t        |t        j                        s|j                  nd  . y wrC   )r   r   r   rU   r5  s     rJ   rK   z-MosaicGridMapping.__init__.<locals>.<genexpr>  s0      & %T8+E+EF

DP&   24c                F    |r| S t        fdt        |       D              S )Nc              3  2   K   | ]  \  }}|vs|  y wrC   rm   )rG   r   idxrt   s      rJ   rK   zHMosaicGridMapping.__init__.<locals>._get_grid_indices.<locals>.<genexpr>'  s!      
AsA\4IC
s   )r   r   )indicesmaybe_include_mapped_dimsrt   s     rJ   _get_grid_indicesz5MosaicGridMapping.__init__.<locals>._get_grid_indices$  s*    $. 
'0
 
 	
rQ   )rC  r}   )%r  rs   r  r  rt   r   r   r`   rb   iterrange_dimension_semanticsinvarsr   r   rV   slice_index_opsslice_block_opsslice_scratch_opsr  maprX   r  r   r   _get_block_shaperW   appendr  r  r  r  r   index_map_grid_avalr  _prepare_mesh_infor"  )rc   r  grid_mappingdimension_semanticsr   r   	user_gridinvarin_avalsscalar_prefetch_avalsoperand_avalsscratch_avalsr  operands_typesr  rV   bmrU   rW   rD  r.  rt   s   `    `              @@rJ   __init__zMosaicGridMapping.__init__  s    !!DI"--DODJ&55D$11D 		* I "*S^;  5H   9~011I''C#.A*B)DF  t  !C(;$<<		A   2t0022I5H4J KII<	  -.N % &s499~&& !D <@::;L;L27 %**-H  %\%A%AB\99:M\;;<M!& ')' "D $) )64)6 $6 (-*,A(	*)+D% Nt':':; /b**2>>:e  >>	 k 34Ek
4d%
H !!+./ ~.D %&: ;D ! D !& &!& !D
 	(+*I*I	
 	DII	DO 	D!$$,
 &&l
 0d*;;dCs   !Nc                2     j                   sd  _        y |t        d      |j                  } j                  9t         fd|D              r%t        d|j                   d j                         t        j                  j                  |       _        y )Nz:Cannot use communication in pallas_call without shard_map.c              3  :   K   | ]  }|j                   v   y wrC   )rs   )rG   arc   s     rJ   rK   z7MosaicGridMapping._prepare_mesh_info.<locals>.<genexpr>9  s     6aQ$//!6s   zECannot shadow axis mesh axis names with grid names. mesh axis names: z, grid names: )	has_communicationr!  rb   
axis_namesrs   rM   pallas_utilsMeshInfo	from_mesh)rc   r   r`  s   `  rJ   rP  z$MosaicGridMapping._prepare_mesh_info/  s    !!dn|
F  J"	6:6	6'~doo5FH
 	
 "**44T:DNrQ   c                     y rC   rm   r   s    rJ   maybe_compress_gridz%MosaicGridMapping.maybe_compress_grid@  s     	rQ   c                     t               }d fd}|j                   | j                                j                  D ]-  }|j                   ||j                  j                               / t        |      S )Nc                J   | j                   D ch c]L  }t        |t        j                        r0j                  r|j
                  j                  vr|j
                  N }}| j                   D ch c]  }t        |t        j                        rd  }}||z  S c c}w c c}w )Ncomms)effectsr   rS   NamedAxisEffectrs   namer   CommsEffect)r  eaxis_name_effectscomms_effectsrc   s       rJ   _get_nonlocal_axis_nameszEMosaicGridMapping.has_communication.<locals>._get_nonlocal_axis_namesI  s     ==8334??affDOO&C &&  ==;223 m 
 ...s   AB1#B )r  r  )setr   r  r  index_map_jaxprr}   )rc   nonlocal_axis_namesrp  rZ  s   `   rJ   r_  z#MosaicGridMapping.has_communicationF  so    %/  7

CD!!   
"2#5#5#;#;
< #$$rQ   c           	         dd}t         j                  j                  t        t         j                  j
                  t        || j                                    S )Nc                    | yd|  dS )Nz##tpu.dimension_semantics<arbitrary>z#tpu.dimension_semantics<r   rm   r)  s    rJ   _get_semanticszAMosaicGridMapping.get_dimension_semantics.<locals>._get_semanticsb  s    	
4(1--rQ   )r)  z
str | Nonerg   str)r)   	ArrayAttrr   rL  r   r   rG  )rc   rw  s     rJ   get_dimension_semanticsz)MosaicGridMapping.get_dimension_semantics`  sE    .
 <<LL 9 9:	
 rQ   N)
r  r  rQ  pallas_core.GridMappingrR  ,Sequence[tpu_core.DimensionSemantics] | Noner   mesh_lib.Mesh | Noner   r   )r   r}  rg   r}   )rg   zir.ArrayAttr)ri   rj   rk   rl   r[  rP  re  	functoolscached_propertyr_  rz  rm   rQ   rJ   r  r    s    **))	66,,$$$$!! ;;JJ33))&&j<j< ,j< H	j<
 !j< %>j<X;"	 % %2rQ   r  c           	        ~| D ]  j                   j                  }j                   j                  }t        |      rat	        j
                  |      }|j                  }t	        j                  ||      j                  t        d |j                  D              z   n|}|j                  t              }t        j                  j                        }|t        j                  j                  k(  rj!                         r|t        j                  j"                  k(  r$fd}	|dk  rt%        d |	       z         |t&        u s|t        j                  j(                  k(  r#j!                         st%        d |	       z         t+        j,                        }
|
d   d   }}|dk\  r|
d   d   }}nd	\  }}|dk\  rR||k(  xs |d
z  dk(  xr ||k(  xs |dz  dk(  }|rd}t+        j.                         rd}t%        d|z    |	       z         |dk(  sJ j                   j                  t0        j2                  k(  rt5        j6                  t8              }nt5        j6                  |      }d|z  }d
|z  }||k(  xs ||z  dk(  }|rt%        d| dt5        j6                  |       d |	       z          y )Nc              3  F   K   | ]  }t        j                  |        y wrC   )r   Blocked)rG   r   s     rJ   rK   z(_check_block_mappings.<locals>.<genexpr>|  s"      4G%&+

a
 4G   !c                     d j                    dj                   d d d j                  j                   d j                  j
                   dS )NzBlock spec for z in pallas_call z has block shape z, array shape z, and index_map z, in memory space z[.
See details at https://docs.jax.dev/en/latest/pallas/grid_blockspec.html#pallas-blockspec)originfunc_src_inforr  r  
block_avalr   )rZ  
debug_infophysical_array_shapephysical_block_shapes   rJ   err_detailsz*_check_block_mappings.<locals>.err_details  so    		{*::;S;S:T U!%&n5I4J K!11778 9 mm889 :ll mrQ      zEThe Pallas TPU lowering currently supports only blocks of rank >= 1. 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).   r  r     r       z In dynamic shape export - your kernel symbolic args must be annotated with constraints where the computation *after* applying any grid mapping is divisible by 8 and 128 respectively. Ex: (mod(floordiv(m_dim, grid_size), 8) == 0))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.     a  The Pallas TPU lowering currently requires that rank 1 block shapes, either 1) the first (and only) dimension of the block shape is equal to the first (and only) dimension of the array shape, or 2) the first (and only) dimension of the block shape is a multiple of the tiling size (z = 128 * (32 // z)) of the array shape. )
array_avalrI   rU   rP   rS   rT   r   rW   r   r`   r   r  r   r   r   r   has_trivial_windowr   rb   r   HBMr   rM  dynamic_shapes_export_enabledrE   r   r   itemsize_bitsr   )r  r   r  rI   array_shaperT   physical_dtyperankr   r  unmapped_bsbs0as0bs1as1evenly_divisible	extra_msgbitwidthpackingtiling_sizerZ  r  r  s     `                 @@@rJ   _check_block_mappingsr  o  s   
  \
bMME--%%K&&<<UC,22n%44[%H^^e 4G*?*E*E4G /G G n(^^#$D 5R]]5O5OPLx++000R5J5J5Lx++555m ax%-() )
 	|x/C/C/G/GG##%8:E-HI I
 ../CDK2 4R 8CqyR"6r":3chc3qy#:'sa '#:%qA  	446N   	
 m
 	
 QYY						)''(89''7hg'Mk*>k(9Q(>2 3> ?$$^45 6 m	
 		
g\
rQ   )r   !dynamic_shape_replacement_enabledc                 34 | j                   j                  d      }|:t        ddd|      r,t        j                         j                  }t        d|       |j                  }	d 3|rt               3	 	 	 	 d03fd}
nd }
t        |j                  | |	       t        |||||
      }|j                          t        j                  j                         }|j                  j                   }t#        j$                  |	j&                        }t        j(                  j+                  |      |d	<   t        j,                  |j                        }t/        ||d
|| j1                         |
||      }|j2                  j5                  |       |j7                  |       g }d }|j8                  }|s't;        d |j                  D              rt=        d      |rt?        |j                        D ]T  \  }}d| }tA        |jB                  jD                        }|tF        u s:|tH        jJ                  jL                  k(  s|tH        jJ                  jN                  k(  r.|j5                  t        jP                  j+                                tS        |jT                  jV                  |jB                  |||| j1                         |
|      }|jY                         sJ |       t[        t]        j^                  |j`                              }|jb                  }|tH        jJ                  jd                  k(  r%|jg                         rt]        jh                  d      }|t[        tk        |jB                  jl                              z  } |
|      }t        jn                  j+                  |      }tq        |t        jr                  j+                  |            }|j`                  D ]_  }tu        |t\        jv                  t\        jx                  t\        jz                  f      r=t=        dt}        |       d|j`                          |j`                  D cg c]  }tu        |t\        jv                         }}t;        |      r|j`                  D cg c],  }tu        |t\        jv                  t\        jx                  f      . }}t        |      st=        d      |j`                  D cg c]*  }tu        |t\        jv                        r|j                  nd, } }t        tZ        t        |        \  }!}"t        j                  j                  d|! d|" d      |d<   |tu        |t\        jh                        st        d| d      |j                  rt=        d      |j                  }#|#dk  s|#dkD  rt        d|# d      |#dk(  rdnd}$t        j                  j                  d |$ d      |d!<   |j5                  t        jP                  j+                  |             |j2                  j5                  |       |j7                  |       W t        j                  j+                  |      |j                   d"<   |D %cg c]  }%|%t\        j                  u rt        n|% }}% |
|      }t        jn                  j+                  |      |j                   d#<   t        j                  j+                  t        j                  j                  d$      t        |j                              |j                   d%<   t        j                  j+                  t        j                  j                  d$      t        |j                              |j                   d&<   |j                         |j                   d'<   |rJ3t        d(      |*|D &cg c]  }&3j                  j+                  |&|&        }'}&ng }'t        tZ        t        j                     |j                  D (cg c]  }(|(j                   c}(      })|)j5                  t        j                  |'t        j                               t        j                  |)      }*i 4|*D ]  }+3j                  |+      4|+<    3j                  j                         D ]I  \  },}-t[        4j                               }.|-|.vs%t        j                  t        j                  d)d*      }/ t        |/t        t        j                         d+   j                        g,      |-ft        |*      g4fd-|*D         j                         }0|*}1t        j(                  j+                  t        |0            |j                  j                   d.t        |,      z   <   dj                  |1      }2t        j(                  j+                  |2      |j                  j                   d/t        |,      z   <   L |S c c}w c c}w c c}w c c}%w c c}&w c c}(w )1NT)optional  r  r  zWPallas TPU requires a libtpu version that's at most a month old. Found version string:
c                ,    t        fd| D              S )Nc              3  n   K   | ],  }t        j                  |      rj                  |      n| . y wrC   )rS   is_dimrf   )rG   rd   "_mosaic_lowering_dynamic_shape_envs     rJ   rK   zNlower_jaxpr_to_module.<locals>.dynamic_shape_replacement_fn.<locals>.<genexpr>  s<        __X& -
;
;H
Es   25)r   )rU   r  s    rJ   r   z;lower_jaxpr_to_module.<locals>.dynamic_shape_replacement_fn  s!        	  rQ   c                    | S rC   rm   r   s    rJ   <lambda>z'lower_jaxpr_to_module.<locals>.<lambda>  s    Q rQ   sym_namemain)mosaic_grid_mappingrk  r{   r~   r   r  r   c              3  >   K   | ]  }|j                            y wrC   )r  )rG   rZ  s     rJ   rK   z(lower_jaxpr_to_module.<locals>.<genexpr>  s!      &("


!!s   zANon-trivial windowing is not supported for grid-free pallas_call.
transform_)rk  r  r{   r~   r   r   )window_boundstransform_indicesz"Unsupported block dimension type: z for block shape: zFAll block dimensions must be Elements or none of them can be Elements.)r   r   z#tpu.element_window<,r   window_kindzUnsupported pipeline mode: r   z=Lookahead is not supported for XLA pipeline emitter lowering.r  z<Only single (1) and double (2) buffering are supported. Got synchronousdouble_bufferedz#tpu.pipeline_mode<pipeline_modewindow_paramsiteration_bounds@   scalar_prefetchscratch_operandsrR  z;Dynamic shape env is None, invariant violated. Unreachable?)rU   dim_vars)static_argnameskeep_unusedr   )	platformsc              3  (   K   | ]	  }|     y wrC   rm   )rG   venvs     rJ   rK   z(lower_jaxpr_to_module.<locals>.<genexpr>  s     /M1A/Ms   z%tpu.dynamic_dimension_mapping_module_z'tpu.dynamic_dimension_mapping_arg_name_)rU   zjax_core.Shaperg   rp   )jmodule_contextget_backendr!   r    platform_versionRuntimeErrorr  rZ   r  r  r  re  r)   Modulecreate	operation
attributesr$   sanitize_name	func_name
StringAttrr   SymbolTablelower_jaxpr_to_funcis_forward_compatbodyrN  insertr  rM   r   r   r   r  r   r   r   r   r  r   DictAttrlower_jaxpr_to_transform_funcrr  r  verifylistr   rM  rW   r  r   r  Bufferedr
  r   DenseI64ArrayAttrdictFlatSymbolRefAttrr   Elementr7  r  r   allpaddingrL  r   r   r   LoweringExceptionuse_lookaheadbuffer_country  dynamic_grid_dimMLIR_DYNAMICr   r   r   r`   r  r  rz  rb   r\   r   rS   r   rH  rV   rE   r@   r"   all_dim_varsrf   itemskeysjaxjitevaluate_shaper#   rx  devicesplatformr   mlir_modulejoin)5r   rQ  r  rR  r{   r   r  r   r  r  r   r  mattrsmodule_namesym_tabfunc_opr  static_gridr  r   rZ  r  r   	mlir_funcrW   r  window_shapeblock_paramsbdis_element_blockis_element_or_squeezed_blockr  pad_lowpad_highr  pipeline_mode_strr:  r%  	grid_varsrT  rH  args_dimvarsrV   placeholderrd   top_level_namesjitted_eval	stablehloarg_namearg_name_strr  r  s5                                                      @@rJ   lower_jaxpr_to_moduler    sP	    ++777F'4T1aI!--/@@
	-.	0  *'+$&)@)B&	 $/  335EzR)
" ))+ii!
++
 
 %"":#7#78+mm''4%
NN1;;''-);;=#?(I	' &&--	..-+		!	!$	# ,8,G,G  K  
<667 ] 2qc"i:
--
$
$& c
!!5!5!9!99!5!5!?!?? 	R[[__.//



"
"
--1!-??A'C	i **55bnnEFk &&m
h2277
7##%#,,Q/ T89Q9QRSSk0=k))--k:l$0044Y?l  "$$k&:&:K<O<OP
 $2b
,R^^,<> 	 %'NN4  %R)<)<= 4 4		 nn(
 rK//1E1EFG(
$ (
 /0#  nn
 %R)<)<=BJJ6I
 
  c7m4&(ll&8&8"7)1XJa8'
]# 
	"-)=)=>!+M?!<  &&#M  %11!|a/!q"  .:Q->MDU(*(:(:!"3!4A6)
_% 2;;??<89ffmmInnY{] | +-,,*:*:=*IG' KOEF[999q@K  /{;K-/-A-A-E-Ek-RG)**,..*<*<nn!!"%s+>+T+T'U+W'&'+->>+=+=nn!!"%s+>+L+L'M,O''( 113 
*+ ')1
G    -
D
D
H
HA
Ni 
 iX!!"U\\$JEUZZ$JF MM(&&y#))<=**62L C J4CCDIc$iJ 
,	C	C	I	I	K, 	SXXZ(o		(gg## 

FCa(8(A(A$B#C
 K|,
 0N/M
 +-	 	
   MMc)n- 	
3c+6FF	
 xx) MMl+ 	
5K8HH	
3,8 
(a4(

@* %Ks$   
!g01g5/g:3!g?7#hh	
)r   c                   t        j                        g j                  j                  } fd}	||	_         t	        j
                  j                  |d|i|	      }
	 |
j                  j                          |
j                  S # t        j                  $ r}t        j                  |      |d }~ww xY w)Nc                 <   t        | g      \  }}j                  |d      }g dgt        |      z  j                  }t	        j
                  j                  j                  d |t        j                         j                  t        j                         
	      }t        |g|| }t        t        j                         sJ        |t#        dt%        t&        j(                              gt        t+        j,                              z  z  }|S )NTrC  rm   ry   r{   r|   r~   r   r   r   r   )r>   r"  r`   r  ro   r  rs   rt   r   	NameStackr!  r$   TracebackCachesjaxpr_subcompr   r   r   r   r   rE   r@   r
  r   )argsgrid_indicesr  jaxpr_indicesarg_block_shapesr   outrV   r   r   r~   r  r{   r  num_grids          rJ   	body_funcz0lower_jaxpr_to_transform_func.<locals>.body_func  s2   $.thZ$@!L/'88 9 M
M"	"		9	9
 '  &&((""$(22--/-%A (% *- *(*CdE--.44. K%6syy%ABCc&t7G  C JrQ   rk  )r`   r  r  ri   r,   FuncOpfrom_py_funcr  r  r)   	MLIRErrorr4    mlir_error_to_verification_error)r  rV   rk  r  r{   r~   r   r   	arg_typesr  r  rm  r  s   `` `````    @rJ   r  r    s     $//0(%%00)   D )=dkk&&	==iH$DLL 
 
 D

9
9!
<!CDs   1B C*C  C)r   r  c          	     V    t        j                        t        j                        g j                  j                  j                  j                  }g j
                  j                  j                   f	d}	||	_         t        j                  j                  |d|i|	      }
|r|
j                  S 	 |
j                  j                          |
j                  S # t        j                  $ r}t!        j"                  |      |d }~ww xY w)Nc                 2  	 t        | g      \  }}}j                  |d      }t        j                  j                  j
                  |t        j                         j                  t        j                         	      }t        |
g|| S )NFr  r  )r>   r"  ro   r  rs   rt   r   r  r!  r$   r  r  )r  r  r  operands_and_scratchr  r   r  r   r   r~   r  r{   r  r  num_scalar_prefetchs         rJ   r  z&lower_jaxpr_to_func.<locals>.body_func  s    :Dx,-;/7L/#7'88 9 M '  &&((""$(22--/-%A %"14H rQ   rk  )r`   r  r  r  r  r  r  r   ri   r,   r  r  r  r  r)   r  r4   r   )r  r  rk  r{   r~   r   r   r  r!  r  r  rm  r  r  r%  s   `` ````     @@@rJ   r  r    s0    $//0(/EEF%%00 (( ((	)77// //
 . )=dkk&&	==iH$&
 <<DLL 
 
 D

9
9!
<!CDs   C: :D(D##D(c                    d fd}|S )Nc           
     p   r
n
fd}t        j                  ||t        j                  d||i             }t	        j
                  || j                        \  }}}|rt        t	        j                  |      }| j                  j                  | j                        }t        ||g|| }	s|	d   S |	S )Nc                      | i |fS rC   rm   )r  kwfuns     rJ   r  z.lower_fun.<locals>.f_lowered.<locals>.<lambda>A  s    #t:Jr:J9L rQ   zmosaic lower_fun)r  rv   r   )lu	wrap_initr   r  petrace_to_jaxpr_dynamicr   r   convert_constvars_jaxprr   r   rv   r  )ctxr  paramsfwrapped_funr  _constsr   r  r*  multiple_resultss             rJ   	f_loweredzlower_fun.<locals>.f_lowered@  s    %LA,,	6&&'91'+R12K 00cllKE1f&&u-E++33%% 4 '
(%
@&
@4
@CVmJrQ   r1  r   rm   )r*  r7  r8  s   `` rJ   	lower_funr:  ?  s    " 
rQ   c                      e Zd Zy)r  Nri   rj   rk   rm   rQ   rJ   r  r  T      rQ   r  c                r    d}t        t        | |            D ]  \  }\  }}||k(  r|dz   } n | |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)old_name_stacknew_name_stackcommon_prefix_idxr   oldnews         rJ   _compute_name_stack_updatesrE  X  s`      NN!KL ma#s
czA#	
 
)*	+^<M<N-O	OOrQ   c                
    |j                   rJ i  i dfd}d fd}d fd}t        |j                  | j                        D ]
  \  }}||<    t	        ||j                  |       | j
                  j                  D cg c]  }|j                   }	}g }
|
j                  |	       |j                  D ]  }t        ||j                        }| j
                  |j                  j
                  z   }t        j                  | |j                  ||j                  j                        }t!        j"                  |j                  j                        5  |5  |j$                  j&                  5  |j                  t(        | j*                     v r|j                  | j*                  ft,        vrJt        ||j                        D cg c])  \  }}t/        |t1        t2        |j4                              + }}}t        ||j                        }t7        | t1        t8        t2           |j                  D cg c]  }|j4                   c}      t1        t8        t2           |j:                  D cg c]  }|j4                   c}      |      }|j                  D cg c]  }|j                   }}t=        |
|      \  }}|}
|D ]  }t?        j@                           |D ]  }t?        jB                  |d        	 t)        | j*                     |j                     |g|i |jD                  }n0t]        d
| j*                   d|j                  j                   d      |j                  j^                  rt	        ||j:                  |       n ||j:                  d   |       d d d        d d d        d d d         t=        |
|	      \  }}|D ]  }t?        j@                           ta        |      dk(  sJ t        ||j:                        }t        ||j:                        D cg c],  \  }}tc        |td        jf                        rti        |      n|. }}}|S c c}w c c}}w c c}w c c}w c c}w # tF        $ r  tH        $ r}tJ        jL                  jN                  s tQ        |      jR                   d| ddz   d	| dz   }tG        |      }|j                  j                  >|j                  j                  jU                         }tW        jX                  |      |_-        ||d }~ww xY w# 1 sw Y   xY w# 1 sw Y   xY w# 1 sw Y   xY wc c}}w )Nc                ^    t        | t        j                        ry j                  | d       S rC   )r   rS   r	   r   )atomblock_shape_envs    rJ   read_block_shapez'jaxpr_subcomp.<locals>.read_block_shapev  s)    $(()tT**rQ   c                Z    t        | t        j                        r| j                  S |    S rC   )r   rS   r	   val)rH  r  s    rJ   read_envzjaxpr_subcomp.<locals>.read_env{  s%    !$(8(89488Hs4yHrQ   c                t    t        |t        j                  t        f      }|sJ t	        |             || <   y rC   )r   r)   ValueKeyScalarBundler   )varrL  is_valid_typer  s      rJ   	write_envz jaxpr_subcomp.<locals>.write_env~  s2    sRXX$?@M#$s)#=CHrQ   
   )messagelevelz: 
zAdditional diagnostics: 
zFailing jaxpr equation: z3Unimplemented primitive in Pallas TPU lowering for z?. Please file an issue on https://github.com/jax-ml/jax/issues.r   )rH  zjax_core.Atom)rQ  zjax_core.Var)5	constvarsr   rH  rv   r;   rw   stackrk  extendeqnsrL  source_infor$   source_info_to_location	primitive	tracebackr   user_contextr1  managerr  r{   r  _ensure_mlir_valuer   r   rV   r   r   outvarsrE  r?   
trace_stoptrace_startr2  r  	Exceptionr   jax_pallas_verbose_errorsr  r   ri   as_python_tracebackr   filter_traceback__traceback__r   r7  r`   r   rS   r	   r   )!r1  r  r  rJ  rM  rS  rT  bsscopeinitial_name_stackcurrent_name_stackeqninvalseqn_name_stacklocr   r  rv   rule_contextrw   poppedpushedr5  rk  ansrm  msg	new_errortboutvalsrQ  rI  r  s!                                  @@rJ   r  r  o  s    __	
#/+
I
 u||S%5%56  ieROE 	)U\\4(030D0DEu

EE"$./ZZ <'c3::&F^^coo&@&@@N

&
&S]]NCOO,E,EC 
'
'(A(A
B 6'C 6'
''//6'	.9	9MM3??+3HH fcjj1!Q !D)<aff$EF&  +SZZ8*-.0LA0LM-.0MA0MN	
 /=.B.BCUejjC
C4
,' 	A
..
	 	2D
//$b
1	2	!s/>#'*zz#& " 3==#5#5"6 755
 	

 
	'	'	3;;,#++a.#&m6' 6' 6' 6'<'~ /,..&& aNN	V		%--(' /
!S #3(8(89k!nq@'  
.] F 1M0M
 D ! 	
 	!1177G$$%Rs"-,-*3%r23# (,) __&&2**>>@B&4&E&Eb&II#q
 	!A6' 6' 6' 6' 6' 6'@s   QUT;-AT.	.Q3>T.	1Q#%T.	)Q(<T.	Q-)AT.	63Q2)A5T.	T;&U"1UT.	2T+B"T&&T++T.	.T83T;;U UU	c                P   t        | t        j                        r| S t        | t              r| S t        | t        j
                  t        j                  t        t        t        j                  f      rt        | t        |j                              S t        dt        |              )Nz1Unsupported argument to a JAX primitive of type: )r   r)   rO  rP  r   genericndarrayr   r   r   TypedNdArrayr   r   rI   r  r   )rL  rV   s     rJ   rb  rb    s~    RXXJ_%J#

BJJU --/ 0s-djj9::

;DI;G rQ   )r   c               ~   t        j                  ||      }t        j                  || j                  dd        }| j                  ^}}t        j                  ||d d f      \  }}	t        j                  ||d d f      }
| j                  |
| j                  d   gd gt        |
      dz
  z        } t        | g|d|	iS )Nr  r   r   rv   	args_tree)	r   tree_unflattenr   tree_flattentree_leavesr   rv   r`   _load_lowering_rule)r1  reftreerA  indexersindexers_avalsref_avalr5  	args_flatr  
avals_flats              rJ   _get_lowering_ruler    s     %%dC0(++D#,,qr2BC.,(Q"//hd0KL)Y$$hd%KL*$$Q'I4&C
Oa4G*HI 	 	# 
S	B9	B		BBrQ   c                  t        j                  ||      }t        j                  || j                  dd        }| j                  ^}}}	t        j                  |||d f      \  }
}t        j                  |||d f      }| j                  || j                  d   gd gt        |      dz
  z        } t        | g|
d|iS )Nr  r   r  r  r  )	r   r  r   r  r  r   rv   r`   _masked_swap_lowering_rule)r1  r  rL  r  rA  r  r  r  val_avalr5  r  r  r  s                rJ   _swap_lowering_ruler    s     %%dC0(++D#,,qr2BC.<<(Hq"//hT0JK)Y$$40* 	$$Q'I4&C
Oa4G*HI 	 	# 
$C	I)	Iy	IIrQ   c                n   t        | t        t        j                  t        j
                  f      r(t        | t        j                  j                               S | j                  t        j                  j                         k(  r| S t        j                  t        j                  j                         |       S rC   )r   r   r   r}  r   r~  r   r)   	IndexTyper   r   r*   
index_castrv  s    rJ   _make_indexr    su    CX%:%:;<q",,**,--VVr||!!H			",,**,a	00rQ   c                R    | rt        |      S t        |t        j                        S )N)rV   )r  rb  r   rO  )cast_to_indexr   s     rJ   _maybe_cast_to_indexr    s"    q>	AK$C$C	DDrQ   c                x   t        | t              rJ t        | t        j                        r1t	        || j
                        }| j                  }| j                  }d}nXt        | t              rt	        ||       }d}d}d}n5t        j                  |       rt        d|        t	        ||       }d}d}d}||||fS )NFr  Tz+Can only use ()-shaped and slice indexing: )r   slicer6   Slicer  startsizestrider   r   rU   rb   )rA  r  r  r  r  squeezes         rJ   _index_to_start_size_strider    s     U##	#X^^$ 		:E88DZZFG#s 4EDFG	xx}DSEJKK 4EDFG	fg	%%rQ   c               2   t        | j                        }g g g g f\  }}}}|D ]  }|t        j                  d x\   t	        |d      }	d}
d}d}n 	 t        t        |      |      \  }	}
}}|j                  |	       |j                  |
       |j                  |       |j                  |        t        |d       }|J | j                  |f       t        d t        ||      D              }t        |      t        |      t        |      t        |      |fS )Nrm   r   r  Tc              3  *   K   | ]  \  }}|s|  y wrC   rm   )rG   r)  r  s      rJ   rK   z0_indexer_to_start_size_stride.<locals>.<genexpr>T  s      .JAw%,   .s   )
rE  rB  r   r7  r  r  r-  rN  r   r   )indexerref_block_shaper  indices_iterstartssizesstridessqueeze_dimsr)  r  r  r  squeeze_dim
next_indexnew_ref_block_shapes                  rJ   _indexer_to_start_size_strider  5  s.    goo&,)+RR&&%, %a
!$]A6	 "
 +F,
(tV[ MM%	LLNN6$% L$'*		?goo??	 .#e\2J . . FmElGnL
 rQ   c                <   g }t        |       dz
  }t        |      dz
  }|dk\  s|dk\  rc|dk\  r||   nd}|dk\  sJ | |   |k(  r|j                  d       |dz  }|dz  }n | |   dk(  sJ |j                  d       |dz  }|dk\  r]|dk\  rc|j                          |S )Nr  r   r  FT)r`   rN  reverse)source_shapetarget_shaperesultsource_indextarget_index
target_dims         rJ   _compute_squeezed_dimsr  _  s     &\"Q&,\"Q&,\Q./;q/@l+bJ1L!Z/mmEalal,'1,,,mmDal 	\Q. 	..	-rQ   c                   |J t        ||d      \  }}}}}t        d |D              st        d      t        j                  j                         }g }	|D ]^  }
t        |
t        j                        s|	j                  |
       /t        |
      x}|	j                  |       N|	j                  |       ` g }g }|D ]o  }
t        |
t        j                        s|j                  |
       /t        |
      x}|j                  |       N|j                  |       |j                  |
       q t        j                  | j                        }|j                         \  }}||k(  s||	v r|}n%t        t        t        j                   |	|      |      }t        j"                  j%                  ||      }t        j                  j%                  ||j&                  ||j(                        }t+        j,                  || ||      }t/        |      r|}~|j                         \  }}t1        |j2                        D cg c]  \  }}||   r| }}}~t5        |j2                  |      }t1        |      D 
cg c]  \  }}
||   r|
 }}}
t        j"                  j%                  ||      }t        j                  j%                  ||j&                  ||j(                        }t+        j6                  ||      }||fS c c}}w c c}
}w )NFr  c              3  2   K   | ]  }|d u xs |dk(    y wNr  rm   r(  s     rJ   rK   z _slice_memref.<locals>.<genexpr>  s      4qa4i!16!4r*  z-Strided slices of references are unsupported.)r  r  r   r)   
ShapedTypeget_dynamic_sizer   rO  rN  _fold_and_get_constant_valuer   r   get_strides_and_offsetsumrL  operatormulStridedLayoutAttrr   element_typer   r?   memref_slicerM   r   rU   r  memref_squeeze)r  r  	ref_dtyper  r  r  r  r  ir_dynamic_sizestatic_startsr)  r  static_sizesdynamic_sizesref_tyref_strides
ref_offsettarget_offset
out_layoutout_tyr  r   dimtarget_sizestarget_stridess                            rJ   _slice_memrefr  t  s    
	$$	$#

 8&%, 
4G4	4
M
NNMM224/- ,aa"1+A.
.!	;1?+, ,- aa"!+A.
.!	;!/*1 =="&"99;+z?"o&F#MHLL-5zM ##''{C*==F''V5H5H& 	fm<#F$;;=K&/&=UFAs\RS_CULU *&,,EL$-k$:RDAq,q/aRNR%%))*nEJ]]	F 

VS
)C	o	 V Ss   ;K:	K:5L L c                v   t        j                  |      }t        j                  |j                        }||k7  r9t        |      dk  rt	        d      |d   t
        j                  u rt	        d      |j                  }t        j                  | j                        }t        j                  j                  |j                  t        |      |j                        }t        |      }	t        |	      dk\  r#|	d   t
        j                  ur|	d   |z  |z  |	d<   t        j                   ||       |t#        |	      fS )Nr  z5Bitcast 1D ref with bitwidth change is not supported.r  zNBitcast a ref whose 2nd minormost dimension is squeezed when bitwidth changes.r   )r   r  rI   r`   r   r   r8  r)   r   r   r   rU   r   r   r  r?   memref_bitcastr   )
r  	bitcasterr  r  src_bitwidthdst_bitwidthnew_ref_dtyper  target_ref_tyr  s
             rJ   _bitcast_memrefr    s6    %%i0,%%ioo6,\!
?a
A  rk222  //-=="&--##oo&&& $ -
 _-	
!#
b
!)=)=
= 	B,.,>  
, 
 rQ   c                   ||j                   k7  rt        d|j                    d|       t        |      dk  rt        d      |d   t        j
                  u s|d   t        j
                  u rt        d      t        j                  |      t        j                  |j                        k7  rt        d| d|j                         t        j                  | j                        }t        j                  j                  |j                  t        |j                         |j                  	      }t        j                   ||       |j                  fS )
Nz!Reshape a ref with dtype change: z vs r  z Reshape 1D ref is not supported.r  r  z=Reshape a ref with squeezed dimension on last two dimensions.z1Reshape a ref with different number of elements: r   )rI   rb   r`   r   r   r8  r   prodrU   r)   r   r   r   r   r   r?   memref_reshape)r  reshaperr  r  r  r  s         rJ   _reshape_memrefr    s9    (.. 

+HNN+;4	{K  	A
@
AAb[111		 4 4	4
G  WW_!88

;O;L Mnn	  =="&--##nn'&& $ - 
,nn
 rQ   c                    |D ]p  }|xt         d x\    t        | |||      \  } }# xt        d x\    t        | |||      \  } }}D t        d x\   t        | |||      \  } }b 	 t        d|        | |fS )Nrm   zUnsupported transform: )	NDIndexerr  r7   r  r8   r  r   )r  r  r  
transforms	transforms        rJ   _transform_refr  
  s     Ii
9;,I 
_  <>*9I+
'Y  =.I 
_  !$;I;"GHHI  
o	rQ   )frozenc                  &    e Zd ZU dZded<   ded<   y)rP  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.
  rp   	key_shapezSequence[ir.OpResult]scalarsN)ri   rj   rk   __doc__rl   rm   rQ   rJ   rP  rP    s       rQ   rP  c                V   |s%g t         j                  | j                        }}||fS t        |d   t               s;t	        j
                  || j                        }t         j                  |      }|}||fS |^ }}|^ }}t        d |j                  D              rt        d      ||fS )Nr  c              3  n   K   | ]-  }t        |t        j                         xr |j                   / y wrC   )r   r2   r  rU   rG   r^  s     rJ   rK   z6_canonicalize_transforms_to_indexer.<locals>.<genexpr>>  s4       !Z--..:177:s   35zCannot do int indexing on TPU)	r  make_trivial_indexerrU   r   r   get_transforms_shaperM   rB  rb   )r  r  transforms_avalsprev_transformsrA  	ref_shaper5  idx_avals           rJ   #_canonicalize_transforms_to_indexerr  /  s    
 
y==hnnMSO 
#	 jni0,,ZHi**95c"o 
#	 !+'n8	 ## 
 899	#	rQ   c                  |j                  |      \  }}}}|j                  | j                        \  }}}}t        |||      \  }	}
|t        | j                  ^}}t        ||j                  ||	      \  }}t        j                  |j                        }t        |j                        dk(  }| j                  \  }t        |j                  t        j                        rEt!        j"                  |j                  j$                        r|st'        d      t)        | g|d|iS t+        |j                        rt-        j.                  |j                        }t1        t2        |
      }
|
j4                  r
t               |j6                  D cg c]  }t9        j:                  d|       }}t3        |
j<                  t?        |      z   |
j6                  |j6                  z   d      }
|j                  }t-        j@                  |j6                  |j                        }n|j                  }|j6                  }|s|st        d      tC        |
|d	      \  }}}}}tE        d
 |D               }|rE| j                  d   j6                  rt'        d      tG        | |tI        jJ                  ||            S t        |j                        dk7  r*d}t        |j                        dk(  rd}t'        d|z         t-        jL                  ||      }|r9tO        jP                  tS        | jT                  jV                  |d      |||      }n7tY        jJ                  tS        | jT                  jV                  |d      ||      }||k7  rr|rBt        jZ                  j]                  |t_        |d            }tY        j`                  ||      }n.tY        jb                  |g dgte        |j6                        z        }tG        | ||      S c c}w )N#tpu.memory_space<smem>z}PRNG keys must be loaded from SMEM. Did you set the memory space to MemorySpace.SMEM in the BlockSpec for the PRNG key input?r  r   rm   rB  rU   int_indexer_shape7Indexing into a ()-shaped Ref not yet supported on TPU.Tr  c              3  2   K   | ]  }|d u xs |dk(    y wr  rm   r(  s     rJ   rK   z&_load_lowering_rule.<locals>.<genexpr>{        ?!d,a1f,?r*  zCan only load scalars from SMEM#tpu.memory_space<vmem>r  #tpu.memory_space<any>8 ANY memory space can only be accessed using async_copy.z3Loads are only allowed on VMEM and SMEM references.rO   r   )3	unflattenr   r  r   rv   r  rI   r)   r   r   rx  r   r   r   r   KeyTy	pl_randomis_pallas_impl_implrb   _prng_key_load_lowering_rulerP   rS   rT   r   r  r  rU   r6   r  rB  r   r   r  r  _maybe_cast_load_to_boolr.   loadr   r?   strided_loadr   r   r   r0   r   r   r   
shape_castextractr`   )r1  r  r  r5  r  r  maskr  r  r  rA  r  ref_typeis_smem_loadaval_outrT   r  
elt_slicesphysical_out_dtypephysical_out_shaper  r  r  need_strideextra	load_avalload_valvec_types                               rJ   r  r  F  s   &00;#z4%.%8%8%F"(a<
,/3 

((/A'	8>>?O# ]]388$(X**+/HH,+8

+	0H0Hnn1  ; < < (MiM9MMhnn-$::8>>J
y#
C
!!,A,G,GI$(q$IJ I
eJ//ii/555C
 /44!00 "!	o
AC C!>	"&%!Q
 ?w???+
}}Q899#C6;;sF3KLL
8  !%>>E
8  !%==He
=E  ""50BC)  ==#	

 		H {{  ==#	

 	H (""#5"34FGK#MNh ""8X6h"qcC	4H.HIh	!#x	::AIs   Pc                  |j                  |      \  }}}}|j                  | j                        \  }}}}t        |||      \  }}	| j                  \  }
t	        |
j
                  t        j                        sJ |
j
                  j                  j                  }| j                  ^}}t        t        |	      }	t        j                  |j                        }t!        ||j
                  ||      \  }}t#        |      dk7  rt%        d      |d   dk7  rt%        d      t'        d |	j(                  D              st%        d      |dd	 |k(  sJ d
|d|       g }t+        |d         D ]l  }t-        d |D              }t        g |	j.                  d||d      }t1        ||d      \  }}}}}|j3                  t5        j6                  ||             n t9        |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.
  r  zSeed key_data must be 1D.r   r  z-Leading dimension of seed key_data must be 1.c              3  &   K   | ]	  }|d k(    ywr  Nrm   r(  s     rJ   rK   z/_prng_key_load_lowering_rule.<locals>.<genexpr>  s     'Q!V's   z$Can only load a single key per load.r  Nzref_block_shape=z key_shape=c              3  F   K   | ]  }|t         j                  us|  y wrC   )r   r8  )rG   r  s     rJ   rK   z/_prng_key_load_lowering_rule.<locals>.<genexpr>  s"      #[5I5I*Is   !!rm   r  Tr  r  r  )r   r   r  r   r   rI   r   r  r  r  rv   r   r  rS   r   r   r  r`   r   r  rU   rF  r   rB  r  rN  r.   r  rP  )r1  r  r  r  r  r5  r  r  r  rA  r  r  r  r   load_opsr   r  
scalar_idxr  s                      rJ   r  r    s    $--i8#z1a%.%8%8	ll&"(a =
,/3 +8	HNNDJJ	//	/nn"",,)((/AY#%%h&9&9:*'	:_o# 	^q
9
::q\Q
M
NN	'SYY'	'
D
EE			*N/?.@),NN	*(1 .a & I $#++$q$!$IJ 7FAq!Q
 OOFKKV,-. 
U95E	FFrQ   c                   |j                   t        j                  k7  r|S t        t              }t
        t        j                     }t        j                  j                  t        j                  j                  d      |      }t        j                  j                  |d      }|j                  rmt        | j                  j                   |d      }t#        j$                  |t        j&                  j)                  ||            }t#        j*                  |||      S t#        j$                  ||      }t#        j*                  |||      S )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.
  r  r   Tr   )rI   rE   r   r   r   _cmpsi_lowering_typesr   ne_pr)   r   r   r   r   rU   r   r   r   r*   r   DenseElementsAttr	get_splatcmpi)	r1  out_avalrL  load_scalar_typepred	predicate
const_zeroload_vector_typevector_zeross	            rJ   r  r    s      ^^syy J&'78	sxx	($nn  !<!<R!@$G)~~!!"2A6*^^&99
 >>
&&'7DL ::il33 0*=J::ij11rQ   c                    |j                   t        j                  k7  r|S t        | j                  j
                  |d      }t        j                  ||      S )zACasts a boolean value back to an integer for storing in a memref.Tr   )rI   rE   r   r   r   r   r*   extui)r1  expected_avalrL  int_out_types       rJ    _maybe_cast_store_to_memref_typer/    sM     CII%J 	77,
 
\3	''rQ   c                
   |j                  |      \  }}}}|j                  | j                        \  }}	}
}t        |||	      \  }}|c|
j                  j                  dk7  rt        d      |
j                  |j                  k7  r&t        d|
j                   d|j                   d      | j                  ^}}t        ||j                  ||      \  }}t        j                  |j                        }t        |j                        }|dk(  }|dk(  }| j                  \  }t!        |t        j"                        s t%        |t'        |
j                              }|s|st        d	      t)        ||d
      \  }}}}}t+        d |D               }|rm|t        d      |
j                  rt        d      t-        j.                  ||      }t1        | |
|      }t3        | |
|      }t-        j4                  |||       |S |sd}|dk(  rd}t        d|z         |
j                  st        d      t7        |j                        }t9        |j:                        D ]2  \  }}t!        |t<        j>                        r!|jA                  |d       4 tC        |      }|D cg c]!  }|tD        jF                  u rdn
tI        |      # }}|jK                  tM        |      tO        jP                               }| jR                  jU                  |j                        }t        jV                  jY                  |t'        |j                  d
            } |rt[        j\                  | |||      }nt_        j.                  | ||      }t3        | |
|      }||k7  r|j                  st        d      t        jV                  jY                  |j                  t'        |j                  d
            }!t_        j`                  |!|      }t        jV                  jY                  |j                  t'        |j                  d
            }"t_        j`                  |"|      }|St        jV                  jY                  |j                  t'        |j                              }#t_        j`                  |#|      }t1        | |
|      }|r'|t        d      t[        jb                  ||||       |S t[        jd                  |||g |       |S c c}w )N   z masked swap with non-32-bit datazDExpected value and mask to have the same shape, but got value shape z vs. mask shape r   r  r  r  r  Tr  c              3  2   K   | ]  }|d u xs |dk(    y wr  rm   r(  s     rJ   rK   z-_masked_swap_lowering_rule.<locals>.<genexpr>7  r  r*  z!SMEM store does not support maskszCan only store scalars to SMEMr  r  r  z>Loads and stores are only allowed on VMEM and SMEM references.zCannot store scalars to VMEMr  )rU   shardingr   zCannot swap scalars to VMEM.zmasked swap with strided store)r  )3r   r   r  rI   itemsizer   rU   rb   rv   r  r)   r   r   rx  r   r   r   rO  r   r   r  r  r.   r  r  r/  storer  r   rB  r2   r  r  rE  r   r8  r-  r   r   rS   get_cur_mesh_shardingr   r   r   r   r?   r  r0   r	  strided_storevector_store)$r1  r  r  r5  r  r  rL  r  r  r  r  	mask_avalr  rA  r  r  r   is_smem_storeis_vmem_storer  r  r  r  r  r  mem_slice_shaper   r^  mem_slice_shape_iterr:  mem_avalmem_aval_shapemem_aval_vec_typeresult_vec_typeval_vec_typemask_vec_types$                                       rJ   r  r    s     )229=#z34=4G4G	ll51(h	 =
,/3 
1$ BCC~~("..))9)//9J!M 
 ((/A'	8>>?O# ]]388$(X**+,";;-";;-+8	C	"
c%6x~~%F
GC	
AC C ;	&!Wa
 ?w???+:;;~~788[[f%F%c8V<F
*3#
>C
LLc6"M	E//He
H5P 
 

3
44(/$ #daa))*Q"# o. 
 $$$a$/C*DD/  __/"X-K-K-M  ( ''DDnn. mm'''4P /fgFF[[*C8F(h<#>>566mm''4@BO7F==$$X^^4@BL


L#
.Cmm''
..+IOO<m }d3d#C6:& @AAc30 
- S#vr5	-Ss   &S;)r   c               B    ~ |D ]  }t        j                  ||      } |S rC   )r?   assume_multiple)r1  rL  valuesmultiples       rJ   _multiple_of_lowering_rulerH    s-     
 -h


c8
,C-	*rQ   c                     d fd}|S )Nc               R   | j                   \  }| j                  d   j                  sfd}t        |d      } || ||      S t	        j
                  |j                  t        j                        rgt        j                     }t        j                     }t        j                  j                  t        | j                  j                  |d      |      }n|j                  t        j                  k(  rdt        j                     }t        j                     }t        j                   j                  t        j"                  j%                  d      |      }nRt	        j
                  |j                  t        j&                        rt)        d	      t)        d
|j                   d      t        | j                  j                  | j                  d         }	t        j*                  j-                  |	|      }
t/        j0                  |	|
      }t3        j4                  ||||      S )Nr   c                   | t         j                  df   } |D cg c]  }|dz   	 }} | |d      } t        j                  |       S c c}w )N.r  T)axiskeepdims)rE   newaxisr  )rL  axesrL  	reduce_fns      rJ   
_proxy_funz@reduce_lowering_rule.<locals>._lowering_rule.<locals>._proxy_fun  sQ    #++s"#%)*Tq**$6 {{3	 +s   A	Fr7  rO  rm   r4  r  z2Reductions over unsigned integers not implemented.zReductions over z not implemented.)r   r   rU   r:  rE   rF   rI   r   r)   r   r   r   r   r   r@   signedintegerr   r   r   unsignedintegerr   r!  r"  r*   r   r0   multi_reduction)r1  r   rO  r   x_avalrQ  proxy_loweringkindrL  out_typeidentityaccrP  type_to_identitytype_to_kinds               rJ   _lowering_rulez,reduce_lowering_rule.<locals>._lowering_rule  s   IV==!!
  !
u.nC..
~~fllCLL1#,,'dS\\*cLL
""??
 c 
	"#++,dS../cNNr~~::2>Dc	c&9&9	:
>   V\\N*;
<> >993==;KH ##--h<H
..8
,C!!$355rQ   r9  rm   )rP  r^  r]  r_  s   ``` rJ   reduce_lowering_ruler`    s    -6\ 
rQ   z-infinf        c               8    d }t        |d      } || ||      S )Nc               d    t        j                  | dd      }t        j                  ||      dkD  S N      ?rb  rL  )rE   whereminargrO  	float_args      rJ   _proxy_reducez0_reduce_and_lowering_rule.<locals>._proxy_reduce  ,    
 		#sC(I7794(3..rQ   FrR  rS  r:  r1  r   rO  rm  rX  s        rJ   _reduce_and_lowering_rulerq    '    / e-.	QT	**rQ   c               8    d }t        |d      } || ||      S )Nc               d    t        j                  | dd      }t        j                  ||      dkD  S re  )rE   rh  maxrj  s      rJ   rm  z/_reduce_or_lowering_rule.<locals>._proxy_reduce  rn  rQ   FrR  rS  ro  rp  s        rJ   _reduce_or_lowering_rulerv    rr  rQ   c                    t        d      )Nz``broadcast_to` is a Triton-specific primitive. Please consider using `jnp.broadcast_to` instead.)r  )r1  r   rU   s      rJ   _broadcast_to_lowering_rulerx    s     	%	 rQ   c                  ~| j                   \  }| j                  \  }|j                  |k(  r|S t        j                  |j
                  t        j                        r;| j                  s| j                  ddd      rd }t        |d      } || |||      S |rdgt        |      z  }	t        ||j                        D ]
  \  }
}||	|
<    t        |	      }t        j                  j                  |t!        |j
                              }t#        j$                  ||      }||j                  k(  r|S t        j                  j                  |j                  t!        |j
                              }t#        j&                  ||      S )	Nr        c               |    t        j                  | dd      }t        j                  j	                  |||      }|dk(  S )Nr  r   )rE   rh  r  r   broadcast_in_dim)rL  rU   broadcast_dimensionsint_val	bcast_vals        rJ   rQ  z3_broadcast_in_dim_lowering_rule.<locals>._proxy_fun  s7    		#q!$g''**7E;OPi!^rQ   FrR  )rU   r~  r  )r   r   rU   rE   rF   rI   r   r~   r!   r:  r`   r   r   r)   r   r   r   r0   r	  	broadcast)r1  rL  rU   r~  r3  aval_inr  rQ  rX  out_shape_listr   r)  	out_shaperZ  s                 rJ   _broadcast_in_dim_lowering_ruler    sT    ||*7+8]]eJ^^GMM399-	 ; ;D!Q G

 zEBNS4H  S3u:%N('--8 1nQn%I}}  $X^^4H 

Hc
*CHNN"j]]nn'7( 
		(C	((rQ   c                   | \  }}|\  }}|\  }}t        t        t        |                  }	t        t        t        |                  }
t        |	t        |      z
  t        |      z
        }t        |
t        |      z
  t        |      z
        }g }t	        t        t        |                  D ci c]  \  }}||
 }}}t	        t        t        |                  D ci c]  \  }}||
 }}}|D ]'  }|j                  d       |j                  ||          ) |D ]'  }|j                  d       |j                  ||          ) |D ]'  }|j                  d       |j                  ||          ) d }|||||||f}ddj                  t        ||             d}t        j                  j                  |      S c c}}w c c}}w )a  Converts a jax dot dimension numbers to a tpu dot dimension numbers.

  Jax dot dimension numbers are given as a tuple of tuples of sequences of ints
  of the form ((lhs_contracting_dims, rhs_contracting_dims), (lhs_batch_dims,
  rhs_batch_dims)).

  TPU dot dimension numbers are given as an MLIR definition of the form
  #tpu.dot_dimension_numbers - which can be found in the tpu dilect definition
  # file, tpu.td .
  r   r  c                >    ddj                  d | D              z   dz   S )N[, c              3  2   K   | ]  }t        |        y wrC   )rx  rG   r   s     rJ   rK   zHjax_dot_dims_to_tpu_dot_dot_dims.<locals>.format_dims.<locals>.<genexpr>i  s     0a3q60r*  ])r  )dimss    rJ   format_dimsz5jax_dot_dims_to_tpu_dot_dot_dims.<locals>.format_dimsh  s!    04000366rQ   z#tpu.dot_dimension_numbers<r  r   )rq  rF  r`   sortedr   rN  r  rL  r)   r   r   )dimension_numbers	lhs_shape	rhs_shapecontracting_dims
batch_dimslhs_contracting_dimsrhs_contracting_dimslhs_batch_dimsrhs_batch_dimslhs_total_dimsrhs_total_dimslhs_non_contracting_dimsrhs_non_contracting_dimsoutput_dim_orderrA  r  lhs_dim_maprhs_dim_mapr  all_dimstpu_dim_numbers_strs                        rJ    jax_dot_dims_to_tpu_dot_dot_dimsr  :  s    $5 Z/?,,#- ..uS^,-.uS^,-.#s/003~3FF $s/003~3FF *3E#i.4I*JKhc3cK+K*3E#i.4I*JKhc3cK+K .cAK,-. & .cAK,-. & .cAK,-.7 ( 	$CHHSh-G$H#IK  
		/	00? LKs   2G "Gc                   |\  \  }}}| j                   \  }	t        | j                  j                  |	      }
t        j
                  j                  |
      sJ t	        j
                  |
      j                  t        fdt        j                  t        j                  t        j                  t        j                  t        j                  fD              r!t        j                  j                  d      }nbt        j                   j                        r!t        j"                  j                  d      }n"t%        | j                   d   j&                        | j(                  \  }}|dk(  r|dk(  r| j(                  d   j*                  d   dk(  rt-        | j(                  d   j*                        dk(  rt-        | j(                  d   j*                        dk(  re|j&                  t.        j0                  k7  s=|j&                  t.        j0                  k7  s | j2                  s| j5                  ddd	      r| j(                  d   j*                  | j(                  d   j*                  k7  rt/        j6                  | j(                  d   j*                  | j                   d   j*                        }t        j8                  j                  t;        |      t=        | j                   d   j&                              }| j(                  d   j*                  |k7  rt?        j@                  ||      }| j(                  d   j*                  |k7  rt?        j@                  ||      }|r|n|j&                  }t        | j                  j                  |jC                  |j*                  d   f|
            }|j&                  |k7  rst        | j                  j                  |jC                  |j*                  |
            }|t.        j0                  k(  rtE        jF                  ||      }nt%        d|      |j&                  |k7  rst        | j                  j                  |jC                  |j*                  |
            }|t.        j0                  k(  rtE        jF                  ||      }nt%        d|      tE        jH                  |t        jJ                  jM                  ||            }t?        jN                  t        jP                  jS                  d      tE        jT                  ||      |dg      }t?        jV                  |
|      S tY        ||j*                  |j*                        }||d   |d   k7  rt%        d      |d   }||tZ        j\                  j^                  k(  rd }nK|tZ        j\                  j`                  k(  r t        jP                  jS                  d      }nt%        d|       tE        jH                  |
t        jJ                  jM                  |
|            }tc        jd                  |
|||||      S )Nc              3  @   K   | ]  }|j                          y wrC   )r   )rG   clsval_types     rJ   rK   z-_dot_general_lowering_rule.<locals>.<genexpr>  s"      		
 
nnX		s   rb  r   r  r  r  r  r  rT  )rU   rI   z#Unsupported preferred_element_type=z#vector.kind<add>z%Per-operand dot precision unsupportedz#tpu.contract_precision<fp32>zUnsupported dot precision: )r  	precision)3r   r   r   r   r)   r  r   r  rM   BF16TypeF32TypeFloat8E5M2TypeFloat8E4M3FNTypeFloat8E4M3B11FNUZTyper   r   r   r   r   rI   r   rU   r`   rE   r   r~   r!   broadcast_shapesr   r  r   r0   r  r   r*   extfr   r!  r"  rV  r   r   mulfr	  r  r   	PrecisionDEFAULTHIGHESTr?   matmul)r1  r   yr  r  preferred_element_typer5  lhs_dimsrhs_dimsr  rZ  rL  lhs_avalrhs_avalbcast_shape	red_dtypered_typelhs_typerhs_typer\  redtpu_dot_dimsprecision_attrout_tiler  s                           @rJ   _dot_general_lowering_ruler  {  s    .8X+8	77( 
	!	!(	++	+]]8$11( 		 ++
**







"
"		 	 ,,

8S
)C	~~  *
..

Xq
)C
cmmA.44
55||(H $
d

,,q/


"a
'
cll1o##
$
)
cll1o##
$
)
..CKK
'^^s{{*##((q"5 ||AQ 5 55((
,,q/

q!1!7!7k MM%%
{
.s}}Q/?/E/EFk 
a		+	-[!,	a		+	-[!,"8hnn  99x~~a02)DH
 ~~" 



;
;
//i/
@h 
ckk	!JJx#!$H1G0I"JKK~~" 



;
;
//i/
@h 
ckk	!JJx#!$H1G0I"JKK
.."&&003?C 
 
 
./

1a	
	C Xs++1, |y|# GHH!I)s}}'<'<<NCMM)))\\'''N  ;I;G
HH^^$$..x=( 
$
 rQ   c                  | j                   }t        j                  |      }t        j                  |      }|t        j                  k(  r,| j                  t        j                        } t        | |      S |t        j                  k(  rt        j                  |t        j                        r%|dk  rh| j                  t        j                        } nHt        j                  |t        j                        r$|dk  r| j                  t        j                        } | t        j                  d| j                         k7  S t        j                  |t        j                        r}|dk  r| j                  t        j                        } t        j                  |t        j                        r$|dk  r| j                  t        j                        } | j                  |      S t        j                  |t        j                        rY|dk  r| j                  t        j                        } t        j                  |t        j                        s| j                  |      S t!        d| d|       )Nto_dtyper  r   rO   zUnsupported cast:  -> )rI   r   r  rE   r   astyper@   _convert_helperrF   r   r   r   asarrayrT  rU  uint32r   )r   r  
from_dtypefrom_bitwidthto_bitwidths        rJ   r  r    s   ww*&&z2-$$X.+399	A1x00 ~~j#,,/		HHS[[!	
CKK	0		HHSYYAQWW---^^J 1 12r
((399
a
~~h-+2B
((3;;
a88H^^J 3 34r
((3::
a>>(CLL1XXh0D
KLLrQ   c                  ~~| j                   d   }| j                  d   }|j                  t        | j                  j
                  |      }k(  r|S j                  dk(  rt        d      fd}fd}	t        j                  }
t        j                  }t        j                  }t        j                  }t        j                        }t        j                        }|dk(  xr |dk(  } ||
      ro |	|
      rg| j                  xs | j!                  ddd	      }||k  r|dk(  s|st#        j$                  ||      S ||kD  ry|dk(  s|sqt#        j&                  ||      S  ||      r |	|      r||k  rB|dk(  r= ||      rt#        j(                  ||      S  ||      rt#        j*                  ||      S ||kD  r|dk(  rt#        j,                  ||      S t        j.                        j0                  t        j.                        j0                  k(  r|S  ||
      r |	|      rt#        j2                  ||      S  ||      r? |	|
      r7| j                  s| j!                  dd
d      r|rLt#        j4                  ||      S t        j6                  k(  r# |	|      r|dk(  rt#        j(                  ||      S  t9        t;        j<                  t>              d      | |      S )Nr   r  z64-bit types are not supportedc                0    t        j                  |       S rC   rD   )rI   	old_dtypes    rJ   r  z5_convert_element_type_lowering_rule.<locals>.<lambda>/	  s    y%8 rQ   c                0    t        j                  |       S rC   rD   )rI   	new_dtypes    rJ   r  z5_convert_element_type_lowering_rule.<locals>.<lambda>0	  s    cnnY6 rQ   r  r  rz           r  FrR  ) r   r   rI   r   r   r   r4  r   rE   r   r   rT  rU  r   r  r~   r!   r*   r  truncfr,  extsitrunciiinfobitsfptosisitofpr   r:  r  partialr  )r1  r   r  r   r3  r$  in_avalrZ  _from_tor   r   signedunsignedold_bitwidthnew_bitwidth
both_32bitforward_compatr  s     `               @rJ   #_convert_element_type_lowering_ruler  	  s    ]]1(LLO'mm)	77( )H1
>
??
8%6#\\(KK'&  (%%i0,%%i0,r!8lb&8*
8_X++ s/J/Ja0N l".ZZ!$$		$.\\(A&&W~#g,l"|r'9	x{{8Q''	v{{8Q''		$);\\(A&&	9		"	"cii	&:&?&?	?hX3v;<<!$$V}X##s'B'B4B'O\\(A&&CII#g,<23E;;x##
+9$$_yI$)
++.
3 3rQ   c           	         |t         t        d |D              rt         | j                  d   j                  sAt	        j
                  t        | j                  j                  | j                  d         |      S | j                  d   j                  s;t	        j                  |g dgt        | j                  d   j                        z        S t	        j                  t        | j                  j                  | j                  d         |      S )Nc              3  $   K   | ]  }|d u  
 y wrC   rm   r  s     rJ   rK   z)_reshape_lowering_rule.<locals>.<genexpr>b	  s     &qd&s   r   )r   rM   r   rU   r0   r  r   r   r   r   r
  r`   r	  )r1  r   	new_sizes
dimensionsr3  s        rJ   _reshape_lowering_ruler  ]	  s     
&I&&
	a		  ==s}}Q?O	
 	
	  
q			>>!R!s3<<?+@+@'A!ABB			



;
;S]]1=M 	
 rQ   c                   ~| j                   \  }| j                  \  }|j                  s`|j                  j                  dk7  rt        d|j                   d      t        j                  |g dgt        |j                        z        S t        j                  t        | j                  j                  | j                  d         |      S )Nr1  zLOnly arrays with 32-bit element types can be converted to scalars, but got: z4. Try casting the input before squeezing the scalar.r   )r   r   rU   rI   r4  rb   r0   r
  r`   r	  r   r   r   )r1  r   r  r  r  s        rJ   _squeeze_lowering_ruler  u	  s    ||*7+8	~~!#~~& ' 
 >>!R!s7=='9!9::			



;
;S]]1=M 	
 rQ   c               2    ~ t        j                  ||      S )N	dimension)r?   concatenate)r1  r  xss      rJ   _concatenate_lowering_ruler  	  s    		y	11rQ   c          
        | j                   \  }t        j                  |j                  t        j                        }t        j
                  |      }t        j                  |      }g }t        || j                        D ]]  \  }	}
|	||<   |j                  t        j                  t        | j                  j                  |
      ||||             ||xx   |	z  cc<   _ |S )NrO   )r   r   r   rU   int64
zeros_like	ones_liker   r   rN  r0   extract_strided_slicer   r   r   )r1  r   r  rL  rW  
slice_sizer  r  outsr  r  s              rJ   _split_lowering_ruler  	  s     ll)6xxBHH5*==$&LL$'	$E3==1 ndHJtKK$$$$AA8 	

 4LDL 
+rQ   c                   t              dk(  r*|dk7  rt        d      fd} t        |d      |       S t        | j                  j
                  | j                  d         }t        j                  ||g      S )Nr  r   z Dimension must be 0 for 1D iota.c                 Z    t         j                  j                  dz   d      } | d   S )Nr  r  )rI   rU   r  r3  r   )r   iota_pbind)iota_2drI   rU   r3  s    rJ   _1d_iota_helperz,_iota_lowering_rule.<locals>._1d_iota_helper	  s5    

e&*Ul*+)1   3g QZrQ   FrR  )r  )	r`   rb   r:  r   r   r   r   r?   iota)r1  rI   rU   r  r3  r  rZ  s    `` `  rJ   _iota_lowering_ruler  	  sw     	Z1_A~9:: >9_u=cBB	77q9I( 
(	{	33rQ   c                  | j                   d   }	| j                   d   }
| j                  d   }t        |	j                        dk7  rt	        d      |	j                  |
j                  d d cxk7  r|j                  k7  rt        d       t        | j                  j                  |      }|
j                  |	j                  dz   k(  sJ t        j                  t        j                  j                  |	j                  |j                  j                        |      }~|dk(  r|t         j"                  j$                  t         j"                  j&                  fv rj|t!        j(                  d	d
d
dd      k(  rt+        j,                  ||dg      S |t!        j(                  d	ddd
d
      k(  rt+        j,                  ||dg      S t	        d      )Nr   r  r  zOnly 2D gather is supportedr  z+Shape mismatch in input, indices and outputr  r  rm   )r   )offset_dimscollapsed_slice_dimsstart_index_mapoperand_batching_dimsstart_indices_batching_dimszUnsupported gather)r   r   r`   rU   r   rb   r   r   r   r0   r	  r)   r   r   r   r  r   GatherScatterModeFILL_OR_DROPPROMISE_IN_BOUNDSGatherDimensionNumbersr?   dynamic_gather)r1  r   rB  r  slice_sizesunique_indicesindices_are_sortedmode
fill_valuer  indices_avalr$  rZ  recovered_indicess                 rJ   _gather_lowering_ruler  	  s    LLO'a,]]1(1
;
<<]]l(("-??
B
CC @ 	77( 
		w}}t3	33	3''mmw||'@'@A
 V




,
,



1
1
 C66!"$(  #4qc::C66!"$(  #4qc::011rQ   c                   t        | j                  j                  | j                  d         }| j                  s| j                  ddd      rt        j                  |||      S t        j                  |||      S )Nr   r  r  r  )	r   r   r   r   r~   r!   r0   	transposer?   )r1  r   permutationrZ  s       rJ   _transpose_lowering_ruler  	  sj    	77q9I( 	s::4AFHa55==1k22rQ   c                   |j                   }|j                   }|j                  r|j                   }n|j                  r|j                   }t        | t        j                        sHt        |dd       x}t        j                  j                         k(  r|}nt        |      }t        | |      } t        |t        j                        sHt        | dd       x}	t        j                  j                         k(  r|	}nt        |      }t        ||      }t        |j                        }
|j                  |j                  k7  r?t        j                  j                  |
t        |            }t        j                  ||       } |j                  |j                  k7  r?t        j                  j                  |
t        |            }t        j                  ||      }| |fS )Nr   )rI   r   r   r)   rO  getattrr  r   r   r   r  rU   r   r0   r  )r   r  rW  y_avalr$  x_dtypey_dtypey_typer   x_typer  x_tyy_tys                rJ   _bcastr'  
  se    LL'LL'llGllG	Arxx	 !VT**r||/?/?/AAi#G,iAy!A	Arxx	 !VT**r||/?/?/AAi#G,iAy!A8>>")\\X^^#==Y(9'(BCDq!A\\X^^#==Y(9'(BCDq!A	
A+rQ   c                   t        ||| j                  d   | j                  d   | j                  d         \  }}| j                  \  }t        j                  |j
                  t        j                        rt        j                  ||      S t        j                  |j
                  t        j                        rt        j                  ||      S t        |j
                        Nr   r  )r'  r   r   rE   rF   rI   r   r*   addir   addfr   r1  r   r  r  s       rJ   _add_lowering_ruler-  *
  s    
 
1cll1os||Aa8H	I$!Q+8^^HNNCKK0::a^^HNNCLL1::aHNN++rQ   c                      e Zd Zy)FoldingErrorNr<  rm   rQ   rJ   r/  r/  8
  r=  rQ   r/  c                @    fd	  | d      S # t         $ r Y y w xY w)Nc                   dk  r
t               t        | j                  dd       }t        t        d}|dk(  rt
        j                  j                  | j                        r6t        j                  | j                  j                  d         j                  S t
        j                  j                  | j                        r6t        j                  | j                  j                  d         j                  S t        d| j                         ||v r) ||   fd| j                  j                  D              S t               )Nr   rk  )zarith.maxsizarith.minsizarith.constantr  zUnsupported constant type: c              3  6   K   | ]  } |d z
          ywr  rm   )rG   r  _foldfuels     rJ   rK   z>_fold_and_get_constant_value.<locals>._fold.<locals>.<genexpr>M
  s     !O%4!8"4!Or2  )r/  r  ownerru  ri  r)   r   r   r   r   r  r  	FloatTyper   rb   operands)r   r4  op_namebinop_foldsr3  s    `  rJ   r3  z+_fold_and_get_constant_value.<locals>._fold=
  s    qyNaggvt,GK ""		"	"166	*~~agg009:@@@<<""166*||AGG..w78>>>6qvvh?@@+![!!Oagg>N>N!OOO
.rQ   rT  )r/  )r   r3  s    @rJ   r  r  <
  s+    &B<	 s    	c                    |S rC   rm   )r5  r   s     rJ   _stop_gradient_lowering_ruler;  V
  s    	
(rQ   )r   r   c                R   t        ||| j                  d   | j                  d   | j                  d         \  }}| j                  \  }t        j                  |j
                  t        j                        rt        j                  ||      S t        j                  |j
                  t        j                        rt        j                  ||      S t        j                  |j
                  t        j                        rt        j                  ||      S t        |j
                        r)  )r'  r   r   rE   rF   rI   rT  r*   maxsirU  maxuir   maximumfr   r,  s       rJ   _max_lowering_ruler@  Z
       
1cll1os||Aa8H	I$!Q+8^^HNNC$5$56;;q!
~~hnnc&9&9:;;q!
~~hnncll3>>!QHNN++rQ   c                R   t        ||| j                  d   | j                  d   | j                  d         \  }}| j                  \  }t        j                  |j
                  t        j                        rt        j                  ||      S t        j                  |j
                  t        j                        rt        j                  ||      S t        j                  |j
                  t        j                        rt        j                  ||      S t        |j
                        r)  )r'  r   r   rE   rF   rI   rT  r*   minsirU  minuir   minimumfr   r,  s       rJ   _min_lowering_rulerF  i
  rA  rQ   c                   | j                   \  }| j                  \  }|j                  t        j                  k7  rt        d      t        |      dk7  rt        d      |t        j                  k7  rt        d      |d   }t        |j                        dk(  }|rt        j                  dg|j                  |j                        }	t        | j                  j                  |	      }
t        j                  dg|j                  |j                        }t        j                  |
|      }|dz  }t        | j                  j                  |      }t!        j"                  ||||      }|rt        j$                  |g dg      S |S )NzOnly float32 is supportedr  z$Only single axis reduction supportedz#Only index_dtype=int32 is supportedr   )r   r   rI   rE   r   r   r`   r@   rU   rS   r   r   r   r   r0   r	  r?   reduce_indexr
  )r1  r   rO  index_dtypereduction_kindrW  r$  rL  is_1d	x_2d_aval	x_2d_typerZ  r  s                rJ   _reduce_index_helperrN  w
  sN   ll)6+8\\S[[ 
9
::Y!^
D
EECII
C
DD	a$
fll
q
 %
$$a%7&,,%7FI999I ##Q$8$8(..IH)Q'AAID	77( Ha~>&
>>&"qc**	-rQ   c           	     Z    t        | |||t        j                  j                  d            S )Nz#tpu.reduction_kind<arg_max>rN  r)   r   r   r1  r   rO  rI  s       rJ   _argmax_lowering_rulerR  
  +    		1dKll78
 rQ   c           	     Z    t        | |||t        j                  j                  d            S )Nz#tpu.reduction_kind<arg_min>rP  rQ  s       rJ   _argmin_lowering_rulerU  
  rS  rQ   c                   t        ||| j                  d   | j                  d   | j                  d         \  }}| j                  \  }t        j                  |j
                  t        j                        rt        j                  ||      S t        j                  |j
                  t        j                        rt        j                  ||      S t        |j
                        r)  )r'  r   r   rE   rF   rI   r   r*   subir   subfr   r,  s       rJ   _sub_lowering_rulerY  
       
1cll1os||Aa8H	I$!Q+8^^HNNCKK0::a^^HNNCLL1::aHNN++rQ   c                   t        ||| j                  d   | j                  d   | j                  d         \  }}| j                  \  }t        j                  |j
                  t        j                        rt        j                  ||      S t        j                  |j
                  t        j                        rt        j                  ||      S t        |j
                        r)  )r'  r   r   rE   rF   rI   r   r*   mulir   r  r   r,  s       rJ   _mul_lowering_ruler]  
  rZ  rQ   c                R   t        ||| j                  d   | j                  d   | j                  d         \  }}| j                  \  }t        j                  |j
                  t        j                        rt        j                  ||      S t        j                  |j
                  t        j                        rt        j                  ||      S t        j                  |j
                  t        j                        rt        j                  ||      S t        |j
                        r)  )r'  r   r   rE   rF   rI   rT  r*   divsirU  divuir   divfr   r,  s       rJ   _div_lowering_rulerb  
  s     
1cll1os||Aa8H	I$!Q+8^^HNNC$5$56;;q!^^HNNC$7$78;;q!
~~hnncll3::aHNN++rQ   c                R   t        ||| j                  d   | j                  d   | j                  d         \  }}| j                  \  }t        j                  |j
                  t        j                        rt        j                  ||      S t        j                  |j
                  t        j                        rt        j                  ||      S t        j                  |j
                  t        j                        rt        j                  ||      S t        |j
                        r)  )r'  r   r   rE   rF   rI   rT  r*   remsirU  remuir   remfr   r,  s       rJ   _rem_lowering_rulerg  
  s     
1cll1os||Aa8H	I$!Q+8^^HNNC$5$56;;q!^^HNNC$7$78;;q!^^HNNCLL1::aHNN++rQ   c                T   | j                   \  }t        j                  |j                  t        j                        rt        j                  |      S t        j                  |j                  t        j                        rt        j                  |      S t        |j                        rC   )
r   rE   rF   rI   r   r-   absir   absfr   r1  r   r  s      rJ   _abs_lowering_rulerl  
  sd    +8^^HNNCKK099Q<^^HNNCLL199Q<HNN++rQ   c                    | j                   \  }| j                  t        j                  d|j                        |fdg| j
                        }t        |t        j                  d|j                        |      S )Nrm   r  r   rO   )	r   r   rS   r   rI   rv   rY  r   r   )r1  r   rW  new_ctxs       rJ   _neg_lowering_rulero  
  sl     ll)6KK$$R6?*))*  ' 
GRXXav||%Da	HHrQ   c                F     t        t        j                  d      | |      S NFrR  )r:  ra  sign_lowering_helperr1  r   s     rJ   _sign_lowering_rulert  
  s(    
''%

 rQ   c                H     t        t        j                  d      | ||      S rq  )r:  ra  nextafter_lowering_helperr1  r   r  s      rJ   _nextafter_lowering_rulerx  
  s*    
,,u
A
 rQ   c                F    |t        d      t        j                  |      S NNot implemented: accuracy)r   r-   rsqrtr1  r   accuracys      rJ   _rsqrt_lowering_ruler     "    
9
::	ArQ   c                F    |t        d      t        j                  |      S rz  )r   r-   sqrtr}  s      rJ   _sqrt_lowering_ruler    "    
9
::	1rQ   c                    t        j                  | j                  d   j                  t         j                        rt        j                  ||      S t        j                  ||      S Nr   )rE   rF   r   rI   r   r*   r\  r  rs  s     rJ   _square_lowering_ruler    sD    ^^CLLO))3;;7::a	Aq	rQ   c                F    |t        d      t        j                  |      S rz  )r   r-   expr}  s      rJ   _exp_lowering_ruler    "    
9
::	!rQ   c                   t        | j                  j                  | j                  d         }t	        j
                  | j                  d   j                  t        j                        rt        j                  ||      }t        |t        j                        s|dk(  rt        j                  |      S t!        ||| j                  d   | j                  d   | j                  d         \  }}t        j"                  ||      S )Nr   r  g       @)r   r   r   r   rE   rF   r   rI   r   r*   r  r   r)   rO  r-   exp2r'  powf)r1  r   r  rZ  s       rJ   _pow_lowering_ruler    s     	77q9I( 	^^CLLO))3;;7Xq!A	Arxx	 Q"W99Q<	1cll1os||Aa8H	I$!Q	1arQ   c               J     t        t        j                  d      | ||      S )NFrR  )r  )r:  lax_internal_integer_powrw  s      rJ   _integer_pow_lowering_ruler  +  s%    	E<,,u	E	1
 rQ   c                    |t        d      | j                  s| j                  ddd      r t        d d      | |      S t	        j
                  |      S )Nr{  r        c                    t        j                  t        j                  t        j                  d      | j
                        | z        S )Nr  )rE   r  r  r   logrI   r  s    rJ   r  z%_exp2_lowering_rule.<locals>.<lambda>9  s+    #''#**RVVAY81<= rQ   FrR  )r   r~   r!   r:  r-   r  r}  s      rJ   _exp2_lowering_ruler  1  sa    
9
::s::4BG9= 
1  
1rQ   c                   |t        d      t        j                  |      }t        j                  |      }| j
                  d   }t        | j                  j                  |      }|j                  st        d|      }nt        j                  |t        d            }t        j                  ||      }t        j                  ||      S )Nr{  r   rf  r  )r   r*   negfr-   r  r   r   r   r   rU   r   r0   r  r+  ra  )	r1  r   r~  neg_x	exp_neg_xr  rZ  onedenoms	            rJ   _logistic_lowering_ruler  ?  s    
9
::
**Q-%hhuo)]]1(	77( 

cX
.C


8[%5
6C
**S)
$%	C	rQ   c                F    |t        d      t        j                  |      S rz  )r   r-   sinr}  s      rJ   _sin_lowering_ruler  Q  r  rQ   c                F    |t        d      t        j                  |      S rz  )r   r-   cosr}  s      rJ   _cos_lowering_ruler  X  r  rQ   c                F    |t        d      t        j                  |      S rz  )r   r-   tanr}  s      rJ   _tan_lowering_ruler  _  r  rQ   c                F    |t        d      t        j                  |      S rz  )r   r-   tanhr}  s      rJ   _tanh_lowering_ruler  f  r  rQ   c                F    |t        d      t        j                  |      S rz  )r   r-   r  r}  s      rJ   _log_lowering_ruler  m  r  rQ   c                F    |t        d      t        j                  |      S rz  )r   r-   log1pr}  s      rJ   _log1p_lowering_ruler  t  r  rQ   c                   |dk(  rt        j                  |      S |dk(  rt        j                  |      S t        d|       )Nr   r  zUnsupported rounding method: )r-   round	roundevenr   )r1  r   rounding_methods      rJ   _round_lowering_ruler  {  sC    ::a=!>>!
 =o=NO
PPrQ   c                ,    t        j                  |      S rC   )r-   ceilrs  s     rJ   _ceil_lowering_ruler        	1rQ   c                ,    t        j                  |      S rC   )r-   floorrs  s     rJ   _floor_lowering_ruler    s    	ArQ   c                ,    t        j                  |      S rC   )r-   ctlzrs  s     rJ   _clz_lowering_ruler    r  rQ   c                x    | j                   d   }|j                  st        d      t        j                  |      S )Nr   z,Population count is not supported on scalars)r   rU   rb   r-   ctpoprk  s      rJ   _population_count_lowering_ruler    s1    ]]1(	
C
DD	ArQ   c                   | t         j                  k(  r)t        j                  t        j                  ||            S | t         j
                  k(  rt        j                  ||      S | t         j                  k(  r)t        j                  t        j                  |      |      S | t         j                  k(  r)t        j                  t        j                  |      |      S | t         j                  k(  r)t        j                  |t        j                  |            S | t         j                  k(  r)t        j                  |t        j                  |            S t        d|        )a  A helper function for lowering comparison operations for boolean inputs.

  Args:
    primitive: A JAX primitive representing a comparison operation, which is
      one of the following: `lax.eq_p` (equals), `lax.ne_p` (not equals),
      `lax.lt_p` (less than), `lax.le_p` (less than or equal to),
      `lax.gt_p` (greater than), or `lax.ge_p` (greater than or equal to).
    x: A boolean array representing the first operand in the comparison.
    y: A boolean array representing the second operand in the comparison.

  Returns:
    A boolean array that is the result of applying the comparison operation
    between `x` and `y` based on the given primitive.

  Raises:
    ValueError: If an unsupported comparison primitive is provided.
  z"Unsupported comparison primitive: )r   eq_prE   logical_notlogical_xorr   lt_plogical_andle_p
logical_orgt_pge_prb   )r^  r   r  s      rJ   _cmp_boolean_lowering_helperr    s    $ #((??3??1a011CHH??1a  CHH??3??1-q11CHH>>#//!,a00CHH??1cooa011CHH>>!S__Q/00
9)E
FFrQ   c                @   t        |||j                  d   |j                  d   |j                  d         \  }}|j                  \  }}|j                  |j                  k7  r%t	        d|j                   d|j                         |j                  }t        j                  |t
        j                        r- t        t        j                  t        |       d      |||      S t        j                  |t
        j                        rt        j                  |t
        j                        }|rt        nt        |    }t         j"                  j%                  t         j&                  j)                  d      |      }	t+        j,                  |	||      S t        j                  |t
        j.                        r]t0        |    }t         j"                  j%                  t         j&                  j)                  d      |      }	t+        j2                  |	||      S t5        d|       )	Nr   r  zMixed dtype operands in cmp: r  FrR  r  zUnsupported dtype in cmp: )r'  r   r   rI   rb   rE   rF   r   r:  r  r  r  r   rU  _cmpui_lowering_typesr  r)   r   r   r   r   r*   r#  r   _cmpf_lowering_typescmpfr   )
r^  r1  r   r  rW  r   rI   is_uintr&  r'  s
             rJ   _cmp_lowering_ruler    s   	1cll1os||Aa8H	I$!Q<<.&&\\V\\!

'~R~F  ,,%^^E399%96	B 
1a 
 	^^E3;;'nnUC$7$78G!(.CD ""2>>#>#>r#BDII::iA&&^^E3<<(	*D""2>>#>#>r#BDII::iA&&8@AArQ   c                z    t        ||g| j                  | j                   \  }}t        j                  ||      S rC   )r'  r   r   r*   andirw  s      rJ   _and_lowering_ruler  
  7     
1	4s||	4cmm	4$!Q	Aq	rQ   c                    | j                   \  }t        | j                  j                  |      }t	        | t        j                  ||            S rC   )r   r   r   r   _not_lowering_ruler?   weird)r1  r   r$  rZ  s       rJ   _is_finite_lowering_ruler    sB    mm)(	77( 
C8Q!7	88rQ   c                z    t        ||g| j                  | j                   \  }}t        j                  ||      S rC   )r'  r   r   r*   orirw  s      rJ   _or_lowering_ruler    s6     
1	4s||	4cmm	4$!Q	1arQ   c                   | j                   d   }t        |j                        }|j                  st	        d|      }ntt        | j                  j                  |      }t        j                  j                  |d      }t        j                  |t        j                  j                  ||            }t        j                  ||      S )Nr   r  )r   r   rI   rU   r   r   r   r   r)   r   r   r*   r   r!  r"  xori)r1  r   r$  out_scalar_type	minus_onerZ  scalar_minus_ones          rJ   r  r  #  s     ]]1(%hnn5/	B0I 998H ~~))/2>"&&00;KLI 
Ay	!!rQ   c                   t        |      dkD  rt        d      | j                  d d \  }}|j                  t	        j                  t        j
                        k7  rOt        | j                  |g|j                  t        j
                        gd g      } t        d d      ||      }|s|S |\  }t        j                  |||      S )	Nr  z+select_n only supported with <= 2 argumentsr  rO   )r   r   rv   c                    | dk7  S r  rm   r  s    rJ   r  z)_select_n_lowering_rule.<locals>.<lambda>H  s
    qAv rQ   FrR  )r`   r   r   rI   r   r   r   r   r   r:  r*   select)r1  r&  r   r  	pred_avalrW  	lower_ctxr  s           rJ   _select_n_lowering_ruler  <  s    Y]
K
LLll2A&)V__**###"((#34V	I ?9%>y$OD	H"!	dAq	!!rQ   c                Z    t        j                  ||       }t        j                  ||      S rC   )rE   maximumminimum)ri  operandru  ress       rJ   _clampr  P  s#    GS!#	S#	rQ   c                6     t        t        d      | |||      S )z0Compute minimum_p(maximum_p(min, operand), max).FrR  )r:  r  )r1  ri  r  ru  s       rJ   _clamp_lowering_ruler  U  s     
36E	23Wc	JJrQ   c          	     F     fd}t        |t        j                        s`t        |t        j                        sF||k(  rAt        |||z         D ]-  }	 |t	        |	t        t        j                              |      }/ |S |dk7  rt        d|d|d      t        |t        j                        }
t        j                  |
t        |t        j                              }t	        dt        t        j                              }t        j                  |
|||      }t        j                   |j"                        5  |j$                  }|j&                  } |||      }t        j(                  |       d d d        |j*                  S # 1 sw Y   |j*                  S xY w)Nc                >   r9j                   j                  j                        }t        |g| | }|S ~ j                   j                  g j                  d t	               j                  t	              dz   d        }t        |g| }|S )Nr+  r  )r   r   rv   r  r`   )r   r  r   r6  r1  has_loop_indexr  s      rJ   	_run_bodyz+_lower_jaxpr_to_for_loop.<locals>._run_body`  s    --55'' 6 )+UFVFQFFd K --55#f+.Fa 12 6  +UCVCdCdKrQ   r  r  zOnly unroll=num_steps=z$ and unroll=1 supported. Got unroll=r   )r   r)   rO  rF  r   r   rE   r@   r   rb  r   rO  r*   r*  r/   ForOpInsertionPointr  induction_variableinner_iter_argsyield_results)r1  r  r  	num_stepsr6  r  unrollr  r  r   lbdubdstepfor_opiv
inner_args	inner_outs   ``  ``           rJ   _lower_jaxpr_to_for_loopr  [  s]   
" UBHH
%BHH-
v
 5%)+, 
a#4SYY#?
@$d Kq[

!yl"GyJL L5+"A"AB#

3*9k6U6UVW#	Q"3CII">	?$99S#tT*&	% 		"	"B''J"j)IJJy	
 

 
s   7FF c                  ~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|	||d}|r't        |t        t        j                              g|}|S )Nr  r   r  r  r  )r`   r   r  r6  ra  pattern_match_scan_to_fori_loopr>   r   rL  rb  r  r   r   rE   r@   )r1  r  linearlengthr  r  
num_consts	num_carry_split_transposer  num_extensivejaxpr_constsr  r6  consts_avals
args_avalsloop_index_startr  s                     rJ   _scan_lowering_ruler    s,    d)j(94---''mWU\\%,,&FFZ% D:,/,&$'zlC,
"t'ABJ!6<8&	z	2$ 	5"F		$2	# v):399)EF
M
MC	*rQ   c          	        t        |||g      \  }}}	|	d d |	dd  c\  }
}}t        | j                  g | j                  d |dz    | j                  |dz   d        ||
t	        j
                  ||
      |g|ddd}||g|S )Nr  r  r+  Tr  )r>   r  r   rv   r*   rW  )r1  
fori_jaxprcond_nconsts
cond_jaxprbody_nconsts
body_jaxprr  r5  body_constscarrylbubfor_outs                rJ   _lower_while_via_forir     s     %TL,+GH![%!9eABi.(2rD$	kk 2,"23q 0 23   jjR  ' b	7	rQ   c          	     2   t        j                  ||||      \  }}|t        | g||||||dS t        |||g      \  }}	}
t        | j                  ||g      \  }}}|
D cg c]  }|j
                   }}t        j                  ||
      } |j                  j                  j                  | }t        j                  j                  |      5  g ||j                  }t        | j                   j#                  g ||      |j$                  g| \  }t        j&                  ||j                         d d d         |j(                  j                  j                  | }t        j                  j                  |      5  g |	|j                  }t        | j                   j#                  g ||      |j$                  g| }|rt        j*                  |       d d d        t-        |j.                        S c c}w # 1 sw Y   xY w# 1 sw Y   /xY w)N)r  r  r  r  r  r+  )ra   pattern_match_while_to_fori_loopr   r>   rv   r   r/   WhileOpbeforeblocksrN  r)   r  at_block_begin	argumentsr  r   r   r  	conditionafterr  r  r  )r1  r  r  r  r  r  r  r5  cond_constsr  r  cond_const_block_shapesbody_const_block_shapescarry_block_shapesr^  carry_typeswhile_opbefore_block	cond_argscondafter_block	body_argsloop_outs                          rJ   _while_lowering_ruler6    sH    ??,
L-*a  	 !!  %/
\<(%!+{E !!L,#?@ G24F "''A'+'[[e,(.''..<,	''5 	07+7 6 67I$$H2H5GH 	% 	
 		
 
FT MM$../	0 -%%,,k:+	''4 
6+6 5 56I$$H2H5GH 	% 	
 		
 
H 	jj
 
h	7 (	0 	0
 
s    #G<A(HAHH
Hc                  |^}}t        |      }|At        | j                  j                  | j                  dd        ||   j
                  g| S t        j                  t        | j                  j                        }t        || j                        }t        j                  t        j                  j                  |t!        d|j"                              }t%        j&                  ||d      }	| j                  j                  | j                  dd        }
t)        j*                  |	j,                        5  t/        |      dkD  r<t1        | t        j2                  |t!        d|j"                              g|d|dd  i}nt        |
|d   j
                  g| }t%        j4                  |       d d d        t)        j*                  |	j6                        5  t        |
|d   j
                  g| }t%        j4                  |       d d d        |	j8                  S # 1 sw Y   lxY w# 1 sw Y   |	j8                  S xY w)Nr  r+  r   T)hasElser  branches)r  r  r   r   rv   r  r  r  r   r   rL  r   r*   r#  CmpIPredicatener   r   r/   IfOpr)   r  
then_blockr`   _cond_lowering_rulerW  r  
else_blockr  )r1  r9  r  r2  indexconstant_indexaval_to_ir_type_with_fn	out_typesr&  if_opr   r  s               rJ   r>  r>    s   ,%$/6.$$#2B2B122F$GR`IaIgIgjn  &--s++HH )3==9)	e[EJJ%?
$ ((4D
1%))11##AB' 2  
))*  8}q

**UK5::6
7  AB<	c *HQK,=,=EEcJJsO 
))* 
((1+*;*;
Cd
CCJJsO 
!  
s   :A:H0H+H(+H?c               |    | j                   j                  | j                        }t        ||j                  g| S Nr+  )r   r   rv   r  r  )r1  r  r  r5  r   s        rJ   _pjit_lowering_rulerG  ;  s9    ))11s?O?O1P	'	<t	<<rQ   c                   |S rC   rm   )r1  r   dst_shardingconcrete_meshs       rJ   _reshard_lowering_rulerK  A  s	     
(rQ   c                   ~|rt         |rt         |j                  rt         | j                  j                  | j                        }t        ||j                  g| S rF  )r   r6  r   r   rv   r  r  )r1  
call_jaxprjvp_jaxpr_funr  symbolic_zerosr  r   s          rJ   _custom_jvp_call_lowering_rulerP  G  sa     ..**11))11s?O?O1P	')9)9	AD	AArQ   c                   |rt         | j                  j                  | j                        }t	        ||j
                  g| S rF  )r   r   r   rv   r  r  )	r1  rM  fwd_jaxpr_thunk	out_treesrO  bwdr  r  r   s	            rJ   _custom_vjp_call_lowering_rulerU  X  sF     **))11s?O?O1P	')9)9	AD	AArQ   c                    ~ ~~g S rC   rm   )r1  r  r   s      rJ   _debug_callback_lowering_rulerW  h  s    	4	)rQ   c                   | j                   j                  t        d| d      t        | j                   j                        }|t	        |      vrt        d| d|       | j                   j                  |   S )Nzprogram id: z- was passed, but user did not provide a grid.%user passed in program id with axis: , but grid only has length: )r   ru   rb   r`   rF  )r1  rL  r  s      rJ   _program_id_lowering_ruler[  o  s     	++3

tfIJ  s##556&	v

/v 68	  
			/	/	55rQ   c               4   t        | j                  j                        }d}t        | j                  j                        D ]  }|t        ||v      z  }||dz   k(  s n& t        d| d| j                  j                         t        j                  |      S )Nr   r  rY  rZ  )	rq  r   rt   rF  r   r   rb   r?   iteration_bound)r1  rL  vmapped_axesseen_user_axesr   s        rJ   _num_programs_lowering_ruler`    s     S))667,.%%//0 ac!</00N!
 
/v 6((223	5  
		Q	rQ   c                   | j                   \  }t        j                  t        | j                  j
                  |      |||      S rC   )r   r?   repeatr   r   r   )r1  r   repeatsrL  r$  s        rJ   _repeat_lowering_rulerd    sB    +8	



;
;X 

 rQ   c                   | j                   \  }t        j                  t        | j                  j
                  |      |||||      S )N)r  stride_dimension)r   r?   dynamic_rotater   r   r   )r1  r   shiftrL  r  stride_axisr$  s          rJ   _roll_lowering_rulerj    sL     +8			



;
;X 
"	
 	rQ   c                   | j                   \  }t        | j                  j                  |      }|dgt	        |      z  }t        j                  |      t        j                  |      z
  }t        j                  |||||      S )z!Lowers a slice to vector dialect.r  )	r   r   r   r   r`   r   r   r0   r  )r1  r   limit_indicesstart_indicesr  r  rZ  r  s           rJ   _slice_lowering_rulern    s|    
 +8	77( _cC&&G
((=
!BHH]$;
;%		%	%=%
 rQ   c                z    t        ||g| j                  | j                   \  }}t        j                  ||      S rC   )r'  r   r   r*   r  rw  s      rJ   _xor_lowering_rulerp    r  rQ   c                z    t        ||g| j                  | j                   \  }}t        j                  ||      S rC   )r'  r   r   r*   shlir1  r   r   s      rJ   _shift_left_lowering_rulert    s7     
1	4s||	4cmm	4$!Q	Aq	rQ   c                z    t        ||g| j                  | j                   \  }}t        j                  ||      S rC   )r'  r   r   r*   shrsirs  s      rJ   %_shift_right_arithmetic_lowering_rulerw    7     
1	4s||	4cmm	4$!Q	Q	rQ   c                z    t        ||g| j                  | j                   \  }}t        j                  ||      S rC   )r'  r   r   r*   shruirs  s      rJ   "_shift_right_logical_lowering_ruler{    rx  rQ   c                F     t        t        j                  d      | |      S rq  )r:  ra  erf_inv_lowering_helperrs  s     rJ   _erf_inv_lowering_ruler~    s(    
**U

 rQ   c                   t        |j                  j                  t        j                        st        d      t        j                  ||      S )NzOnly float32 is supported.)approx)r   r   r  r)   r  rb   r?   
reciprocal)r1  r   r  s      rJ   _reciprocal_lowering_ruler    s7    	AFF''	4
1
22	&	))rQ   c                  t        |j                  j                  t        j                        st        d      |t        j                  t        j                  t        j                  t        j                  fvrt        d      | j                  \  }}t        j                  j                  |j                  t        j                   t        j"                  |                  }t%        j&                  |||      S )Nz Only float32 input is supported.zaOnly bfloat16, float8_e5m2, float8_e4m3fn, and float8_e4m3b11fnuz are supported as target dtypes.)r   r   r  r)   r  rb   rE   bfloat16float8_e5m2float8_e4m3fnfloat8_e4m3b11fnuzr   r   r   rU   r$   r   rI   r?   stochastic_convert)r1  r   random_bitstarget_dtyper5  r  rZ  s          rJ   _stochastic_round_lowering_ruler    s     
AFF''	4
7
88	ll	oo			  	*  ,,-1g]]mmT**399\+BC( 
		![	99rQ   c                    | t         j                  k(  r|t         j                  k(  ry | t         j                  k(  r2|t         j                  t         j
                  t         j                  fv ry t        d|  d| d      )Nz!Unsupported elementwise packing: r  z8. Only f32 <-> bf16 and i32 <-> i16/i8/i4 are supported.)rE   r   r  r@   int16int8int4rb   )unpacked_dtypepacked_dtypes     rJ   !_check_elementwise_packing_dtypesr    sr    s{{"|s||'C
syy \	ii3886 & ).)9l^ L? ?	 rQ   c                  | j                   d   }t        |j                  |       t        |      }t        j
                  j                  |j                  t        t        j                              }t        j                  |||      S )Nr   )target_type)r   r  rI   r   r)   r   r   rU   rE   r  r?   pack_elementwise)r1  r  r  r  packed_ir_typerZ  s         rJ   _pack_elementwise_lowering_ruler    sg     LLO'#GMM<@$\2.]]mm&szz2( 
		h	GGrQ   c                    | j                   d   }t        ||       t        j                  j	                  |j
                  t        |            }t        j                  ||t        |      |      S )Nr   )source_typer@  )	r   r  r)   r   r   rU   r   r?   unpack_elementwise)r1  r   r@  r  r  r  rZ  s          rJ   !_unpack_elementwise_lowering_ruler  &  sd     LLO'#NLA]]mm&~6( 
		0>e
M MrQ   c                   ~| j                   \  }t        j                  t        | j                  j
                  |      |      S rC   )r   r?   bitcastr   r   r   )r1  r   tyr$  s       rJ   _bitcast_lowering_ruler  3  s?    +8	



;
;X 	
 rQ   c               *   | j                   \  }| j                  \  }t        j                  |j                        }t        j                  |      }||k7  rt        d      t        j                  t        | j                  j                  |      |      S )Nz!Changing bitwidths not supported.)r   r   r   r  rI   r   r?   r  r   r   r   )r1  r   r  r  r$  r  r  s          rJ   #_bitcast_convert_type_lowering_ruler  ?  s     +7+8%%gmm4,%%i0,\!
A
BB	



;
;X 	
 rQ   c                  t        | t        j                        rt        j                  | j
                  t        j                        rd| j                  t        j                  k(  sJ t        |j                  j                  | t        j                        }t        j                  |      S t        |j                  j                  | d| j                        }t        |t         j"                        sJ t%        j&                  |g g       S t        | t(        j*                        rEt        |j                  j                  | t        j                        }t        j                  |      S t-        dt/        |        d      )Nr   T)r   r   r   r   )r   r   r   rE   rF   rI   r   r   r   r   r   r   r   r   r?   	sem_allocr)   r   r.   allocar   r   r   r   )rV   r1  memref_types      rJ   _alloc_valuer  S  s,    e''(
~~djj+"="=>.":"::::#



;
;
%//k
 ]];''#



;
;
!((	k R]]333]];B//$223!99#--K
 ==%%.tDzl!<==rQ   )alloc_fnc               F    |rt        d       j                  D cg c]"  }t         j                  j                  |      $ }}t        j                  |      }|j                  D cg c]  }|j                   }	} j                  j                         5  t        j                  |      }d d d        t        j                  |j                        5  t         fd|	      }
t!        d |	D              }t!        t        t"        |	|            } j                  j%                  g  j&                  |       t)         |g||
 }t        j*                  |       d d d        |j,                  S c c}w c c}w # 1 sw Y   xY w# 1 sw Y   |j,                  S xY w)Nz4run_scoped lowering does not support collective axesc                     |       S )N)r1  rm   )rV   r  r1  s    rJ   r  z+_run_scoped_lowering_rule.<locals>.<lambda>  s    HTs3 rQ   c              3  l   K   | ],  }t        |t        j                        r|j                  nd  . y wrC   )r   r   r   rU   r  s     rJ   rK   z,_run_scoped_lowering_rule.<locals>.<genexpr>  s0      , %/q%2C2C$D$N ,r>  r+  )r   r   r   r   r   r?   RegionOprH  rV   r   r.  r0  r)   r  r  rL  r   rX   r   rv   r  r  r  )r1  r  collective_axesr  r6  rV   rZ  regionr  rU  r  rv   r  s   `  `         rJ   _run_scoped_lowering_ruler  r  s~    
T
UU --
 c**GGN(  <<!&#ll+aff+(+
--/ .&&u-E.	% 
3X>D ,"*, ,L;%|5 6L



&
&7s''7,7 ' C U
3V
3d
3CJJsO
 
'
 ,. .
 
s$   'E6*E;F BF F	F c                $    t        j                   j                  j                  || fd      \  }}d } j                  j                  x}r0t        |      dkD  rt        d      |d   }|j                  |d       }|rt        d|       ||fS )Nc                    t        |       S )N)	axis_name)_axis_index_rule)rk  r1  s    rJ   r  z'_device_id_to_logical.<locals>.<lambda>  s    #C48 rQ   r  z@Unable to determine core axis name if grid_names is more than 1.r   z Unrecognized axes in device_id: )	r2   device_id_to_logicalr   ry   rs   r`   r   poprb   )r1  	device_iddevice_id_typelogical_device_idnon_mesh_axes
core_indexrs   core_axis_names   `       rJ   _device_id_to_logicalr    s     &0%D%D	''8	&"] *''222Z2
:
L   ]N"">48J

*=/:  
J	&&rQ   c          	        t        j                  || j                        \  }}t        j                  ||dt
        j                  t        j                  t        j                  t        j                  h       t        j                  ||      \  }}t        ||j                  |j                  |      \  }}t        j                  |      S )Nread)allowed_semaphore_types)r   r  r   r2   check_sem_avalsr   r   r   r   r   SEMAPHORE_INTERPRET_DTYPEr  rI   rU   r?   sem_read)r1  r  r  sem_avalsem_transforms_avalssemr  r5  s           rJ   _semaphore_read_lowering_ruler    s     $-#;#;Is||#T ( 

 
 




'
'

/
/		
 ,,Y=/#z#x~~x~~zJ&#q	c	rQ   c               L   t        j                  || j                        \  }}}}}t        j                  ||      \  }}}}	}
t        ||j                  |j
                  |      \  }}|	!t        | |	|      \  }	}||
t        d      |}
t        j                  |||	|
       g S )NzBCannot specify both `core_index` and the core axis in `device_id`.r  core_id)
r   r  r   r  rI   rU   r  rb   r?   
sem_signal)r1  r  r  r  r  r5  r  r  r  r  r  r  s               rJ   _semaphore_signal_lowering_ruler    s     #11)S\\J(Aq!Q2;2J2J3/#z5)Z #x~~x~~zJ&#q.sI~NIw		P
 	
 j..ey*E	)rQ   c                  t        j                  || j                        \  }}}}t        j                  ||      \  }}}}|st        d      t	        ||j
                  |j                  |      \  }}t        j                  ||       g S )Nz'Non-decrementing wait is not supported.)	r   r  r   r   r  rI   rU   r?   sem_wait)	r1  r  r  r  r5  r  r  r  	decrements	            rJ   _semaphore_wait_lowering_ruler    s|      ..y#,,G(Aq!&/&>&>y$&O##z5)	
G
HH#x~~x~~zJ&#q,,sE	)rQ   c          	        |rt        d      t        j                  ||      \	  }}}}	}
}}}}t        j                  || j                        \	  }}}}}}}}}|j                  t
        j                  k(  rt        d      t        j                  || j                        }|d   |d   }}t        ||j                  ||      \  }}|%t        ||j                  |j                  |      \  }}t        ||j                  ||	      \  }}t        |
|j                  |j                  |      \  }
}d }|t        | ||      \  }}t        j                  |||
||||       g S )Nz#DMA with add=True is not supported.z(DMAs with bool dtypes are not supported.r   r  )source_semaphorer  r  priority)r   r   r  r   rI   rE   r   rv   r  rU   r  r?   enqueue_dma)r1  r  r  r  r  r  src_refsrc_transformsdst_refdst_transformsr  sem_transformssrc_semsrc_sem_transformsr  src_ref_avalr5  dst_ref_avalr  src_sem_avalrv   src_ref_block_shapedst_ref_block_shaper  s                           rJ   _dma_start_lowering_ruler    s    	
C
DD tT*
	 tS\\2 F<L!Xq,1 399$
H
II))$0@0@A,-9!_l1o*|!!#6*'1 ##\%7%79KJGQ |!!#6*'1 #x~~x~~~N&#q'.sI~NIw//	 
)rQ   c          	     ~   t        j                  ||      \	  }}}}}}	}
}
}t        j                  || j                        \	  }}
}}
}}
}
}
}
t        j                  || j                        }|d   }t	        ||j
                  |j                  |      \  }}
t	        ||j
                  ||      \  }}
t	        ||j
                  |j                  |	      \  }}
d }|t        | ||      \  }}| j                  s| j                  ddd      rt        j                  ||||       g S t        j                  |||||       g S )Nr  r  r     )r  r  )r   r  r   rv   r  rI   rU   r  r~   r!   r?   	wait_dma2)r1  r  r  r  srcr  dstr  r  r  r5  r  src_avaldst_avalr  rv   r  r  s                     rJ   _dma_wait_lowering_ruler  !  s?    tT*
			5>5M5M
CLL628Q!Xq!Q ))$0@0@A, O/#x~~x~~~N&#q#x~~
K&#q#x~~x~~~N&#q'.sI~NIws::4BGMM#sC1 
) MM#sC9gF	)rQ   c               
   | j                   j                  }|r ||v rt        | |j                  |            S t	        j
                         }| j                   j                  }|t        d      |j                  }|j                  }|j                  |      }t        ||         }t        t        j                  ||dz   d  t        j                              }	t        j                  t        j                   ||	      |      S )Nrg  zMesh context is not set.r  rO   )r   rs   r[  r@  r?   r  ry   rb   
mesh_shaper`  r   r   r  r@   r*   rd  r_  )
r1  r  rs   r  ry   r  r`  
axis_index	axis_sizeminor_divisors
             rJ   r  r  C  s    ##..*I+$Sz/?/?	/JKKmmo)%%22,
/
00&&*&&*	***Z01)ggja)*"((;- 
U[[M:I	FFrQ   c                    t        | j                  j                  | j                  d         }t	        j
                  |      S r  )r   r   r   r   r?   sem_barrier)r1  r  s     rJ   _get_barrier_semaphore_ruler  X  s9      	77q9I+ 
	%%rQ   c                0    t        j                  |       g S rC   )r?   delay)r1  nanoss     rJ   _delay_ruler  b  s    ))E	)rQ   c                  ~~|rt        d      t        j                  ||	|      \  }
}|rt        d      | j                  D cg c]  }|j
                    }}t        |      }t        |      dk(  xr |d    }|s|st        d| j                         |r|rut        j                  |g|
  t        d |
D              st        d      dj                  d	 t        t        j                         j                  |            D              }t!        j"                  |
||
       yt        | j                        dk7  rt        d      | j                  \  }|
\  }|r|j%                  d      st        d      |d d }t!        j&                  d      }t)        j*                  |j,                        5  t/        |j0                        }t(        j2                  j5                  |j
                  |t(        j6                  j                  d            }t9        j:                  |g g       }t(        j<                  j5                         }t?        j@                  |d      }|gt        |j
                        z  }tC        jD                  |||       t!        jF                  ||j
                  |       t!        jH                  g        d d d        yc c}w # 1 sw Y   yxY w)Nz/Ordered debug_print is not supported on Pallas.zAOnly positional arguments are supported by debug_print on Pallas.r  r   zJAll inputs to debug_print must be all scalars or a single vector, but got c              3     K   | ]C  }t        |j                  t        j                        xr |j                  j                  d k(   E yw)r  N)r   r   r)   r   r   )rG   rk  s     rJ   rK   z$_debug_print_rule.<locals>.<genexpr>  s<       SXXr~~
.
G388>>R3G
G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.r  c              3  F   K   | ]  \  }\  }}}}|	| d| | n|  y w)N$rm   )rG   rA  textfieldspecr5  s         rJ   rK   z$_debug_print_rule.<locals>.<genexpr>  s>      )c)D%q $)#4TF!D6#
$
>r  )	formattedrm   z2Only one vector input to debug_print is supported.z{}z5For vector input, the format string must end with {}.r  r  r   )%r   r   merge_callback_argsrb   r   rU   r  r`   r2   check_debug_print_format	TypeErrorr  r   string	Formatterr   r?   r  endswithr  r)   r  r  r   rI   r   r   r   r.   r  r  r*   r   r0   r5  
log_bufferr  )r1  fmtorderedpartitionedin_treestatic_argsnp_printoptionshas_placeholderslogging_recorddyn_argsr  r   rV   is_scalar_inputsis_all_scalarsis_single_vectorrk  r  r  r  r  
index_typezerorB  s                           rJ   _debug_print_ruler  h  s    ?
O
PP..w+N,$
K  25>$**n>>'(.)*a/K8H8K4K
,
	~	  ))#55   ;
 	
 GG -6 &&s+. c GGD#!12 	!
<  LL'4&3	d!3
L
MMCR#<<&	% $TZZ0L}}  

\\''(AB ! H
 --"b
)C!!#J>>*a(Dfs4::&G
LLc7#NN3

C(JJrN 
G ?h 
s   K6DKKc                <   ~ t        |      dk(  r7t        |d   t              r$t        j                  |d   j
                         g S t        d |D              }|s'|D cg c]  }|j                   }}t        d|       t        j                  |       g S c c}w )Nr  r   c              3  d   K   | ](  }t        |j                  t        j                         * y wrC   )r   r   r)   r   )rG   seeds     rJ   rK   z+_prng_seed_lowering_rule.<locals>.<genexpr>  s     MtZ		2>>:Ms   .0z+All seed data must be scalar integers. Got )	r`   r   rP  r?   prng_set_seed_32r  r  r   rb   )r1  seedsall_integersr  
seed_typess        rJ   _prng_seed_lowering_ruler    s    	 	Z1_E!Ho>q))*I MuMM,	(-.$)).J.
B:,O
PPu	) /s   Bc                   t        |      dk  rt        d      | j                  d   }t        | j                  j
                  |      }t        j                  |      S )Nr  z*random_bits only supports rank>=2 outputs.r   )r`   r   r   r   r   r   r?   prng_random_bits)r1  rU   r$  rZ  s       rJ   _prng_random_bits_lowering_ruler    sU    Z1_
J
KK]]1(	77( 
		h	''rQ   c               B    t        |j                  d      } || |      S rq  )r:  r  )r1  r  implseed_lowerings       rJ   random_seed_loweringr     s    DII>-	sE	""rQ   c               2   |dk(  sJ d       | j                   \  }t        |j                  t        j                        sJ |j                  j
                  j                  }t        j                        sfd}|}t        |d      } || |||      S )Nr  zOnly 32-bit PRNG supported.c                    t         j                  j                  |       j                  t        j
                        } j                  | ||      S rC   )r  r5   key_datar  rE   r  r  )key	bit_widthrU   r  s      rJ   new_loweringz*random_bits_lowering.<locals>.new_lowering  s;    JJ$++CJJ7cc9e44rQ   FrR  )r%  rU   )
r   r   rI   r   r  r  r  r  r  r:  )	r1  r  r%  rU   rV   	_proxy_fnr&  bits_loweringr  s	           @rJ   random_bits_loweringr)    s    	b777
,,%$	DJJ

	++	+			$)		!	!$	'5 II>-	sDIU	CCrQ   c                   | j                   \  }}t        |j                  t        j                        sJ |j                  j
                  }t        |j                  d      }t        j                  |      r
 || ||      S t        j                  | t        j                  |      |gt        t        j                  | j                              }  || ||      S )NFrR  )r   r   )r   r   rI   r   r  r  r:  fold_inr  r  r   r   rS   r   rL  r   )r1  r  msgs	keys_aval	msgs_avalr  fold_in_lowerings          rJ   random_fold_in_loweringr0    s    )Y	IOOTZZ	00	0			$t||eDd#Ct,,


c"*"8"8"CY!O"%h&<&<cmm"LNC Ct,,rQ   c                    | j                   d   }t        |j                  t        j                        sJ |j                  j
                  }t        j                  |      s|S t        d      )Nr   zIkey_data not support for Pallas PRNG keys. Use split_pallas_seed instead.)	r   r   rI   r   r  r  r  r  rb   )r1  r$  r-  r  s       rJ   random_unwrap_loweringr2    s[    ll1o)	IOOTZZ	00	0			$		!	!$	'J$	 rQ   c               H    ~ t        j                  |      s|S t        d      )NzMwrap_key_data not support for Pallas PRNG keys. Use wrap_pallas_seed instead.)r  r  rb   )r1  r#  r  s      rJ   random_wrap_loweringr4    s*    			!	!$	'O#	 rQ   c                    |j                   S rC   )r  )r1  r#  s     rJ   _split_key_lowering_ruler6    s     
		rQ   c                   t        j                  |      st        d|       S t        |t	        |j
                              S )Nz$Can only join Pallas keys. Got impl=r  )r  r  rb   rP  r   r  )r1  r  r  s      rJ   _join_key_lowering_ruler8  &  s9    		!	!$	'<TFCDD	E$..4I	JJrQ   c                  ~ |st        d      t        j                         sg S t        j                  j                  ||      }|j                  j                         \  }|j                  j                         \  }|j                  j                         \  }t        j                  j                  ||      }t        |t        j                        sJ t        |t        j                        sJ t        dt        t        j                               }	t#        j$                  ||	      }
t'        j(                  |
|j*                         g S )NzmNon-debug checks are not supported by the Mosaic backend. Functionalize them via `jax.experimental.checkify`.r  )r   pallas_helpersdebug_checks_enabledr  r  r   _predrF  	_metadata_payloadr   r   FailedCheckErrorr   r   rE   r}   r*   r  r+   assert_
fmt_string)r1  err_treedebugerr_argserrorr&  exception_treepayload	exceptionr  not_preds              rJ   _check_lowering_rulerJ  -  s    
	
	?  
	,	,	.I
((

Xx
0%;;&4__++->nn##%)7hh  9)	Ix88	99	9	Ix88	99	9 "/9:)ZZi((**Xy++,	)rQ   c                :    d }t        |d      } || ||||      S )Nc                    t        j                  d      5  t        j                  | |||d      }d d d        |S # 1 sw Y   S xY w)Nthreefry2x32F)use_rolled_loops)r  named_scoper   _threefry2x32_lowering)k1k2m1m2r  s        rJ   
_lower_funz*_threefry2x32_lowering.<locals>._lower_funM  sC    		( P''BBOcPJPJs	   ;ATrR  ro  )r1  rQ  rR  rS  rT  rU  threefry_lowerings          rJ   rP  rP  K  s)    
  
TB	3BB	//rQ   c               4   t        j                  |      }|t        j                  t        j                        j
                  kD  r:t        dt        j                  t        j                        j
                   d      d }t        |d      } || |      S )NzIota with >z items.c                J   t        j                  | t         j                        }d}t        t	        |       dz
  dd      D ]8  }t        j                  t         j                  | |      }|||z  z  }|| |   z  }: t        j                  | t         j                        }||fS )NrO   r  r  )rI   rU   r  )rE   zerosr@   rF  r`   r   broadcasted_iota)rU   	iota_data
multiplierr  	counts_lo	counts_his         rJ   rU  z-_iota_2x32_shape_lowering.<locals>._lower_fun\  s    		%syy1IJSZ\2r* &&		#i 9z))iE#Jj 		%syy1IirQ   TrR  r4  )r   r  r  rE   r@   ru  r   r:  )r1  rU   total_elementsrU  iota_lowerings        rJ   _iota_2x32_shape_loweringra  V  sr    775>.bhhsyy)---
BHHSYY,?,C,C+DGL
MM
  J>-	s%	((rQ   c                  
 |\  
|d   }t        | j                  j                  | j                  d         }t	        |t
        j                        st        d      t        |      D ]y  \  }\  }}}|dk(  r|dk(  r|dk(  r|f
fd	}	|dk7  rt        j                   |	|      
g|      
|dk7  rt        j                  
 |	|      g|      
|dkD  spt        d       
S )Npadding_configr   z Only vector types are supported.c                   t        j                  j                        }| ||<   t        j                  j                  |j                  j                        }t        t        j                        rt        j                  |      }|S t        j                  j                  j                  j                        }t        j                  |t        j                  j                  ||            }|S rC   )r  r   rU   r)   r   r   r  r   OpResultr0   r  r   r*   r   r!  r"  )rL  rL  rU   pad_vec_typepadscalar_attrr  padding_values         rJ   _padz _pad_lowering_rule.<locals>._pad{  s    7<<%%&eeDk]]&&

,,
#
#l
 
M2;;	/|]; j ll&&w||'@'@-Pnn  **<E
 jrQ   r  z!Not implemented: interior padding)r   r   r   r   r   r)   r   r   r   r?   r  )r1  r  r   rc  rZ  rL  lowhighinteriorrj  r  ri  s             @@rJ   _pad_lowering_rulern  l  s    '=*+.	77a( 
Hbmm	,
@
AA%.~%> E!d!S$
axDAI(a- $ axcG 4Egqy$t* 5Fg!| CDD;E> 
.rQ   c               d    t        |      D ]  \  }}d|v s|t        |      c S  t        d      )Nmosaicz2No mosaic or default platform indexing rule found.)r   r   r   )r1  r  r   pss       rJ   _platform_index_loweringrr    sC     # ea2~^
 	:	 rQ   c                   | j                   j                  |f      d   }t        |t        t        j
                              S )Nr   r  )r   r   r   r   rE   r@   )r1  r  r  s      rJ   _dim_as_value_loweringrt    s6    $$AA3&I!L+	[,=cii,H	IIrQ   c                
    ~ ~g S rC   rm   rs  s     rJ   _touch_lowering_rulerv    s    	1	)rQ   r~  )r   AnyMemorySpace | Nonerg   zTPUMemorySpace | Literal[ANY])r   rw  rg   zir.Attribute)F)rI   r:   r   r}   rg   ir.Type)NNFT)r   rw  r   r}   r   r}   rC   )r   r   r   zir.Type | Nonerg   rh   )r  zjax_core.Primitiver   zCollection[tpu_core.KernelType]r   r}   rg   zCallable[[T], T])r   r   rV   r   rU   ztuple[int, ...] | Nonerg   rx  )r  z%str | tpu_core.GridDimensionSemanticsrg   rx  )r  r  r   mlir.LoweringRuleContextr  zjax_core.DebugInforg   None)r   ry  rQ  r{  r  r  rR  r|  r{   rz   r   r}  r  r}   rg   z	ir.Module)r  r  rV   jax_core.AbstractValuerk  rx  r  r  r{   rz   r~   r}   r   
Any | Noner    DynamicShapeReplacementFn | Nonerg   func.FuncOp)r  r  r  r  rk  rx  r{   rz   r~   r}   r   r|  r   r}  r  r}   rg   r~  )r*  r   r7  r}   rg   r   )r@  	list[str]rA  r  rg   ztuple[list[str], list[str]])r1  ro   r  r  r  rh   rg   zlist[ir.Value])rL  objectrV   r   rg   r   r9  )rA  z+tuple[indexing.Slice | int | ir.Value, ...]r  r}   rg   z*tuple[ir.Value, int | ir.Value, int, bool])r  r  r  &tuple[int | pallas_core.Squeezed, ...]r  r}   rg   ztuple[tuple[ir.Value, ...], tuple[int | ir.Value, ...], tuple[int, ...], tuple[bool, ...], tuple[int | pallas_core.Squeezed, ...]])r  Sequence[int]r  r  rg   zSequence[bool])
r  rh   r  r  r  r:   r  r  rg   z7tuple[ir.Value, tuple[int | pallas_core.Squeezed, ...]])
r  rh   r  r7   r  r:   r  r  rg   zBtuple[ir.Value, DTypeLike, tuple[int | pallas_core.Squeezed, ...]])
r  rh   r  r8   r  r:   r  r  rg   z tuple[ir.Value, tuple[int, ...]])r1  r   rg   rP  )r1  r   rL  rh   rg   ztuple[ir.Value, jnp.dtype])r1  r   rL  rh   rg   rh   )r1  r   rU   r  )r   r9   r  r   rg   r9   )r   ir.Value | objectr  r  rW  r   r   r   r$  r   rg   ztuple[ir.Value, ir.Value])r5  r   )r   r9   r  r9   )r1  r   r  r  r  int | ir.Valuer  r  r  r}   r  r   )r1  r   r  jax_core.ClosedJaxprr
  ztuple[bool, ...]r  r   r  r}   r  z
bool | intr  r   r  r   r  r}   )
r1  r   rM  r  rN  zlu.WrappedFunr  r   rO  r}   )r1  r   rL  r   )rV   r{  r1  r   rg   rh   )r1  r   r  primitives.DeviceIdType)r1  r   r  r  r  r   r  r}   )r1  r   r  r   )r1  r   r  r   )r1  r   r  rx  )r1  r   r#  rP  )r1  ry  r  r'   )r1  r   r   z	jax.Array(  r  
__future__r   collections.abcr   r   r   r   r   r   r  r  r  typingr   r	   r
   r   r   r   r  r   r   r   jax._srcr   r   r   r   rS   r   r   r   r   r,  r   r   mesh_libr   r   r   r   r   r    jax._src.cloud_tpu_initr!   jax._src.exportr"   jax._src.export._exportr#   jax._src.interpretersr$   r%   r.  jax._src.laxr&   r  jax._src.lax.control_flowr'   jax._src.libr(   jax._src.lib.mlirr)   jax._src.lib.mlir.dialectsr*   r+   r,   r-   r.   r/   r0   jax._src.pallasr   r1   r:  r2   r3   ra  jax._src.pallas.mosaicr   r4   tpu_primitivesr5   r  jax._src.stater6   state_primitivesjax._src.state.typesr7   r8   jax._src.typingr9   r:   jax._src.utilr;   r<   r=   r>    jax.experimental.mosaic.dialectsr?   	jax.numpynumpyrE   r   r  r   r   AnyMemorySpacer   r   r   rI   r   r  r  r@   ru  r_   ra   r  rL  
unsafe_mapr   r?  r   rN   rP   rX   rZ   r   DimSizer   r   	dataclassro   r   r   r   r   r   r   r   
KernelTyper  rq  r  r   TCr  r
  r  r  r  r  r  r  r  r:  rf  r  rE  r  rb  get_pr  swap_pr  r  r  r  r  r  r  r  r  r  rP  r  load_pr  r  r  r/  r  multiple_of_prH  r`  r   CombiningKindMAXIMUMFrT  MAXSIrU  MAXUIREDUCE_MAX_KINDSr   ri  REDUCE_MAX_IDENTITY_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_preduce_and_prq  reduce_or_prv  broadcast_to_prx  broadcast_in_dim_pr  r  dot_general_pr  r  convert_element_type_pr  	reshape_pr  	squeeze_pr  concatenate_pr  split_pr  r  r  gather_pr  transpose_pr  r'  add_p	add_any_pr-  r/  r  stop_gradient_pr;  max_pr@  min_prF  rN  argmax_prR  argmin_prU  sub_prY  mul_pr]  div_prb  rem_prg  abs_prl  neg_pro  sign_prt  nextafter_prx  rsqrt_pr  sqrt_pr  square_pr  exp_pr  pow_pr  integer_pow_pr  exp2_pr  
logistic_pr  sin_pr  cos_pr  tan_pr  tanh_pr  log_pr  log1p_pr  round_pr  ceil_pr  floor_pr  clz_pr  population_count_pr  r  r:  eqr   r;  r  sltr  sler  sgtr  sger  ultuleugtuger  CmpFPredicateOEQONEOLTOLEOGTOGEr  r  r  r  and_pr  is_finite_pr  or_pr  not_pr  
select_n_pr  r  clamp_pr  r  scan_pr  r   while_pr6  cond_pr>  jit_prG  	reshard_prK  custom_jvp_call_prP  custom_vjp_call_prU  debug_callback_prW  program_id_pr[  num_programs_pr`  repeat_prd  roll_prj  slice_prn  xor_prp  shift_left_prt  shift_right_arithmetic_prw  shift_right_logical_pr{  	erf_inv_pr~  reciprocal_pr  stochastic_round_pr  r  pack_elementwise_pr  unpack_elementwise_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  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_pr0  random_unwrap_pr2  random_wrap_pr4  split_key_pr6  
join_key_pr8  check_prJ  threefry2x32_prP  iota_2x32_shape_pra  pad_prn  platform_index_prr  dim_as_value_prt  touch_prv  )r{   s   0rJ   <module>r4     s   B " D D      > > 
       % '   &  %   %  #  ; & * & 4 % , 7 #   , ) + + - * - / 5 & 1 3 1 ? 6 # 9 : , ! " " $ 0  
 	%%((>9!!288G$  $ "((288$((


CZCZ (778 <2 22 %
8S !"E#s(O3 
 " " "P(  > > >,@'@"@:F+7F 27*.;B6 *.$!%4" (	4"
 4" 4"n%" 6>5H5HIk+r/I  CL 6>5H5H5K5K4M#	
 2 	
   %)
";


 "
 	
= E"y y #yxb
8b
.b
 #b
 
	b
X "&.3o.o)o o
 Fo %o o (,o ov FJ88
 8 	8
 +8 %8 8 8 #C8 8F FJ.3;; +; 	;
 %; ; ; #C; (,; ;|*	 PPP !P.f	f!/f8@ffR (..5IC	C JC  (//EJJ	J KJ,1E&	4&EI&/&0'';' 	'
'T*G	GG G <	G
 =GT%	%% % <	%
 H%P!	!! ! <	!
 &!H( d#! ! $! . 
))eD^; E^;@/Gd#2	#2-5#2#2L(	(2:(( 
))eDp	p Epf +AX-@-@+A/f LL&&&//v++11--33  LL%-xrxx)--  1GG24  ( s'' ()B C LL&&&//v++11--33  LL%,xrxx)--  1GG24  ( s'' ()B C LL&&&**v++//--11  LL#q  1GG24  ( s'' ()B C (()
+ *
+ (
+ )
+ (778	(5 9 )?8+>+>)?')	')')T>1B ))*z	z +zzMB -Cx/B/B-C=3	=3=3@ 4Jh6I6I4JK L. 4Jh6I6I4JK L( ))8N(:M:M8NO2 P2
 2HH4G4G2HI	 J2 

#4 $4$ %;2	;2 &;2| (3 )3        	 
 "   F II2h112u ))eD, E,9 4 ++, - II%6L8K8K6L	,	, II%6L8K8K6L	,	,	> ? @ ? @ II2h112u,, II2h112u,, II2h112u	,	, II2h112u	,	, 		0F(2E2E0FG, H, II2h112uII 

1G83F3F1GH I ( ) $ % 

# $ % & 		0F(2E2E0FG H 		e< = ))* +
 

u=
 >
 '  ( " 		" # 		" # 		" # 

# $ 		" # $ % $Q %Q 

# $ $ % 		" # ../ 0 HHe!!$$HHe!!$$HHe!!%%HHe!!%%HHe!!%%HHe!!%%  HHe!!$$HHe!!$$HHe!!%%HHe!!%%HHe!!%%HHe!!%%  HHe!!%%HHe!!%%HHe!!%%HHe!!%%HHe!!%%HHe!!%% $GDB> XXsxx388SXXsxxH DC,Bh.A.A,BCi*D1 II2h112u
 (9 )9 HH1X001e
 		0F(2E2E0FG" H"0 5Kx7J7J5KL" M"&
 $K %K
,$2,;I,(6, .2, &)	,^ JJ3x223*	*  * 	*
 * * * * * **Z	: 2HH4G4G2HI9 	9  J9 x 

1G83F3F1GH# I#L 

1G83F3F1GH= I=
 ' (
 *<<=B	B %B !	B
 B B >B  *<<=B	B >B 	223 4 *@H,?,?*@66 ,Bh.A.A,B   //0	 1	 --.	 /  2HH4G4G2HI	 J  II2h112u
 '8&&'


   '8&&'


 '8&&'


 & ' 
//0* 1* 99::	: ;:,
 99:	H		H ;	H ;;<	M		M =	M 001 2 -Cx/B/B-C	">
 >*=>>> 
//0 	 1>'	'+'0 .D0C0C.D	, !!0F(2E2E0F	 ,	. .D0C0C.D 2233	3 ,	3
 3 
3 43l 112,C 3B ((7M9L9L7MNG OG( **9O8;N;N9O&& 
**+ ,
 	//0X	X 
X 1Xv 223 4" 99:( ;( **+# ,#
 **+D ,D --.- /- ,,-	 .	 **+ , 223	(7 4 112K 3K ((7M9L9L7MN	 O: ++,0 -0 ../) 0)* 		") #)X 556	! ! 7 
112J 3J
 ../ 0ey Js   #
Ac*