
    Vpf                   
   U d dl mZ d dlZd dlZd dlmZmZmZmZ d dl	Z	d dl
Z
d dl
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Zd dlZd dlmZmZmZmZmZ d dl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-m0Z0 d dl1m2Z2m3Z3 d dl4m5Z6 d dl7m8Z9 d dl7m: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lCmDZD e*jE        eFcZFZGe*jH        eIcZIZJ ejK        d          ZLeZMd ZN ejO        d! ejP        d"d#          d$%          ZQ ejO        d& ejP        d'd(          d)%          ZRe"jS        ZSd*eTd+<   ee=jM        eUe=jM        d,f         f         ZVdd1ZWdd3ZXe=jY        jZ        Z[dd7Z\dd9Z]d: Z^d; Z_dd?Z`d@ Zaee=jb        eUe=jb        d,f         f         ZcddAZdi  eje        e jf                   ee=jg        jh        dB           eje        eji                   ee=jg        jh        dB           eje        e jj                   ee=jg        jh        dC           eje        ejk                   ee=jg        jh        dD           eje        ejl                   ee=jg        jh        dE           eje        ejm                   ee=jg        jh        dF           eje        ejn                   ee=jg        jh        dG           eje        e jo                   ee=jg        jp        dC           eje        ejq                   ee=jg        jp        dD           eje        ejr                   ee=jg        jp        dE           eje        ejs                   ee=jg        jp        dF           eje        ejt                   ee=jg        jp        dG           eje        e ju                  e=jv        jZ         eje        e jw                  e=jx        jZ         eje        e jy                  e=jz        jZ         eje        e j{                  e=j|        jZ         eje        e j}                  e=j~        jZ         eje        e j                  e=j        jZ         eje        ej                  e=j        jZ         eje        ej                  e=j        jZ         eje        ej                  e=j        jZ         eje        ej                  dH  eje        ej                  dI iZdJeTdK<   e j        ae j        J  ee=jg        jh        dL          e eje        e j                  <    ee=jg        jp        dL          e eje        e j                  <   ddPZddSZddUZi ZdVeTdW<   ddZZeeej        <   eeej        <   d[ eej        <   eeej        <   dd]Z G d^ d_e          Zi Zd`eTda<   ddeZddfZddiZddkZdl Z eej        j        e           ddmZ eej        e           ejk        ejl        ejm        ejn        ejq        ejr        ejs        ejt        ej        ej        ej        ej        ej        eji        ej        e j        fD ]Z eee           dn Ze j                                        D ]\  ZZe ee eeee                     do Z eej        e           ddtZddwZdd{ZddZ e=j                    ZeBr eBj        e            ej                    ZddZddZddZdddZddZddZee(j        e(j        e(j        f         Z G d d          Z e	j        d           G d d                      Ze	j         G d ds                      Ze	j         G d dv                      Ze	j         G d d                      ZeNs G d de          ZneZi ZdeTd<   deTd<    ej        eŦ          Z	 dddZǐddZ eɦ            Zd ZːddZ eɦ            ZʐddZ͐ddZΐddZ ej        d          ZѐddZҐddZӐddZԐddZՐddÄZ G dĄ de          Zg dƢZؐddɄZِdd˄Zڐdd΄ZېddτZddddddddBdBdddddМddZd Ze=jM        ZeAj        jZ        ZeAj        Z G d d          Zd dddddddddddddddd̐dZd͐dZdΐdZdϐdZdАdZdѐdZdҐd$ZdӐd%ZdԐdՐd)Z	 d֐d*Zdאd-Zddd.dؐd/Zdd0dِd1Z eej         eed23                      eej         eed3                     d4 Zd5 Zdڐd6Zdېd;Zdڐd<Zdܐd=Zdܐd>Zdܐd?Zdܐd@ZdݐdBZdސdCZdD Z eej        e            eej        dE            ddߐdHZ dI Z eeeAj        dJ          Z eeeAj        dK          ZdِdLZ	 	 	 dddRZ eedSdT          Z eedU          Z	 eedV          Z
ddWZddYZdd\Zd] Zdd`Z	 dddfZddgZdBZdLZdhZddiZddddoZddddrZdddd~ZddZd ZddddddZddZd#d dddLddddd	ddZddZddZdS (      )annotationsN)CallableIterableIteratorSequence)partial)Any
NamedTupleProtocolUnioncast)ad_util)config)core)dtypes)effects)linear_util)path)pickle_utilsharding)sharding_impls)source_info_util)util)
xla_bridge)partial_eval)xla)
AutoLayoutDeviceLocalLayout)Sharding)
xla_client)xla_extension)dialects)ir)func)hlo)register_jax_dialects)AbstractRefTFjax_dump_ir_toJAX_DUMP_IR_TO zPath to which the IR that is emitted by JAX should be dumped as text files. If omitted, JAX will not dump IR. Supports the special value 'sponge' to pick the path from the environment variable TEST_UNDECLARED_OUTPUTS_DIR.)helpjax_include_debug_info_in_dumpsJAX_INCLUDE_DEBUG_INFO_IN_DUMPSTruezDetermine whether or not to keep debug symbols and location information when dumping IR code. By default, debug information will be preserved in the IR dump. To avoid exposing source code and potentially sensitive information, set to falsezeffects_lib.EffectTypeSetlowerable_effects.xIrValuesreturnboolc                8    t          | t          j                   S )z,Returns true if `x` is not a block argument.)
isinstancer$   BlockArgumentr2   s    Z/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/_src/interpreters/mlir.py_is_not_block_argumentr;   ^   s    2+,,	,,    ir.DenseIntElementsAttrc                z    t           j                            t          j        | t          j                            S N)r$   DenseIntElementsAttrgetnpasarrayint64xss    r:   dense_int_elementsrG   c   s'    		 	$	$RZBH%=%=	>	>>r<   rF   Sequence[bool]ir.DenseElementsAttrc                   t          j        t          j        | t           j                  d          }t	          |           dk    r9t          j        |                                dk    rdndt           j                  }t          j        	                    |t          j
                            d          t	          |           g          S )Nlittlebitorder   r      typeshape)rB   packbitsarraybool_lenitemuint8r$   DenseElementsAttrrA   IntegerTypeget_signless)rF   as     r:   dense_bool_elementsr]   h   s    k"(2rx((8<<<! 	WW\\
affhh!mmrx88A			!	!bn))!,,SWWI 
" 
? 
? ?r<   ir.DenseBoolArrayAttrc                @    t           j                            |           S r?   )r$   DenseBoolArrayAttrrA   rE   s    r:   dense_bool_arrayra   q   s    			"	"2	&	&&r<   c                |    t           j                            t           j                            d          |           S )N    r$   IntegerAttrrA   rZ   r[   is    r:   i32_attrrh   t   )    **2>+F+Fr+J+JANNNr<   c                |    t           j                            t           j                            d          |           S )N@   rd   rf   s    r:   i64_attrrl   u   ri   r<   sizes#Sequence[int | ir.RankedTensorType]ir.RankedTensorTypec           	        t          t          j        dt          j                            t          t          j        dt          j                            fd}t          ||           }|sDt          t          j        t          t          j
        g t          j                                      S t          |          dk    r|d         S t          j        |t          d                    S )NrN    c                    t          |           t          u r-t          t          j        | gt          j                            S | j         k    rt          j        |           } t          j        |           S r?   )	rQ   intir_constantrB   rT   int32r&   convertreshape)di32_typeint1ds    r:   	lower_dimzshape_tensor.<locals>.lower_dim{   s`    Aww#~~1#rx00111	
8		K!$$["""r<   rN   r   )aval_to_ir_typer   ShapedArrayrB   rv   map	type_castr$   RankedTensorTyperu   rT   rV   r&   concatenaterl   )rm   r|   dsrz   r{   s      @@r:   shape_tensorr   w   s    
$*4::
;
;%T-b"(;;<<(# # # # # # 9e"	 ,R(+bhr286L6L*M*MNNN
2ww!||a5L?2x{{+++r<   c                f     | j         di |} ||g|R  }|                     |j                   |S )zSide-effects on `ctx`rr   )replaceset_tokens_out
tokens_out)ctxlowering_funargsctx_override_kwargsctx_newouts         r:   delegate_loweringr      sM    CK..-..'W$t$$$#W'(((	*r<   c                    t          | t          j                  rdS t          | t                    o+t	          |           dk    ot          d | D                       S )z8Returns true if `x` is an ir.Value or tuple of ir.ValuesTrN   c              3  J   K   | ]}t          |t          j                  V  d S r?   )r7   r$   Value.0vs     r:   	<genexpr>z _is_ir_values.<locals>.<genexpr>   s.      55a*Q))555555r<   )r7   r$   r   tuplerV   allr9   s    r:   _is_ir_valuesr      s]    28 4
Q

 63q66Q; 6551555557r<   rN            rc   rk   c                 x    t           j                            t           j                                                  S r?   )r$   ComplexTyperA   F32Typerr   r<   r:   <lambda>r      s"    "."4"4RZ^^5E5E"F"F r<   c                 x    t           j                            t           j                                                  S r?   )r$   r   rA   F64Typerr   r<   r:   r   r      s"    2>#5#5bjnn6F6F#G#G r<   z%dict[np.dtype, Callable[[], ir.Type]]_dtype_to_ir_type   dtype!core.bint | np.dtype | np.genericir.Typec                   t          | t          j                  rt          j        t          j                  } t          | t          j        t          j        f          sJ t          |                       t          j        |           } 	 t          |          }n%# t          $ r}t          d|            |d }~ww xY w |            S )Nz'No dtype_to_ir_type handler for dtype: )r7   r   bintrB   r   rv   genericrQ   r   KeyError	TypeError)r   ir_type_factoryerrs      r:   dtype_to_ir_typer      s    ty!!  HRXE	EBHbj1	2	2??DKK???
(5//%D'.OO	 D D D
9%99; ;@CDD 
		s   B 
B;#B66B;aval$core.ShapedArray | core.DShapedArrayc                    t          j        |           } t          j        | j                  st	          |           S t
          j                            | j        t          | j	                            S r?   )
r   physical_avalis_constant_shaperR   _dynamic_array_ir_typesr$   r   rA   r   r   )r   s    r:   _array_ir_typesr      s[    		D	!	!$		
	+	+ )"4(((			 	 -=dj-I-I	J	JJr<   core.ShapedArrayc                    t           j                                        fd| j        D             }t           j                            |t          | j                            S )Nc                D    g | ]}t          |          t          u r|nS rr   )rQ   rt   )r   ry   dyn_sizes     r:   
<listcomp>z+_dynamic_array_ir_types.<locals>.<listcomp>   s+    
A
A
AQ311H
A
A
Ar<   )r$   
ShapedTypeget_dynamic_sizerR   r   rA   r   r   )r   rR   r   s     @r:   r   r      sU    ]++--(
A
A
A
Adj
A
A
A%			 	 (8(D(D	E	EEr<   z8dict[type[core.AbstractValue], Callable[[Any], IrTypes]]ir_type_handlerscore.AbstractValueIrTypesc                    	 t          t          |                    |           S # t          $ r%}t          dt          |                      |d}~ww xY w)zConverts a JAX aval to zero or more MLIR IR types.

  In general, a JAX value may be represented by multiple IR values, so this
  function may return a tuple of types.z"No ir_type_handler for aval type: N)r   rQ   r   r   )r   r   s     r:   r}   r}      sc    
PDJJ'---	 P P P
EdEE
F
FCOPs   !$ 
A AAc                >    t           j                                        S r?   )r&   	TokenTyperA   )_s    r:   r   r      s    1B1B1D1D r<   tuple[ir.Type, ...]c                ^    t          |           }t          |t          j                  r|fn|S r?   )r}   r7   r$   Type)r   typs     r:   aval_to_ir_typesr      s,    #c27++	4#4r<   c                      e Zd ZddZdS )ConstantHandlervalr	   r4   r3   c                    dS )zlBuilds an IR representation for a constant `val`.

    A JAX value is represented by zero or more IR values.Nrr   )selfr   s     r:   __call__zConstantHandler.__call__         r<   Nr   r	   r4   r3   __name__
__module____qualname__r   rr   r<   r:   r   r      s(        = = = = = =r<   r   zdict[type, ConstantHandler]_constant_handlerstype_rQ   handler_func                    |t           | <   d S r?   r   )r   r   s     r:   register_constant_handlerr      s    )Ur<   c                    t           |          S r?   r   )r   s    r:   get_constant_handlerr     s    	E	""r<   r   r	   c                v   t          |           j        D ]U}t                              |          }|r7 ||           }t	          |          sJ t          |           |f            |c S Vt          | d          r!t          |                                           S t          dt          |                      )zTranslate a Python `val` to an IR constant, canonicalizing its dtype.

  Args:
    val: a Python value to be translated to a constant.

  Returns:
    A representation of the constant as an IR value or sequence of IR values.
  __jax_array__zNo constant handler for type: )	rQ   __mro__r   rA   r   hasattrru   r   r   )r   thandlerr   s       r:   ru   ru     s     99  a $$Q''G GCLLc311$s))S!1111jjj S/"" ,s((**+++>499>>???r<   np.ndarray | np.genericc                "   t          | j                  }| j        }| j        t          j        k    rt          j        | d          } t          j        |           } t          j        	                    | ||          }t          j        |          S )NrK   rL   rP   )r   r   rR   rB   rU   rS   ascontiguousarrayr$   rY   rA   r&   constant)r2   element_typerR   attrs       r:   _numpy_array_constantr     sx    !!'**,
'%W
A)))A1!			!	!!,e	!	D	D$	d		r<   c                      t          d          )Nznumpy masked arrays are not supported as direct inputs to JAX functions. Use arr.filled() to convert the value to a standard numpy array.
ValueError)r   kwargss     r:   _masked_array_constant_handlerr   !  s     V 	W 	W Wr<   c                   | j         t          j        k    r2t          t	          j        | j        t          j                            S d| j        v r| j	        dk    rt	          j
        t	          j        d| j                            \  t	          j
        t	          j        d| j                            \  }| t          fdt          | j                  D                                }t!          j        t$          j                            | j        t+          |j                             t          |          t-          |                    }|S t          |           S )az  Constant handler for ndarray literals, handling zero-size strides.

  In most cases this function calls _numpy_array_constant(val) except it has
  special handling of arrays with any strides of size zero: for those, it
  generates appropriate calls to NumpyArrayConstant, Broadcast, and Transpose
  to avoid staging in large literals that might arise from np.zeros or np.ones
  or the output of lax.broadcast (which uses np.broadcast_to which in turn
  uses size-zero strides).

  Args:
    val: an ndarray.

  Returns:
    An XLA ComputationDataHandle / XlaOp representing the constant ndarray
    staged into the XLA Computation.
  r   r   c              3  B   K   | ]}|v rd nt          d          V  dS r   N)slice)r   axzero_stride_axess     r:   r   z,_ndarray_constant_handler.<locals>.<genexpr>=  sL       9 9"$ $&)9#9#9aauT{{ 9 9 9 9 9 9r<   )r   r   float0r   rB   zerosrR   rU   stridessizewhereequal	not_equalr   rangendimr&   broadcast_in_dimr$   r   rA   r   dense_int_array)r   
other_axescollapsed_valr   r   s       @r:   _ndarray_constant_handlerr  '  s6   " 	Y&- #)28!D!D!DEEECKCHqLL!S[!9!9::(2<3;7788KJ 9 9 9 9(-ch9 9 9 9 9 :M


I'(;<<	> 	>m,,
##	% %C
 J %%%r<   c                F    t          t          j        ||                     S r?   )r   rB   rT   )r   r   s     r:   _python_scalar_handlerr  Q  s    	rxU33	4	44r<   c                (    t          j                    S r?   )r&   create_token)r   s    r:   _token_constant_handlerr	  W  s    				r<   	file_namestrcachesTracebackCachesc                    |j                             | d           }||S t          j        j        }|rt          j        |d|           } | |j         | <   | S )Nr,   )canonical_name_cacherA   r   &hlo_source_file_canonicalization_regexvalueresub)r
  r  canonical_file_namepatterns       r:   get_canonical_source_filer  ]  s`    377	4HH$9?' /wI..I+4&i(	r<   r   ModuleContextc                    | j         j                            |d           }||S t          j        |          }|| j         j        |<   |S r?   )traceback_cachesis_user_file_cacherA   r   is_user_filename)r   r
  is_userr   s       r:   _is_user_filer  h  sM     377	4HH'N))44#7:#))4	*r<   tbxc.Tracebackir.Locationc                   | j         j                            |d          }||S g }t          j        j        }|dk    r|nd}|                                \  }}t          |          D ]\  }}t          | |j	                  s||         }	||	f}
| j         j
                            |
d          }|t          j        ||	          }t          j                            t!          |j        | j                   |j        |j                  }t          j                            |j        |          }|| j         j
        |
<   |                    |           t/          |          |k    r nt/          |          }|dk    rt          j                                        }n=|dk    r	|d         }n.t          j                            |d         |dd                   }|| j         j        |<   |S )z8Converts a full traceback to a callsite() MLIR location.Nr   i  childLocrN   )r  traceback_cacherA   r   traceback_in_locations_limitr  
raw_frames	enumerater  co_filenamelocation_cacher   raw_frame_to_framer$   Locationfiler  r
  
start_linestart_columnnamefunction_nameappendrV   unknowncallsite)r   r  loc
frame_locsframes_limitcodeslastisrg   codelasti
code_lastiframefile_locns                 r:   _traceback_to_locationr?  p  s   ,00T::#_J*4:,!-!2!2,--//-%5!!  gad.// 1IEuJ


-
1
1*d
C
CC
{1$>>e!!
#EOS5I
J
J



 h
 KU08DDc8;c)*5c
:,&&e ' 
*oo!!VV
+



CCAvv
Q-CC
+

z!}jn
=
=C-0#&r*	*r<   	primitivecore.Primitiveparamsdict[str, Any]source_infosource_info_util.SourceInfoc                   |j          dt          j        ||           }t          j        j        r<|j        t          j        	                                }nt          | |j                  }nxt          j        |          }|t          j        	                                }nCt          j                            t          |j        | j                  |j        |j                  }t          j                            ||          }|S )N/r"  )
name_stackr   str_eqn_compactr   $include_full_tracebacks_in_locationsr  	tracebackr$   r+  r2  r?  r   
user_framer,  r  r
  r  r-  r.  r/  )r   r@  rB  rD  eqn_strr4  r<  s          r:   _source_info_to_locationrN    s     & : :$Y77: :'06 C$K!!cc"3(=>>cc'44E}K!!ccK6u7:7KM M"-u/AC Cc 	3//#	*r<   module	ir.Module
stage_name
str | Nonec                   t           j        }|sdS |dk    r1t          j                            dd          }|st          d          t          t                    }| j        j	        d         }t          j        |          j        }d| dt          |           d| d	}t          j        |          }|                    d
d
           ||z  }|                    t#          |                      |S )au  Dumps the `module` IR to a file.

  Dumps the module if JAX_DUMP_IR_TO is defined.

  Args:
    module: The module to dump
    stage_name: A name to distinguish different stages of a module, will be
      appended to the `module.name`.

  Returns:
    The name of the file containing the dump if JAX_DUMP_IR_TO is defined and
    the module was dumped, `None` otherwise.
  NspongeTEST_UNDECLARED_OUTPUTS_DIRr,   zFJAX_DUMP_IR_TO='sponge' but TEST_UNDECLARED_OUTPUTS_DIR is not definedsym_namejax_irr   z.mlirT)parentsexist_ok)_JAX_DUMP_IR_TOr  osenvironrA   r   next_ir_dump_counter	operation
attributesr$   
StringAttr_make_string_safe_for_filenamer   Pathmkdir
write_textmodule_to_string)	rO  rQ  out_dir_nameidrV  module_namer/  out_dir	full_paths	            r:   dump_module_to_filerl    s	    !&,	 4X:>>"?DDL E D E E E "(4(h''-+	V"	V	V5kBB	V	VZ	V	V	V$Il##'	--t-,,,n)'//000	+r<   c                6    t          | |          }|rd| dS dS )NzThe module was dumped to .z)Define JAX_DUMP_IR_TO to dump the module.)rl  )rO  rQ  	dumped_tos      r:   dump_module_messagerp    s0    !&*55) 73y333366r<   sc                .    t          j        dd|           S )Nz
[^\w.)( -]r,   )r  r  rq  s    r:   rb  rb    s    	r1	%	%%r<   c                    t          j                    }|(t                              t          j                  }|dv}| j                            ||           |                                S )N)false0)r,  enable_debug_info)	ioStringIOr  lower _JAX_INCLUDE_DEBUG_INFO_IN_DUMPSr  r_  printgetvalue)rO  rw  outputenable_debug_flags       r:   rf  rf    sa    ;==&		"B"HII)?f8IJJJ			r<   bytesc                    t          j                    }| j                            |           |                                S )N)r,  )rx  BytesIOr_  write_bytecoder}  )rO  r~  s     r:   module_to_bytecoder    s7    :<<&!!v!...			r<   
ir.Contextc                    t          j                    } |                     t                     |                                  |                     d           t          j        rt          j                            |            t          j	        
                    |            t          j                            |            t          j                            |            | S )z,Creates an MLIR context suitable for JAX IR.F)r$   Contextappend_dialect_registryupstream_dialectsload_all_available_dialectsenable_multithreadingr#   sdyregister_dialectmhloregister_mhlo_dialectchlor&   )contexts    r:   make_ir_contextr    s    JLL'	!!"3444	%%''' 
&&&\ +L!!'***
-%%g...
-  )))
,(((	.r<   c                  6    e Zd ZU ded<   ded<   ded<   d
dZd	S )ShapePolyLoweringStatetuple[str, ...]dim_varsr5   uses_dim_varshas_platform_index_argumentlowering_platformstuple[str, ...] | Nonec                    |-t          |          dk    rdt          |          z   }d| _        nd| _        t          |          dk    | _        || _        d S )NrN   )_platform_indexTFr   )rV   r   r  r  r  )r   r  r  s      r:   __init__zShapePolyLoweringState.__init__   s^     %#.@*A*AA*E*E%h7h)-d&&).d&h--!+DDMMMr<   N)r  r  r  r  r   r   r   __annotations__r  rr   r<   r:   r  r    sT          
  $###	 	 	 	 	 	r<   r  T)frozenc                  H    e Zd ZU dZded<   dZded<   dZded<   dZded<   dS )	LoweringParametersNz1tuple[tuple[core.Primitive, LoweringRule]] | Noneoverride_lowering_rulesFr5   global_constant_computation
for_export#export_ignore_forward_compatibility)r   r   r   r  r  r  r  r  rr   r<   r:   r  r  +  sc         
 PTSSSS ',++++ * /4%333333r<   r  c                  >    e Zd ZU ded<   ded<   ded<   ded<   d	 Zd
S )r  zdict[xc.Traceback, ir.Location]r$  z-dict[tuple[types.CodeType, int], ir.Location]r)  zdict[str, str]r  zdict[str, bool]r  c                >    i | _         i | _        i | _        i | _        d S r?   )r$  r)  r  r  r   s    r:   r  zTracebackCaches.__init__H  s(    DD "D Dr<   Nr  rr   r<   r:   r  r  A  sU         2222????&&&&%%%%! ! ! ! !r<   c                     e Zd ZU dZ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<   ed4d            Zd d d d d d d d!d5d(Zed6d*            Zd7d,Z	d8d0Z
d9d2Zd3 Zd S ):r  z2Module-wide context information for MLIR lowering.r  r  rP  rO  zir.InsertionPointipzir.SymbolTablesymbol_tablestr | xb.XlaBackend | Nonebackend_or_nameSequence[str]	platformsAxisContextaxis_context	list[Any]
keepalivesIterator[int]channel_iteratorhost_callbacksr  shape_poly_statezdict[Any, func_dialect.FuncOp]cached_primitive_loweringsr  r  r  lowering_parametersr4   sharding_impls.AxisEnvc                    | j         j        S r?   )r  axis_envr  s    r:   r  zModuleContext.axis_envh  s    %%r<   N)r  rO  r  r  r  r  r  ir.Context | Noneir.Module | Noneir.InsertionPoint | Noneir.SymbolTable | None%None | dict[Any, func_dialect.FuncOp]None | TracebackCachesc               D   |pt                      | _        |	pAt          j                            t          j                            | j                            | _        |
pt          j        | j        j	                  | _
        |pt          j        | j        j                  | _        || _        || _        || _        |i n|| _        |t%                      n|| _        || _        || _        || _        |pt/          dt1          |                    | _        || _        d S )N)r4  rr   )r  r  r$   Modulecreater+  r2  rO  InsertionPointbodyr  SymbolTabler_  r  r  r  r  r  r  r  r  r  r  r  r   r  r  )r   r  r  r  r  r  r  r  r  rO  r  r  r  r  r  s                  r:   r  zModuleContext.__init__l  s   $ /o//DLSBI,,1D1DT\1R1R,SSDK7B%dk&677DG$Mt{7L(M(MD*DDN$D-G-Orr,F 	#2B2J_..."2 	,D DO(DF0U95E5EFF 	2Dr<   xb.XlaBackendc                    t          | j                  dk    rt          d          | j        t	          | j        t
                    rt          j        | j                  S | j        S )NrN   zaccessing .backend in multi-lowering setting. This can occur when lowering a primitive that has not been adapted to multi-platform lowering)rV   r  NotImplementedErrorr  r7   r  xbget_backendr  s    r:   backendzModuleContext.backend  si     4>Q	   #z$2F'L'L#^D0111r<   rt   c                *    t          | j                  S r?   )r]  r  r  s    r:   new_channelzModuleContext.new_channel  s    %&&&r<   host_callbackr	   Nonec                :    | j                             |           d S r?   )r  r1  )r   r  s     r:   add_host_callbackzModuleContext.add_host_callback  s    }-----r<   	keepalivec                :    | j                             |           d S r?   )r  r1  )r   r  s     r:   add_keepalivezModuleContext.add_keepalive  s    O9%%%%%r<   c                &    t          j        | fi |S r?   dataclassesr   r   kws     r:   r   zModuleContext.replace      +"5d"A"Ab"A"AAr<   )r4   r  )r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  rO  r  r  r  r  r  r  r  r  r  )r4   r  )r4   rt   )r  r	   r4   r  )r  r	   r4   r  )r   r   r   __doc__r  propertyr  r  r  r  r  r  r   rr   r<   r:   r  r  N  s        ::---- !!!!**** =<<< $###))))& & & 8& $(!%%),0LP15!"3 "3 "3 "3 "3 "3H 	  	  	  8	 ' ' ' '. . . .& & & & BAAAAr<   c                      e Zd ZU dZded<   ded<   ded<   ded	<   d
ed<   ded<   ded<   dZded<   dZded<   dZded<   dZded<   ddZ	d Z
d dZdS )!LoweringRuleContextz/Per-rule context information for MLIR lowering.r  module_contextsource_info_util.NameStackrH  zcore.Primitive | Noner@  Sequence[core.AbstractValue]avals_inr	   	avals_outTokenSet	tokens_inzTokenSet | Noner   Nzdict[core.Var, ir.Value] | Noneaxis_size_envrr   Sequence[ir.Value]dim_var_valuescore.JaxprEqnContext | Nonejaxpr_eqn_ctxSequence[str] | Noner  c                6    | j         
J d            || _         d S )Nz"Should only set `tokens_out` once.)r   )r   r   s     r:   r   z"LoweringRuleContext.set_tokens_out  s$    ?""$H""" DOOOr<   c                &    t          j        | fi |S r?   r  r  s     r:   r   zLoweringRuleContext.replace  r  r<   r4   r5   c                8    | j         j        }|j        o|j         S )zOReturns true if the lowering parameters are in forward compatibility mode.
    )r  r  r  r  )r   r  s     r:   is_forward_compatz%LoweringRuleContext.is_forward_compat  s*     -A& 	H#GGr<   )r   r  )r4   r5   )r   r   r   r  r  r  r  r  r  r   r   r  rr   r<   r:   r  r    s         77((((""""((((...37-7777 (*.))))/3-3333 %))((((! ! ! ! BAA     r<   r  c                      e Zd Zd	dZdS )
LoweringRuler   r  r   ir.Value | Sequence[ir.Value]r4   'Sequence[ir.Value | Sequence[ir.Value]]c                    dS )z.Converts a JAX primitive invocation into MLIR.Nrr   )r   r   r   r  s       r:   r   zLoweringRule.__call__  r   r<   N)r   r  r   r  r4   r  r   rr   r<   r:   r  r    s(        ; ; ; ; ; ;r<   r  z"dict[core.Primitive, LoweringRule]
_loweringsz-dict[str, dict[core.Primitive, LoweringRule]]_platform_specific_loweringsprimruleplatformc                n    ||t           | <   n't          j        |          D ]}|t          |         | <   |S r?   )r  r  expand_platform_aliasr  )r  r  r  ps       r:   register_loweringr    sI    Jt %h// 3 3.2"1%d++	+r<   Iterable[IrValues]list[ir.Value]c                    g }| D ]G}t          |t          j                  r|                    |           2|                    |           H|S )z@Concatenates/flattens a list of ir.Values or ir.Value sequences.)r7   r$   r   r1  extendrF   r   r2   s      r:   flatten_ir_valuesr    sQ    
#  a!RX 	jjmmmm	jjmmmm	*r<   c                <    t          |           dk    r| d         n| S )NrN   r   )rV   r9   s    r:   _unwrap_singleton_ir_valuesr	    s    3q66Q;;1Q44A$Er<   Iterable[IrTypes]list[ir.Type]c                    g }| D ]G}t          |t          j                  r|                    |           2|                    |           H|S )z>Concatenates/flattens a list of ir.Types or ir.Type sequences.)r7   r$   r   r1  r  r  s      r:   flatten_ir_typesr    sQ    
#  a!RW 	jjmmmm	jjmmmm	*r<   Iterable[ir.Type]nsSequence[int]list[IrTypes]c                ~    t          |           fd|D             }t          t                    t          u sJ |S )zSplits `xs` into subsequences of lengths `ns`.

  Unlike `split_list`, the `sum(ns)` must be equal to `len(xs)`, and if n == 1
  then types are not wrapped in a singleton list.c                    g | ]?}|d k    rt                    n't          fdt          |          D                       @S )rN   c              3  6   K   | ]}t                    V  d S r?   r]  r   r   xs_iters     r:   r   z0unflatten_ir_types.<locals>.<listcomp>.<genexpr>  s=       4$ 4$ 59MM 4$ 4$ 4$ 4$ 4$ 4$r<   )r]  r   r   )r   r>  r  s     r:   r   z&unflatten_ir_types.<locals>.<listcomp>  sq     1 1 1() #$q&&ge 4$ 4$ 4$ 4$((4$ 4$ 4$ /$ /$ 1 1 1r<   iterr]  _unflatten_done)rF   r  unflattenedr  s      @r:   unflatten_ir_typesr  
  sY    
 HH'1 1 1 1-/1 1 1+	g	'	'?	:	:	:	:	r<   rt   c                X    t          | t          j                  rdnt          |           S )NrN   )r7   r$   r   rV   r9   s    r:   len_ir_typesr    s#    BG$$	0#a&&0r<   Iterable[ir.Value]ysSequence[IrTypes]list[IrValues]c                ~    t          |           fd|D             }t          t                    t          u sJ |S )zSplits `xs` into subsequences of lengths `ns`.

  Unlike `split_list`, the `sum(ns)` must be equal to `len(xs)`, and if n == 1
  then values are not wrapped in a singleton list.c           
         g | ]`}t          |t          j                  rt                    n4t	          fd t          t          |                    D                       aS )c              3  6   K   | ]}t                    V  d S r?   r  r  s     r:   r   z<unflatten_ir_values_like_types.<locals>.<listcomp>.<genexpr>"  s)      <<tG}}<<<<<<r<   )r7   r$   r   r]  r   r   rV   )r   yr  s     r:   r   z2unflatten_ir_values_like_types.<locals>.<listcomp>!  sp        #-Q"8"8 =g<<<<eCFFmm<<<<<  r<   r  )rF   r   r  r  s      @r:   unflatten_ir_values_like_typesr'    s[     HH'     + 
g	'	'?	:	:	:	:	r<   z[^\w.-]r   JSharding | Nonec                    || S t          | t          j                  r| S t          | t          j        t          j        f          st
          |                      |j        | j                            S )z5Returns the new aval sharded based on sharding proto.)	r7   r   AbstractTokenr~   DShapedArrayr  updateshard_shaperR   )r   r   s     r:   sharded_avalr.  *  sn     Kd()) K	D4+T->?	@	@ 
	)X)$*55	6	66r<   rR   
core.Shapetuple[int | Value, ...]c           	         t           j        j        r$ j        J t	           fd|D                       S                      dt          j                    gt           j	        j
        j                  z  d             t          t          t          j        | j	        j
        j                  d           g j        R  }t	          d t!          |t#          |                    D                       S )Nc              3  N   K   | ]}j                             ||          V   d S r?   )r  rA   )r   ry   r   s     r:   r   z%eval_dynamic_shape.<locals>.<genexpr>:  s6      <<"&&q!,,<<<<<<r<   eval_dynamic_shape)r@  r  r   T)multiple_resultsc              3  p   K   | ]1\  }}t          j        |          rt          j        |          n|V  2d S r?   )r   is_constant_dimoperatorindex)r   ry   d_irs      r:   r   z%eval_dynamic_shape.<locals>.<genexpr>D  sa       D DQ '+&:1&=&=G"""4 D D D D D Dr<   )r   dynamic_shapesr  r  r   r   r   dim_value_avalrV   r  r  r  	lower_funr   evaluate_shaper  zipr  )r   rR   ress   `  r:   r3  r3  6  s%     D(((<<<<e<<<<<<
++&%''(3s/A/R/[+\+\\   C
)#UC,>,O,XYY  "9%(%79 9 9C  D D #E+<S+A+A B BD D D D D Dr<   tuple[Value, ...]c                ^    ddt          fdt          | |          D                       S )z-Evaluates the dynamic shapes as int32 values.ry   int | Valuec                $   t          |           t          u r-t          t          j        | t          j                            S t          t          j        dt          j                            }| j         |k    rt          j
        ||           S | S )Nr   rr   )rQ   rt   ru   rB   rT   rv   r}   r   r~   r&   rw   ry   rz   s     r:   convert_dimz/eval_dynamic_shape_as_vals.<locals>.convert_dimK  sp    Aww#~~!28444555 !1"bh!?!?@@h	
8		{8Q'''r<   c              3  .   K   | ]} |          V  d S r?   rr   r   r   rE  s     r:   r   z-eval_dynamic_shape_as_vals.<locals>.<genexpr>T  +      FF!{{1~~FFFFFFr<   )ry   rB  r   r3  r   rR   rE  s     @r:   eval_dynamic_shape_as_valsrK  H  sG        
FFFF'9#u'E'EFFF	F	FFr<   c                ^    ddt          fdt          | |          D                       S )z7Evaluates the dynamic shapes as int or ir.int32 values.ry   rB  r4   int | ir.Valuec                    t          |           t          u r| S t          t          j        dt
          j                            }| j         |k    rt          j        ||           S | S )Nrr   )	rQ   rt   r}   r   r~   rB   rv   r&   rw   rD  s     r:   rE  z0eval_dynamic_shape_as_ivals.<locals>.convert_dim[  sV    Aww#~~h !1"bh!?!?@@h	
8		{8Q'''r<   c              3  .   K   | ]} |          V  d S r?   rr   rG  s     r:   r   z.eval_dynamic_shape_as_ivals.<locals>.<genexpr>d  rH  r<   )ry   rB  r4   rM  rI  rJ  s     @r:   eval_dynamic_shape_as_ivalsrP  W  sG        
FFFF'9#u'E'EFFF	F	FFr<   r   c                <    t          t          | |                    S )z4Evaluates the dynamic shapes as one 1d int32 tensor.)r   r3  )r   rR   s     r:   eval_dynamic_shape_as_tensorrR  f  s     
(e44	5	55r<   c                  8    e Zd ZU ded<   ded<   ded<   ded<   d	S )
LoweringResultrP  rO  
Any | Noner  r  r  r  r  N)r   r   r   r  rr   r<   r:   rT  rT  k  sB         ******r<   rT  )cpucudarocmtpuaxis_ctxsharding_impls.SPMDAxisContextc                   | j         }t          |t          j                  rL|j         j        |j        k    r7t          j                            |j         |j        |j        | j                  S t          j	         |j
        |          |          d         }t          j                            |||j        | j                  S )N)memory_kind_manual_axesr   )meshr7   r   NamedShardingrR   _from_parsed_pspec_parsed_pspecr]  manual_axesparse_flatten_op_sharding_to_xla_hlo_sharding)rZ  r   r   r_  parsed_pspecs        r:   add_manual_axesrg  u  s    	$>788 
)mTZ'''::x-8;O) ; + + + ";#h#D))41 1124L'::
Lh&:' ; ) ) )r<   0xc.OpSharding | sharding.SdyArraySharding | Nonec                   |d S t          |t                    sJ t          |t                    rt          | |j        |          S t          |t
          j        t
          j        f          sJ t          j	        |j
        t          j                  r)t          j        ||          }t          j        |          }| j        }t          |t          j                  r|j        rt%          |||j                  }t(          j        j        r |j        |j                  S  |j        |j                                                  S r?   )r7   	JShardingr(   _to_physical_op_sharding
inner_avalr   r~   r+  r   
issubdtyper   extendedr   physical_shardingr   r  SPMDAxisContextrc  rg  r   r   use_shardy_partitionerr  _to_sdy_shardingre  to_proto)r   r   r   rZ  s       r:   rk  rk    s'    4	Hi	(	((((k"" D#C(CCC	D4+T->?	@	@@@@tz6?33 $/h??Hd##D(>9:: >>x49==H"( 0$8$TY///	&	&ty	1	1	:	:	<	<<r<   layout%DeviceLocalLayout | None | AutoLayoutc                    | dS t          | t                    rdS |t          j        u rdS |                     |j                  S )Ndefaultauto)r7   r   r   abstract_token_to_xla_layoutr   )rt  r   s     r:   rz  rz    sL    ^9
## 6	T   9			tz	*	**r<   c                F    | d S t          | t                    sJ | j        S r?   )r7   rj  r]  rs  s    r:   _get_mem_kindr|    s*    Y4	Ay	!	!!!!	
r<   )replicated_argsarg_shardingsresult_shardings
in_layoutsout_layouts	arg_namesresult_namesnum_replicasnum_partitionsall_default_mem_kindinput_output_aliasespropagated_out_mem_kindsmesh_shape_tupleri  jaxprcore.ClosedJaxprordered_effectslist[core.Effect]r  r  r  r  r  r  rH  r  donated_argsr}  Sequence[bool] | Noner~  !Sequence[JSharding | None] | Noner  r  6Sequence[DeviceLocalLayout | None | AutoLayout] | Noner  r  Sequence[str | None] | Noner  r  r  r  r  r  None | tuple[int | None, ...]r  tuple[None | str, ...] | Noner  r  "tuple[tuple[str, int], ...] | Nonec               	  () t          t          t          j        |                    }|	|j        nt          t
          |j        |	          }|
|j        nt          t
          |j        |
          }|rd}d}n2|	t          t          |	          nd}|
t          t          |
          nd}d}d |D             }|rt          |          t          |          k    r'|st          |          rt          d| d| d          |dk    r|
t          d |
D                       r|}|t          ||||||          \  }}t                              |j                  }|rt!          d|j                   |dt          |          rUd	 t#          ||          D             }d
}|sd| d| }|r-t%          j        dd                    |           d|            ||J ||J ~t+          j        d          }g } g }!t.          j        j        sOd |j        D             }"t          t5          t7          j        d |"t;                                                    }#nd}#t=          |||| ||!|t?          |#|                    }$|$j         5  tB          j"        #                    |$j                   5  |$j$        j%        j&        }%t.          j'        j        rm|J |$j$        j(        )                    tT          j+        ,                    dtT          j+        j-        .                    d |D                                            t^          0                    d|           } tB          j1        .                    |           |%d<   te          |          |%d<   te          |          |%d<   tg          |$d||f|d||	|
|||||||||d ddd           n# 1 swxY w Y   ddd           n# 1 swxY w Y   	 |$j$        j%        4                                s%t!          dtk          |$j$        d          z             nq# tB          j6        $ r_}&dg)()fd(|&j7        D ]}' (|'           t!          d                     )          d z   tk          |$j$        d          z             |&d}&~&ww xY wtq          |$j$        |$j9        |$j:        |$j;                  S )!zLowers a top-level jaxpr to an MLIR module.

  Handles the quirks of the argument/return value passing conventions of the
  runtime.
  Nc                $    g | ]}|t           v |S rr   )_platforms_with_donation)r   r   s     r:   r   z)lower_jaxpr_to_module.<locals>.<listcomp>  s-     ? ? ?1 !%= = =  = = =r<   zeIn multi-platform lowering either all or no lowering platforms should support donation. Lowering for z of which only z support donationrN   c              3     K   | ]}|d u V  	d S r?   rr   r   rq  s     r:   r   z(lower_jaxpr_to_module.<locals>.<genexpr>  s&      'L'LaT	'L'L'L'L'L'Lr<   !Cannot lower jaxpr with effects: c                6    g | ]\  }}|t          |          S rr   r  )r   r\   ry   s      r:   r   z)lower_jaxpr_to_module.<locals>.<listcomp>  s'    LLL41a!LALLLr<   zTSee an explanation at https://jax.readthedocs.io/en/latest/faq.html#buffer-donation.z Donation is not implemented for z.
z&Some donated buffers were not usable: z, c                n    g | ]2}t          |d           |j        D ]}t          j        |          |3S rR   )r   rR   r   r6  )r   r   ry   s      r:   r   z)lower_jaxpr_to_module.<locals>.<listcomp>  si     H H H$GD'4J4J H!ZH Ht/CA/F/FHA H H H Hr<   c                P    |                      |                                          S r?   )union	_get_vars)accnews     r:   r   z'lower_jaxpr_to_module.<locals>.<lambda>	  s    cii>X>X r<   rr   )r  r  r  r  r  r  r  r  r_  c                ^    g | ]*\  }}t           j        j                            ||          +S rr   )r#   r  MeshAxisAttrrA   )r   r/  r   s      r:   r   z)lower_jaxpr_to_module.<locals>.<listcomp>  sA     6 6 6 dD <,00t<< 6 6 6r<   r   rV  zmhlo.num_replicaszmhlo.num_partitionsmainT)rH  publicr}  r~  r  r  xla_donated_argsr  r  arg_memory_kindsresult_memory_kindsarg_layoutsresult_layoutsr  z)Cannot lower jaxpr with verifier errors. verificationz(Cannot lower jaxpr with verifier errors:c                                         d| j                                         d| j                    | j        D ]} |           d S )N	z		at )r1  messagelocationnotes)ry   r>  emit_diagnostic_info	msg_liness     r:   r  z3lower_jaxpr_to_module.<locals>.emit_diagnostic_info=  sp    'AI''(((---...w    !Q   r<   
)<r   r   r  canonicalize_platformin_avalsr.  	out_avalsr|  rV   anyr  r   _set_up_aliasesr1   filter_not_inr   r   r>  warningswarnjoin	itertoolscountr   r:  r  sorted	functoolsreducesetr  r  r  r$   r+  r2  rO  r_  r`  rq  r  r1  r#   r  MeshOpMeshAttrrA   _module_name_regexr  ra  rh   lower_jaxpr_to_funverifyrp  	MLIRErrorerror_diagnosticsrT  r  r  r  )*ri  r  r  r  r  r  rH  r  r}  r~  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  platforms_with_donationunlowerable_effectsunused_donationsmsgchannel_iterr  r  all_dim_polyr  r   attrsery   r  r  s*                                           @@r:   lower_jaxpr_to_moduler    s   < C0)<<==) - 5enn,>> "2":u<2BCC  E )4 M=999:>  .9 }.>????C  ? ?	 ? ? ? 1
"##s9~~55 6-- 6	;1:	; 	;'	; 	; 	;< < <  C'L'L;K'L'L'L$L$L %+:
)\
/,1 ,1(L *77FF J
HHH
I
II#l"3"3LL3x+F+FLLL
`C" CByBBSBBc ?m >		"233> >8;> > ? ? ? !'''%###  ##,* . 
		$ H H%. H H HLVI,-X-X-9355B B C C D DHH Ho )!+'3%3*='=h	'R'R	T 	T 	T# { ; ;BK''44 ; ; J +E$* 9)))	jo
,

l#''6 6$46 6 67 78 89 9 9 %((k::K))+66E*!),!7!7E
#+N#;#;E
 VUO;'#)1)!)/"!9; ; ; ;!; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ;BM:&&(( ;
5
cj.
9
9:; ; ;; 
 
M 
M 
M;<I           
    1
TYYy))D0(^DDE F FKLM
M 

CNC4F,
. 
. .sQ   %O%4DOO%O	O%O	O%%O),O)1AP5 5R#ARR#c                
   | d gt          |          z  } nt          |           } d }t          ||          }t          ||          }|4|2t          d |D                       st          d |D                       r&d gt          |          z  }d gt          |          z  }t	          j        t          j                  }t          t          ||||                     D ]+\  }\  }	}
}}|r|||	|
f         	                    |           ,t          |          }t          t          ||                    D ]F\  }\  }	}|	|f}|
                    |d          r$||                                         }|| |<   d||<   G| |fS )Nc                *    |                                  S r?   )strip_weak_type)r\   s    r:   r   z!_set_up_aliases.<locals>.<lambda>S  s    Q..00 r<   c              3     K   | ]}|d u V  	d S r?   rr   r   r\   s     r:   r   z"_set_up_aliases.<locals>.<genexpr>^  s&      	.	.!t)	.	.	.	.	.	.r<   c              3     K   | ]}|d u V  	d S r?   rr   )r   rs     r:   r   z"_set_up_aliases.<locals>.<genexpr>_  s&      	1	1!t)	1	1	1	1	1	1r<   rr   F)rV   listr   r  collectionsdefaultdictdequer'  r>  r1  rA   popleft)r  r  r  r  r  r  strip_metadata	donationsrg   r   amdonatedaliasedout_donated_argsrmkeyinput_ids                    r:   r  r  K  s   ! 6CMM1 455 10.**(.),,) "5"=		.	.-	.	.	... #>		1	10	1	1	111 #> vH-&3y>>1%k&788))2	($l4HII*K *K & &%a	%$GW &7?r
""1%%%,'' Y0C!D!DEE ) )ma$ *C}}S" )3''))h'(8$#(x 	/	//r<   c                  j    e Zd ZU dZded<   d Zd Zdd
Zedd            Z	ddZ
ddZddZddZdS )r  aY  An immutable container of tokens to be used to lower effectful jaxprs. When lowering
  effectful jaxprs, we need to thread HLO tokens to sequence them. Each effect
  will need its own token that will be threaded in and out of the effectful
  primitives. A `TokenSet` encapsulates a set of HLO tokens that will be
  used by the lowering rules.
  z+collections.OrderedDict[core.Effect, Token]_tokensc                2    t          j        |i || _        d S r?   )r  OrderedDictr  )r   r   r   s      r:   r  zTokenSet.__init__  s    *D;F;;DLLLr<   c                *    t          | j                  S r?   )rV   r  r  s    r:   __len__zTokenSet.__len__  s    t|r<   effectcore.Effectr4   Tokenc                    | j         |         S r?   r  )r   r  s     r:   rA   zTokenSet.get  s    <r<   r   Sequence[core.Effect]c                T    d |D             }t          t          ||                    S )z?Creates a `TokenSet` corresponding to a list of `core.Effect`s.c                *    g | ]}t                      S rr   )r  r   r   s     r:   r   z#TokenSet.create.<locals>.<listcomp>  s    ...lnn...r<   )r  r>  )clsr   tokenss      r:   r  zTokenSet.create  s/     /.g...FC(()))r<   #Sequence[tuple[core.Effect, Token]]c                N    t          | j                                                  S r?   )r   r  itemsr  s    r:   r  zTokenSet.items  s    ##%%&&&r<   set[core.Effect]c                N    t          | j                                                  S r?   )r  r  keysr  s    r:   r   zTokenSet.effects  s    t|  ""###r<   c                :     t           fd|D                       S )zHReturn a subset of the `TokenSet` restricted to a set of `core.Effect`s.c              3  6   K   | ]}|j         |         fV  d S r?   r  )r   effr   s     r:   r   z"TokenSet.subset.<locals>.<genexpr>  s/      @@S$,s+,@@@@@@r<   )r  )r   r   s   ` r:   subsetzTokenSet.subset  s&    @@@@@@@@@@r<   r  c                    g }|                                  D ]P}||j        v r#|                    ||j        |         f           .|                    || j        |         f           Qt          |          S )zRReturns a new `TokenSet` with tokens replaced with ones from the input `TokenSet`.)r   r  r1  r  )r   r  
new_tokensr  s       r:   update_tokenszTokenSet.update_tokens  s    J||~~ 4 4			3s 3455553S 123333Jr<   N)r  r  r4   r  )r   r  r4   r  )r4   r  )r4   r  )r  r  r4   r  )r   r   r   r  r  r  r  rA   classmethodr  r  r   r  r  rr   r<   r:   r  r  y  s           7666< < <          * * * ;*
' ' ' '$ $ $ $A A A A           r<   r  jit)r  r}  r~  r  use_sharding_annotationsr  r  api_namer  r  r  r  r  r  r  r/  r   r  r  r  Sequence[int | None] | Noner  r  r  r  r  r  func_dialect.FuncOpc               x   WX t           j        j                  }t          j        dt          j        t          j                            g|z  }t          t          |          }t          t          |j                  }t          t          |j                  }t          |          Xd |D             }t          j        gXz  }||z   |j        z   }g |||}t          j        gXz  |j        z   }g ||}|
dg|Xz   z  }g ||
}
Xfd|
D             }
|dg|Xz   z  }g ||}|dgXz  }g ||}|dg|Xz   z  }g ||}|dg|Xz   z  } g | |}|dgXz  } g | |}|dg|Xz   z  }!g |!|}|dgXz  }!g |!|}|g dg|Xz   z  |}t          |          }"t          |          }#t          j                            |"|#          }$t%          j        ||$ j                  }%t          j                            |rdnd          |%j        d	<    j                            |%           d}&|/t3          j         fd
t7          |||          D                       }&d}'|,t3          j        d t7          ||          D                       }'d}(|-t3          j        d t7          |||          D                       }(d})|,t3          j        d t7          ||          D                       })d}*|/t3          j         fdt7          |||          D                       }*d}+d},||dt          |          z  }g g }.}-t7          |||          D ]}\  }/}0}1|/)|0'|-                    |/gt;          |1          z             n&|-                    |0gt;          |1          z             |.                    |0gt;          |1          z             ~t3          j        |-          }+t3          j        |.          },d}2|-t3          j        d t7          |||          D                       }2||&|'|(|
|)|Xdk    s|dk    rWd t=          t          |"                    D             }3|fd t7          ||          D             }4t7          |3t3          j        |4                    D ])\  }5}6|6r"t          j                            d          |5d<   *|	rP|&Nt7          |3|&          D ]=\  }5}7|76t@          j!        j"        rtG          |7          |5d<   +tG          |7          |5d<   >|':t7          |3|'          D ])\  }5}8|8"t          j                            |8          |5d<   *|(:t7          |3|(          D ])\  }5}9|9"t          j                            |9          |5d<   *|):t7          |3|)          D ])\  }5}:|:r"t          j                            d          |5d<   *|
t3          j$        tK          t=          t          |#                              t          t:          |                    };g }<t7          ||
          D ]I\  }=}>|>'|<&                    dgt;          |=          z             .|<&                    |;|>                    Jt7          |3|<          D ]\  }5}>|>tO          |>          |5d<   |dk    rKt7           j        j        |3d|                   D ]'\  }?}5t          j                            |?          |5d<   (n3 j(        j)        r'|3D ]$}5t          j                            d          |5d<   %Xdk    r4|3||Xz            }@|@D ]$}5t          j                            d          |5d<   %t          j*                            d |3D                       |%_+        d  t=          t          |#                    D             }AXdk    r1|AdX         }B|BD ]$}5t          j                            d          |5d<   %|rb|AXd         }Ct          |C          t          |          k    r8t7          |C|          D ]'\  }5}Dt          j                            |D          |5d!<   (|	rP|*Nt7          |A|*          D ]=\  }5}7|76t@          j!        j"        rtG          |7          |5d<   +tG          |7          |5d<   >|+:t7          |A|+          D ])\  }5}E|E"t          j                            |E          |5d<   *|2:t7          |A|2          D ])\  }5}9|9"t          j                            |9          |5d<   *t          j*                            d" |AD                       |%_,        |rt          j-        .                                g|Xz   z  }F|D ]T}G|F                    |Grt          j-        /                    |G          nt          j-        .                                           U|%0                    |F          }Hn|%0                                }Ht          j1        |H          5  |Hj2        }It3          j3        |I|Xg          \  }J}K}Kti           |dg dtj          6                    g           dd|J#	  	        W|	s|&Wfd$t7          |I|&|          D             }I|&#|d%k    rWfd&t7          |I|&|          D             }It3          j3        to          |I|          |Xg          \  }K}L}Mtk          t7          ||L                    }N|M}O|)|&                    t3          j8        ||                    }Pn|}Pd' |j9        D             }Qtu           |j;        |P|N|Qg|OR d(|Ji\  }R}Sg }T|D ]*}U|T                    |S                    |U                     +|T&                    |R           ty          |T          }V|	s|*Wfd)t7          |V|*|          D             }V|,!|d%k    rd* t7          |V|,|          D             }V|*#|d%k    rWfd+t7          |V|*|          D             }Vt%          j=        |V           ddd           n# 1 swxY w Y   |%S ),a  Lowers jaxpr and its callees to an IR function.

  Assumes that an MLIR context, location, and insertion point are set.

  Args:
    ctx: the lowering context.
    name: the function name. The name will be uniquified by the symbol table,
      so it is ok to use the same name multiple times.
    jaxpr: the jaxpr to lower.
    effects: a sequence of `core.Effect`s corresponding to an ordering of tokens
      that will be created in or used by the lowered function.
    public: if true, the function's visibility is set to "public".
    replicated_args: if present, annotates arguments as replicated.
    arg_shardings: sharding annotations for each argument (optional).
    result_shardings: sharding annotations for each result (optional).
    use_sharding_annotations: if True, use "mhlo.sharding" annotations on
      parameters and return values to express sharding. If False, use
      hlo.custom_call operators with sharding annotations.
      TODO(b/228598865): remove this option when "mhlo.sharding" annotations are
      propagated on non-entry functions during MLIR->HLO conversion.
    input_output_aliases: optional sequence that maps argument numbers to the
      corresponding output that should alias them.
    xla_donated_args: optional sequence of args to set donation annotations.
    api_name: The name of the higher level primitive which should show up in the
      name stack.
  Returns:
    MLIR func op
  rr   c                *    g | ]}t                      S rr   
token_typer  s     r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  s    ///!///r<   Nc                "    g | ]}|d n|z   S r?   rr   )r   r\   
num_tokenss     r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  s6     ; ; ; ! %&IDD!"Z; ; ;r<   Fr  r  privatesym_visibilityc                \    g | ](\  }}}t          ||          gt          |          z  )S rr   rk  r  r   r\   rq  typesr   s       r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  sQ     	J 	J 	JQ5 #31
-
-	.e1D1D	D 	J 	J 	Jr<   c                :    g | ]\  }}|gt          |          z  S rr   r  )r   mkr  s      r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  s+    XXX	E"U##	#XXXr<   c                X    g | ]'\  }}}t          ||          gt          |          z  (S rr   rz  r  r   lr\   r  s       r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>!  sN     	H 	H 	HQ5 A

	,u"5"5	5 	H 	H 	Hr<   c                :    g | ]\  }}|gt          |          z  S rr   r  )r   
is_donatedr  s      r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>'  s,    hhh0A
E*U++	+hhhr<   c                \    g | ](\  }}}t          ||          gt          |          z  )S rr   r  r  s       r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>,  sQ     	O 	O 	OQ5 #31
-
-	.e1D1D	D 	O 	O 	Or<   r?   c                X    g | ]'\  }}}t          ||          gt          |          z  (S rr   r   r!  s       r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>E  sN     	M 	M 	MQ5 A

	,u"5"5	5 	M 	M 	Mr<   r   c                    g | ]}i S rr   rr   r  s     r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>S  s%     02 02 0202 02 02r<   c                :    g | ]\  }}|gt          |          z  S rr   r  )r   
replicatedr  s      r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>W  s>     B B BCT:uZL<+>+>> B B Br<   Tz!mhlo.is_same_data_across_replicassdy.shardingmhlo.shardingzmhlo.memory_kindzmhlo.layout_modezjax.buffer_donorztf.aliasing_outputzjax.global_constantr,   z	jax.tokenc                L    g | ]!}t           j                            |          "S rr   r$   DictAttrrA   r   r  s     r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  s&    777E		777r<   c                    g | ]}i S rr   rr   r  s     r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  s%     11 11 11b11 11 11r<   zjax.result_infoc                L    g | ]!}t           j                            |          "S rr   r-  r/  s     r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  s&    888%r{u888r<   )	r  rH  r@  r  r  r  r   r  r  c                D    g | ]\  }}}||nt          |||          S r?   wrap_with_sharding_op)r   r\   rq  a_avalentry_lowering_ctxs       r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  sN     M M MaF y!!34F6STUUM M Mr<   r  c                    g | ]M\  }}}|t           j        ur7t          j        |j        t          j                  r|t          ||          n|NS r?   r   ry  r   rm  r   rn  replicate_trailing_dimsr   orq  r\   r6  s       r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  sr        aA t***99 +>?i ""4a
;
;
;NO  r<   c                P    g | ]#}t          t          j        |                    $S rr   ru   r   canonicalize_dtype)r   r2   s     r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  s+    KKKk#03344KKKr<   r  c                D    g | ]\  }}}||nt          |||          S r?   r3  )r   r;  rq  o_avalr6  s       r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  sN     T T TaF y!!34F6STUUT T Tr<   c                @    g | ]\  }}}||nt          |||          S r?   )wrap_with_memory_kind)r   r;  r  r@  s       r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  sI     O O OaV z!!4QFCCO O Or<   c                    g | ]M\  }}}|t           j        ur7t          j        |j        t          j                  r|t          ||          n|NS r?   r8  r:  s       r:   r   z&lower_jaxpr_to_fun.<locals>.<listcomp>  sr        aA t***99 +>?i ""4a
;
;
;NO  r<   )>rV   r  r  r   r~   r   r>  rB   rD   r   r}   r  r  ry  r  r$   FunctionTyperA   func_dialectFuncOpr  ra  r`  r  insertr   flattenr>  r1  r  r   BoolAttrr   rq  r  get_sharding_attr	unflattenr  r  rh   r  r  	ArrayAttr	arg_attrsresult_attrsr+  r2  r/  add_entry_blockr  	arguments
split_listr  r  r  r'  	wrap_nameconstsjaxpr_subcompr  r  return_)Yr   r/  r  r   rH  r  r}  r~  r  r  r  r  r  r  r  r  r  r  r  r  num_dim_varsdim_var_avalsdim_var_typesinput_typesoutput_typestoken_typestoken_avalsinput_avalsoutput_avalstoken_input_output_aliasestoken_shardingstoken_replicated_argstoken_memory_kindstoken_layoutsflat_input_typesflat_output_typesftypefunc_opir_arg_shardingsir_arg_memory_kindsir_arg_layoutsir_donated_argsir_result_shardingsir_result_memory_kinds"custom_call_ir_result_memory_kindsr?  custom_call_respomr  r  ir_result_layoutsrM  replicated_ir_argsr  r)  r   r]  rt  r$  
output_idsaliasesitypesaliasvar_nametoken_arg_attrsrN  token_result_attrsnamed_result_attrsname_mem_kindarg_locsr>  entry_block	flat_argsr  r   
token_argsunflattened_argsr  r   callee_name_stackrS  out_valsr   outsr  flat_outputsr6  r  sY   `                                                                                      @@r:   r  r    s   j S)233,#B(A"((K(KLLMP\\-o}55- OU^44+_eo66,7||*//w///+$%
2++en<+<-<+<<+%&3eoE,.;..,%"&<*+D!EO7O:NO; ; ; ;%9; ; ; fz 9:O6o66M!fz)O<<+;< "G|j'@A@-@@O!<*#<=?+?.>?$*,E.E1DEFlZ78M0M0K0KFZ'M6}6~6N!T5'\J%>?TCST%k22&|44
/

.0A
B
B%e777')+):):'hhi*) *)'%&'"""|	J 	J 	J 	J]KHH	J 	J 	JK K !,XXS9I;5W5WXXXZ Z .\	H 	H[+FF	H 	H 	HI IN /!lhhSIY[fEgEghhhj jO !,	O 	O 	O 	O.>MM	O 	O 	OP P  '+$$'!(3/B+C+C!CrC68K*, , 	9 	9R	RZ

C5<...////

B4,u---... bTL$7$778888!\#..)-o)F)F&	M 	MlKK	M 	M 	MN N
 !		%		(		#		)		$			a			02 02#.//0002 02 02I "B B"?K@@B B B"9dl;M.N.NOO M M
% 	M79{t7L7L%3
4 A$4$@ ,<== A A/%*0 A$5h$?$?E.!!%6x%@%@E/"& #I/B C C E E
%"&(m&7&7&D&D%"
#!y.99 @ @-%&(m&7&7&?&?%"
#""9o>> < <
% 	<&(kood&;&;%"
#'>U3())**++S|-L-LN Nj"$g{,@AA , ,-&%=
..$,v"6"66
7
7
7
7
..E*
+
+
+
+i11 8 8,%(0%$
%a !5!>!*=L=!9; ; C C/(E')}'8'8'B'B#$$C 
	 	< = = =%')}'8'8'<'<#$$A~~!,|j/H"HIo" 3 3%[__T22k((77Y7779 9G11 11-..//11 11 11, !^^%kzk2# 1 1;??400eK <%jkk2
#l"3"3330,?? < <,%#%=#4#4U#;#;   ?"5"A|-@AA ? ?x		(. 	?"3H"="=%

#4X#>#>%
 '|-CDD @ @x		$&M$5$5h$?$? !"\+<== > >v		$&M$5$5f$=$= !))88<888: :'  ,##%%&,*CDH K KooQIbk&&q)))BK4G4G4I4IJJJJ))(33KK))++K	%% B' B'%I  ?9|Z6PQQNAq,zTt//"%%$>	; ; ;
 $ M(8(DM M M M!)-={KKM M Mi #    Y(8+FF	  i '+o&y+>>	z"'$ '$#Az# Wj1122I+D$++DN4,J,JKK$KKelKKKF(U[+Y66 6 6&46 6Hj D ' '
kk*..%%&&&&KK$T**L# T(;(GT T T T!,0C\RRT T Tl *5$&..O O"> N  NO O Ol
 &46>>    \+>MM	  l &&&EB' B' B' B' B' B' B' B' B' B' B' B' B' B' B'H 
.s   Hp//p36p3ir.Valuer]  aval_outc                J   || j         }n5t          |          }t          |t          j                  s
J |            |}t          d|g| gdd          }dt          j                            |          i}t          j                            |          |j	        d<   |j
        S )Nannotate_device_placementTrN   )result_typesoperandshas_side_effectapi_version_xla_buffer_placementmhlo.frontend_attributes)rQ   r}   r7   r$   r   custom_callra  rA   r.  r`  result)r2   r]  r  result_typer   op	dict_attrs          r:   rB  rB    s    &KK
(
#
#Cc27##((S(((K.k]C1F F F"&(9(9+(F(FG).0kooi.H.H"-*+	r<   c                    t          | ||t          j                                                                        t          t          |j                                      S )N)unspecified_dims)r4  xcHloSharding	replicaters  r  r   r   )r   r   r   s      r:   r9  r9    sS     
	3bn..0099;;5++,,
. 
. 
. .r<   c           	        t          |j        j        j                  }t	          t          j        dt          j        t          j
                                      g|z  }t          t          |j                  }t          t          |j                  }t          |j                                                  }d |D             }g |||}g ||}t#          |          }t#          |          }	t$          j                            ||	          }
|j        J t-          j        |j        j        |
|j        j                  }t$          j                            d          |j        d<   |j        j                            |           |                                }t%          j        |          5  tA          |j!        |          }tE          j#        ||t          |j                  g          \  }}}|$                    tK          tM          ||                    |           | g|R  }j'        rg fd|D             |}t-          j(        tS          |                     ddd           n# 1 swxY w Y   |S )	z<Emits the contents of a lowering rule as a private function.rr   c                *    g | ]}t                      S rr   r  r  s     r:   r   z._emit_lowering_rule_as_fun.<locals>.<listcomp>,  s    ,,,!,,,r<   Nr  r  r  )r  r  c                D    g | ]}j                             |          S rr   r   rA   )r   r  sub_ctxs     r:   r   z._emit_lowering_rule_as_fun.<locals>.<listcomp>A  s*    <<<"&&s++<<<r<   )*rV   r  r  r  r}   r   r~   r   r>  rB   rD   r   r  r  r  r  r   r  r$   rD  rA   r@  rE  rF  r/  r  ra  r`  r  rG  rO  r  r'  rP  r   rQ  r   r  r>  r   rU  r  )lowering_ruler   rV  rX  rY  rZ  effsr[  rd  re  rf  rg  r~  r  r  r  r  r  s                    @r:   _emit_lowering_rule_as_funr     s    S'8ABB, D$R)B28)L)LMMNN- OS\22+_cm44,	cm##%%	&	&$,,t,,,+<-<+<<+.;..,%k22&|44
/

.0A
B
B%		"	"	" 2E#&#5#8: : :')+):):9)E)E'%&!((111''))+	%% 	2 	25[* *37?CSVbdghkhudvdvUw3x3x0NJ 0kkHSz-B-B$C$C)7  9 9G=4#3444D DC<<<<t<<<CdCd*400111	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 
.s   6B8I::I>I>
core.Jaxprr  rS  Sequence[IrValues]r   r  r  #tuple[Sequence[IrValues], TokenSet]c               `    d j         vsJ d fdd!d}d"fd}d# fd}	i t          d |D                       s
J |            t          d |D                       s
J |            t          |t          j                  sJ t          |                      t          |          t          |j                  k    sJ ||f            t          |          t          |j                  k    sJ ||f            t           j	        j
                  t          |          k    sJ  j	        j
        |f            t          ||j        |           t          ||j        |           t          j        |          }
|j        D ]}t          |j                  }|j                            ||j        j        z             }t%           |j        |j        |          }t          j        |j        j                  5  |5   |	|j                  }i }d}||}nt/          |j                  p j         D ]b}|j        t2          |         v rt2          |         |j                 ||<   2|j        t4          j        |         v rt9          |j                  ||<   c|j        t:          v rt:          |j                 }n'|j        t4          j        v rt9          |j                  }t?          t@          j!        "                    |j#                            }|$                    |          }t          ||j                  }tK           |j        |j        |t          ||j&                  |d|j        |	  	        }tN          j(        j)        r$fd|D             }|                    |          }t          d |D                       sJ ||f            tU          |tW          |j                  |||j#        g|R i |j        }|r|j,        }| t[          d|j         d|j#         d          |#                                |#                                k    r[t[          d|j         dt]          |#                                           dt]          |#                                                     |/                    |          }ddd           n# 1 swxY w Y   ddd           n# 1 swxY w Y   	 t]          |          }n(# t`          $ r}t[          d| d|           |d}~ww xY wt          |          t          |j&                  k    sJ ||f            t          ||j&        |           t          j1        ||
           t]          fd|j&        D                       |fS )$a  Lowers a jaxpr into MLIR, inlined into an existing function.

  Assumes that an MLIR context, location, and insertion point are set.

  dim_var_values: the list of dimension variables values in the current
    IR function, in the order of ctx.shape_poly_state.dim_vars.
  gpur   	core.Atomr4   r3   c                    t          |           t          j        u r&t          t	          j        | j                            S t          | t          j                  sJ |          S r?   )	rQ   r   Literalru   r   r>  r   r7   Var)r   envs    r:   readzjaxpr_subcomp.<locals>.readU  sR    Aww$,/6677748$$$$$Vmr<   r   c                x    t          |           t          j        u rt          j        | j                  S | j        S r?   )rQ   r   r  r   abstractifyr   r   )r   s    r:   r   zjaxpr_subcomp.<locals>.aval\  s.    Aww$,_QU###Vmr<   core.Varnodec                    |J t          |t          j                  r|}nGt          |          dk    r%t	          j        dt          d           |d         }nt          |          }|| <   d S )NrN   zJAX lowering rules should not wrap singleton values in tuples. It will be an error to wrap a singleton value in a tuple in a future version of JAX.r   )
stacklevelr   )r7   r$   r   rV   r  r  DeprecationWarningr   )r   r  wr  s      r:   writezjaxpr_subcomp.<locals>.writeb  s    $!! 
aa	Ta% 1		. 	. 	. 	.
 G$KKCFFFr<   r@  rA  LoweringRule | Nonec                X    j         j        d S j         j        D ]\  }}| |u r|c S d S r?   )r  r  )r@  r   r  r   s      r:   get_override_lowering_rulez1jaxpr_subcomp.<locals>.get_override_lowering_rules  sI    
6>T*B  4	a 
4r<   c              3  4   K   | ]}t          |          V  d S r?   r   r   s     r:   r   z jaxpr_subcomp.<locals>.<genexpr>}  s*      ,,!]1,,,,,,r<   c              3  4   K   | ]}t          |          V  d S r?   r  r   s     r:   r   z jaxpr_subcomp.<locals>.<genexpr>~  s*      ..!]1......r<   )rH  N)	r  r@  rH  r  r  r  r   r  r  c                    i | ]N}t          |          t          j        u r1|j        D ])}t          |          t          j        u | |          *OS rr   )rQ   r   r+  rR   r  )r   r\   ry   r  s      r:   
<dictcomp>z!jaxpr_subcomp.<locals>.<dictcomp>  si     B B Bd1gg9J.J.J"#' /K.JT!WW-@-@ DDGG-@-@-@-@r<   )r  c              3  4   K   | ]}t          |          V  d S r?   r  r   s     r:   r   z jaxpr_subcomp.<locals>.<genexpr>  s*      44aq!!444444r<   zLowering rule for `z4` needs to set `tokens_out` because it has effects: rn  z4` returns incorrect set of output tokens. Expected: z vs. Actual: -Output of translation rule must be iterable: , got output c              3  .   K   | ]} |          V  d S r?   rr   )r   r   r  s     r:   r   z jaxpr_subcomp.<locals>.<genexpr>  s+      ..1ttAww......r<   )r   r  r4   r3   )r   r  r4   r   )r   r  r  r3   )r@  rA  r4   r  )2r  r   r7   r   	NameStackrQ   rV   invars	constvarsr  r  r   r   	last_usedeqnsrD  r   rH  rN  r@  rB  user_contextrK  _platforms_for_eqn_ctxr   r  r   _backend_specific_translationsxla_fallback_loweringr  _translationsr  effects_libr  	filter_inr   r  r  outvarsr   r:  r  lower_per_platformr  r   r   r   r  r   clean_up_dead_vars)r   r  rH  r  rS  r  r   r   r  r  r  eqnin_nodesrD  r4  override_ruleplatform_rulesdefault_ruler   r   r  r  rule_ctxr  ansr   	out_nodesr  r  r  s   `                           @@r:   rT  rT  F  s    
cm	#	#	#	#             "      #%#	,,t,,,	,	,22d222	..v...	.	.66666	J 0 :	;	;MMT*=M=MMMM	Tc%,''	'	'	'%	'	'	'	VEO,,	,	,	,ufo	,	,	,	S!*	+	+s>/B/B	B	B	BSEYEbdrDs	B	B	BeU_f%%%eU\4   nU##)Z E1 E1c4$$H/)) :: * < <K
"3sz;
O
OC		&s'@	A	A 62 623 62 6200??m02n*.l		"$ (00ACM 	E 	EA]:1=== <Q ? NN1} B1 EEE 5cm D DN1=J&&#CM2,,]c///.s}==,[0::3;GGHHg--((iT3:&&h$ +ck**iQ Q Qh 
		$ AB B B B"*B B B ##-#@@44844444EEsHoEEExS]););-|"{8 !)8 8 8 -0J8 8c
 
 2 (
8CM 8 8),8 8 89 9 9 9#4#4#6#666bCM b b !2!2!4!455b bDI*J\J\J^J^D_D_b bc c c %%j11m62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62p:**ii : : : 22 2,/2 2 3 389:: s88s3;'''''#s'''s{I&&&Ci0000	.......	.	.	66sO   S+KSS+SS+SS++S/	2S/	7T
T,T''T,eqn_ctxr  r  c                (    | dS | j         dk    rdS dS )zFReturns platforms to override based on compute type of jaxpr equation.Nrr   device_host)rV  )compute_type)r  s    r:   r  r    s&     _2]**8	r<   descriptionr  dict[str, LoweringRule]r  r  effects_lib.Effects	rule_argsir.Value | tuple[ir.Value, ...]c           
     
    t           j                  p j        p j        j        }t	          |          dk    r9|                    |d         |          }|t          d| d|d                    g }	i |                                D ]1\  }
}|
|vr
t	          |	          |
<   |	                    |           2fd|D             }|rC|t          d| d|           |D ]}
t	          |	          |
<   |	                    |           |	sJ t	          |	          dk    rG |	d          g|R i |}t           fdt          t          t          |                               |S t	          |          dk    rt	          |	          d	k    sJ ||	f            t	           j                  dk    s
J d
             j        d         }t          t          j        dt"          j                            }|j        |k    rt)          j        ||          }t)          j        |g|t	          |                    }t/          |          D ]\  }
|j                 j                                        }t5          j        |          5  t)          j        t;          t#          j        |
                             g           ddd           n# 1 swxY w Y   t<          j                             |          }t          j!        gt	          |          z   j"        z   }t          t          |          }t)          j        tG          |          |t	          |	                    }t/          |	          D ]J\  }fd                                D             } $                    |          |j                 j                                        }t5          j        |          5   |g|R i |}	 t          |          }n(# tJ          $ r}tM          d| d|           |d}~ww xY wt           fdt          t          |                     j'        8t	          |          t	          j'                  k    sJ fd|D             |z   }t)          j        |           ddd           n# 1 swxY w Y   L|j(        }|r~tS          j*        tW          ||          t	          |          g          \  }} j,        -                    t]          t_          ||                              } 0                    |           |S )a  Emits code for a primitive for the current lowering platform(s).

  For example, given
      platform_rules = dict(tpu=rule0, cpu=rule0)
      default_rule = rule1

  and
      ctx.module_context.lowering_parameters.platforms = ("cpu",)

  emits:
      rule0(ctx, *rule_args, **rule_kwargs)

  In case of multi-platform lowering, e.g., if
      ctx.module_context.lowering_parameters.platforms = ("cpu", "cuda", "tpu")

  emits:
    rule_idx = case current_platform_idx:
                   0: return 0  # cpu rule index
                   1: return 1  # cuda rule index
                   2: return 0  # tpu rule index
    output = case rule_idx
               0: return rule0(*rule_args, **rule_kwargs)
               1: return rule1(*rule_args, **rule_kwargs)

  Args:
   ctx: lowering context.
   description: a string to include in error messages.
   platform_rules: map platform names, e.g., "cpu", "cuda", to
     `LoweringRule`s, for the platforms that have non-default lowering.
   default_rule: an optional rule to use for platforms not in `platform_rules`.
   effects: the set of effects for the current primitive.
   rule_args: the args of the lowering rules.
   rule_kwargs: the kwargs of the lowering rules.
  rN   r   Nz%MLIR translation rule for primitive 'z' not found for platform c                    g | ]}|v|	S rr   rr   )r   r   platform_to_kept_rules_idxs     r:   r   z&lower_per_platform.<locals>.<listcomp>  s1     %M %M %M1()1K(K(K &'(K(K(Kr<   z' not found for platforms c                .    t          | j                  S r?   wrap_compute_type_in_placeownerr;  r   s    r:   r   z$lower_per_platform.<locals>.<lambda>%  s    ,S!':: r<   r   z#Must have a platform_index variablerr   r   )r8  num_branchesc                &    g | ]\  }}|k    |S rr   rr   )r   r   rule_idxrg   s      r:   r   z&lower_per_platform.<locals>.<listcomp>A  s.     1 1 1#.1h"*a--  !"/--r<   r  r  r  c                .    t          | j                  S r?   r  r  s    r:   r   z$lower_per_platform.<locals>.<lambda>N  s    .sAG<< r<   c                D    g | ]}j                             |          S rr   r  )r   r  	inner_ctxs     r:   r   z&lower_per_platform.<locals>.<listcomp>S  s:     1 1 1 )--c22 1 1 1r<   )1r  r  r  r  rV   rA   r  r  r1  r   filterr;   r  r  r}   r   r~   rB   rv   rQ   r&   rw   CaseOpr'  regionsblocksr$   r  rU  ru   r  r  r  ry  r  r  r   r   r   r   resultsr   rQ  r'  r  r  r  r>  r   )r   r  r  r  r   r  rule_kwargsr  r  
kept_rulesr   pruleplatforms_without_specific_ruler~  current_platform_idxrz   rule_idx_opbranchr  rule_out_avalsrZ  case_opplatforms_for_this_ruler  r  r  r  r   rg   r  r  s   `                           @@@r:   r  r    s   R 5S5FGG L!mL/2/A/K  	^^qilL99D|	- 	- 	-'l	- 	-. . .
 $&*/1 &&((  ha	$'
OOq!e%M %M %M %M	 %M %M %M!$ $	A 	A 	A>	A 	AB B B - 6 6&)*oo ##l###	__Z]3::::k::F::::%'8'@'@AA   M	Y!		J1 4 4 4y*6M 4 4 4	S	 	 A	%	%	%'L	%	%	% +A.T-bAAABB((**;x1EFF
H:!5(+I8 8 8+ 	"" J Jda #*1133F		6	"	" J J	k;rx(B1(EFFGGHIIIJ J J J J J J J J J J J J J J/99'BB/'(3+?+??#-O._n55,J'55'#&z??4 4 4' :&&  ga1 1 1 12L2R2R2T2T1 1 1 &=>>I_Q&--//F		6	"	"  tI9	999[99fF%f--		 F F F >&> >5;> > ? ?DE	FF 

<
<
<
<
'
3
3   
		)?##s9+?'@'@@@@@1 1 1 1 /1 1 13<=		k)              " O' #o$Wl;;
? OFG ,,Xc/:@7B 7B .C .C D DJz"""	.sI   ;KK	 K	2R= PR=
P5P00P55A<R==S	S	c                t    d | D             }d |                                 D             fd| D             S )Nc                .    i | ]}t          |          |S rr   rh  )r   consts     r:   r  z_ir_consts.<locals>.<dictcomp>b  s     8882e99e888r<   c                X    i | ]'\  }}|t          t          j        |                    (S rr   r=  )r   id_r  s      r:   r  z_ir_consts.<locals>.<dictcomp>c  sA       
#u 
;s-e4455  r<   c                :    g | ]}t          |                   S rr   r  )r   r  	ir_constss     r:   r   z_ir_consts.<locals>.<listcomp>g  s$    	3	3	35)BuII
	3	3	3r<   )r  )rS  unique_constsr  s     @r:   
_ir_constsr	  a  s\    88888- %++--  ) 
4	3	3	3F	3	3	33r<   funr   r4  c                     d fd}|S )zConverts a traceable JAX function `fun` into a lowering rule.

  The returned function does not use `avals_out`, so callers may pass any value
  as `avals_out`.r   r  c           
        rnfd}t          j        ||          }| j        t          j                    n| j        j        }|5  t          j        j        r| j	        J g | j	        
                                |R }d t          | j	                  D             t          j        dt          j        d                    }|dfgt!          | j	                  z  }fd| j        D             }t          j        |g ||R           }t'          j        |          \  }	}
}n t'          j        || j                  \  }	}
}\   | j        !| j                            | j                  }n| j        }t3          ||	| j        | j        t9          |          g|R d| j        i\  }}|                     |           |cd d d            S # 1 swxY w Y   d S )	Nc                      | i |fS r?   rr   )r   r  r
  s     r:   r   z.lower_fun.<locals>.f_lowered.<locals>.<lambda>p  s    ##t:Jr:J:J9L r<   c                >    i | ]\  }}|t          j        |          S rr   )r   DBIdx)r   rg   ry   s      r:   r  z0lower_fun.<locals>.f_lowered.<locals>.<dictcomp>~  s&    IIIDAqq$*Q--IIIr<   rr   rv   Fc                    g | ]U}t          |          t          j        u r4|                    t	          fd |j        D                                 n|dfVS )c              3  D   K   | ]}                     ||          V  d S r?   rA   )r   ry   idxs     r:   r   z:lower_fun.<locals>.f_lowered.<locals>.<listcomp>.<genexpr>  s/      .N.Nswwq!}}.N.N.N.N.N.Nr<   r  T)rQ   r   r+  r,  r   rR   )r   r\   r  s     r:   r   z0lower_fun.<locals>.f_lowered.<locals>.<listcomp>  sz     / / / "!WW(999 ((.N.N.N.Nag.N.N.N)N)N(OOO?@$H / / /r<   r  r  )lu	wrap_initr  
contextlibnullcontextmanagerr   r:  r  r  valuesr'  r   r~   rB   r   rV   r  annotatepetrace_to_jaxpr_dynamic2trace_to_jaxpr_dynamicr  r  r   rT  rH  r  r	  r  r   )r   r   rB  fwrapped_funr  i32_avalimplicit_argsexplicit_argsr  r   rS  sub_contextr   r  r  r
  r4  s                  @r:   	f_loweredzlower_fun.<locals>.f_loweredo  si   L%L%L%L%LA,q&))K+.+<+Dz%''' (  
  			$ T  ,,,3"))++3d33IIIc6G,H,HIII#B(9(9::"E*+c#2C.D.DD/ / / /!$/ / / k+/O/O/O/OPP5kBBq&&!8clSSq&" 
	"(003=0II(!
ucncm
V

-#- - - +- -kc6 
   =                 s   	E6GGGr   r  rr   )r
  r4  r$  s   `` r:   r<  r<  j  s0    
$ $ $ $ $ $ $L 
r<   c           
         |j         s`||cxu rYn nV||j        t          |          f}	 | j        |         }nF# t          $ r# t          | ||||||          }|| j        |<   Y nw xY wt          | ||||||          }|S )Nr  r  )rS  r  r   r  r   r  )	r   fn_name
call_jaxprr   rH  r  r  r  rg  s	            r:   _lower_jaxpr_to_fun_cachedr*    s    		 #yL@@@@@@@@J$eGnn
5C4.s3gg 4 4 4"
w
GZ9#% % %g -4c$S)))	4 !Wj':!# # #G 
.s   8 *A%$A%inner_backendr  c                    | d S |^}}|rt          d          | |k    r,|t          j        |           vrt          d| d|  d          d S d S )Nz>Multi-platform lowering when a backend= parameter is specifiedz Outer-jit backend specification z5 must match explicit inner-jit backend specification rn  )r  r  r  r   )r+  r  outer_backendmore_lowering_platformss       r:   check_backend_matchesr/    s     
F,>)-) J
HJ J J}$$23MBBBB
	<= 	< 	<+8	< 	< 	<= = = %$BBr<   r'  c          	     &   ~t          |t          j                  rt          j        |          }t          ||j                   t                                                    }t          t          |          }t                      gt          |          z  |z   }t          |          }t          || ||||	|
          j        j        }fd|D             }g |||R }t#          j        |t&          j                            |          t-          |                    }t/          |j        |          }t3          j        |t          |          g          \  }}                    t9          t;          ||                              }||fS )Nr'  c                :    g | ]}                     |          S rr   r  )r   r  r  s     r:   r   z!call_lowering.<locals>.<listcomp>  s%    2223IMM#222r<   )r7   r   Jaxprr  close_jaxprr/  r  r  r   r   r}   r  rV   r  r*  r/  r  rE  CallOpr$   FlatSymbolRefAttrrA   r  r'  r  r   rQ  r  r  r>  )r(  rH  r)  r  r   r  r  r  r  r  r  r   r   rZ  re  symbol_namer  callr  r   s          `            r:   call_loweringr8    s   
 
DJ'' ,
++J///""$$%%'_i00,,,#g,,.=,&|44*	7Jy! ! !!%e  3222'222&	*>	*F	*T	*	*$		.155kBB.t44
6 
6$ -T\<HH)oi#g,,@@&)&&xGV0D0D'E'EFF*	J	r<   )r  c          
         t          || j        ||| j        | j        | j        | j        g|R d| j        i\  }}|                     |           |S )Nr  )r8  rH  r  r  r  r  r  r   )r   r/  r  r)  r   r  r  s          r:   core_call_loweringr:    sn    #
CNJ1C	lCM3=)37) ) ) ') ))V V	r<   	core_call)r/  c                @    | dk    rdS | dk    rdS t          d          )Nr  hostdevicedensezVInvalid compute type received. Current supported values are `device_host` and `device`r   )c_types    r:   map_compute_typerA    s9    }67 4 	5 	5 5r<   c                    | j         r| j         j        hdt          j                            t          | j         j                            i}t          j                            |          |j        j        d<   d S d S d S )N_xla_compute_typer  )	r  r  r$   ra  rA   rA  r.  r_  r`  )r   r  r  s      r:   r  r    sw    "s'8'E'Q$bm&7&7*788': ': ;I:<+//):T:TBL6777 #"'Q'Qr<   c                  t          j        j        t           j                  r~j        j                            j                  j        }fdt          t          |                    D             }g ||}t          j
                  }t          | |||          S t          j        j                  sGt          | j                  }t          j        t!                    ||t#          |                    }nXt%          d j        D                       s
J             t          j        t!                    |t#          |                    }t'          | |j                   |S )Nc                $    g | ]}j         |z   S rr   )r   )r   rg   r  s     r:   r   z$broadcast_in_dim.<locals>.<listcomp>  s     FFF1X]Q&FFFr<   broadcast_dimensionsc              3  Z   K   | ]&}|t           j                                        k    V  'd S r?   )r$   r   r   r   ry   s     r:   r   z#broadcast_in_dim.<locals>.<genexpr>  sI       * * bm44666 * * * * * *r<   )r   rm  r   rn  _rulesphysical_element_avalrR   r   rV   r   r   r   r   rR  r&   dynamic_broadcast_in_dimr}   r  r   r  r  )	r   r  r  rG  	elt_shapetrailing_dimsphysical_aval_outrR   r   s	     `      r:   r   r     s   
 x~v77 %;;  FFFFc)nn0E0EFFFMB1BMB*844R"9MO O O O !(.11 1*3??e(
(
#
#R

.
/
/ cc  * *".* * * * * 4 4+34 4 4 
(
#
#R
.
/
/1 1c sCI...Jr<   ops	ops_avalsr  	out_shapec                   g }t          ||          D ]\  }}|j        }t          j        ||          r|                    |           7t          |          t          |          k    sJ ||f            t          t          t          |          t          |          z
  t          |                              }|                    t          | |t          j	        ||j
                  |                     |S )z)Broadcasts multiple ops to the out_shape.rF  )r>  rR   r   definitely_equal_shaper1  rV   r  r   r   r~   r   )	r   rP  rQ  rR  r   r  op_avalop_aval_shaperG  s	            r:   multi_broadcast_in_dimrW    s   
 	#i(( 	N 	Nkb'MM"=)<< N	jjnnnn3y>>111M93M111!%I]9K9K(KSQZ^^"\"\]]	jj!#r"&"29gm"L"L7KM M M N N N N 
*r<   c                   t          j        |          }t          j        |j                  s8t	          | |j                  }t          j        t          |          ||          S t          j        t          |          |          S r?   )	r   r   r   rR   rR  r&   dynamic_reshaper}   rx   )r   r  r  rR   s       r:   rx   rx   "  sx    ))(			/	/ 6(hn==E!!2u   ;x00"555r<   c                  t          j        |j        t           j                  r|j        j                            |j                  j        }dgt          |          z  }dgt          |          z  }g ||R }g ||R }g ||R }t          j	        |          }	t          | ||	|||          S t          d |||fD                       rUt          | |          }t          | |          }t          | |          }t          j        t          |          ||||          S t          j        |t#          |          t#          |          t#          |                    S )Nr   rN   )start_indiceslimit_indicesr   c              3  @   K   | ]}t          j        |           V  d S r?   r   r   r  s     r:   r   zslice_op.<locals>.<genexpr>:  s0      
Z
ZQt%a(((
Z
Z
Z
Z
Z
Zr<   )r   rm  r   rn  rJ  rK  rR   rV   r   r   slice_opr  rR  r&   real_dynamic_slicer}   r   r  )
r   r2   r  r[  r\  r   rM  trailing_zerostrailing_onesrO  s
             r:   r_  r_  ,  s   x~v77 1%;;  S3y>>)NS3y>>)M5m5n55M0m0i00M((-((G*844C-]"/B B B B 
Z
Z=-QX2Y
Z
Z
ZZZ 123FFm23FFm,S'::g#!!	=-2 2 2 Yq&}55&}55&w//1 1 1r<   c                  | j         d         }t          j        |j        t          j                  r|j        j                            |j                  j        }| j         dd          }t          j        |r|d         j        nd          }t          t          j        d|                    gt          |          z  }g ||R }t          j        |          }t          j        |          }|j        }	t          j        |	          st!          | |	          }	t#          j        t'          dgt          |          z            t'          |          t#          j        t!          | |j                  |	                    }
t#          j        t-          |          ||
t#          j        |
|	          t'          dgt          |          z                      S t#          j        ||t3          |	                    S )Nr   rN   rD   )r  r   rm  r   rn  rJ  rK  rR   r>  ru   rB   rT   rV   r   r   r   rR  r&   clampr   subtractr`  r}   adddynamic_slicer  )r   r  r2   r[  x_avalrM  index_avalsr   ra  slice_sizesclamped_starts              r:   rg  rg  G  s   <?&x~v77 	(%;;  ,qrr"K% +8A: :E!"(1e"4"4556YGN5m5n55M!(++H''F+			,	, M /sK@@KIA3]+++,,=!!	l$S&,77  M !!!1{++aS3}---..	   Q{/K/KLLLr<   c                  t          j        |j        t           j                  r|j        j                            |j                  j        }| j        dd          }t          j        |r|d         j        nd          }t          t          j        d|                    gt          |          z  }g ||R }t          j        |          }	t          | |	|||          S t!          j        |||          S )Nr   r   rD   r   )r[  )r   rm  r   rn  rJ  rK  rR   r  r>  ru   rB   rT   rV   r   r   dynamic_update_slicer&   )
r   r  r2   r,  r[  rM  ri  r   r   rO  s
             r:   rm  rm  j  s    x~v77 >%;;  ,qrr"K% +8A: :E!5111223c)nnDE,m,e,,M*844%66.;= = = = #Av}===r<   c           	     d   t          d |||fD                       r?t          j        ||t          |          t          |          t          |                    S t	          | |          }t	          | |          }t	          | |          }t          j        t          |          |||||          S )Nc              3  >   K   | ]}t          j        |          V  d S r?   r^  r  s     r:   r   zpad.<locals>.<genexpr>~  sE       	N 	Nq		"	" 	N 	N 	N 	N 	N 	Nr<   )r   r&   padr  rR  dynamic_padr}   )r   r  r2   padding_valuepadding_lowpadding_highpadding_interiors          r:   rp  rp  {  s     	 	N 	N[-9;K-M 	N 	N 	N N N G71m";//"<00"#3446 6 6
 /sK@@K/\BBL3C9IJJ?!!	=+|5EG G Gr<   	dimensionc                  t          j        |j                  sEt          | |j                  }t	          j        t          |          |t          |                    S t	          j        t          |          t          |                    S r?   )	r   r   rR   rR  r&   dynamic_iotar}   rl   iota)r   r  rv  rR   s       r:   ry  ry    s{    			/	/ D(hn==E!!   8OH--x	/B/BCCCr<   c                    t          t          j        |t          j        |j                                      }t          | ||d          S )zAReturns an IR constant shaped full of `value` shaped like `aval`.rr   rF  )ru   rB   rT   r   r>  r   r   )r   r  r   zeros       r:   full_like_avalr|    s?    	RXeV%>tz%J%JKK	L	L$	#tT	C	C	CCr<   c                    t          | j        d         xt          j                  rAt	          j        j        t          j                  r t          fd          | ||          S t          j
        ||          gS )Nr   c                R    j         j                            j         | |          gS r?   )r   rJ  rf  )r2   r&  r\   s     r:   r   z&add_jaxvals_lowering.<locals>.<lambda>  s"    17>#5#5agq!#D#D"E r<   )r7   r  r   r~   r   rm  r   rn  r<  r&   rf  )r   r2   r&  r\   s      @r:   add_jaxvals_loweringr    sv    cl1o%t'788 R11RF9EEEEFFsAqQQQ
'!Q--r<   c                    |gS r?   rr   )r   r2   s     r:   r   r     s    1# r<   	directioncomparison_typec                6   |Ft          j        | j                  j        }t	          |t           j                  r|j        rdnd}nd}t          j        | |t          j	        
                    |          t          j        
                    |                    S )zCreates CompareOp.NUNSIGNEDSIGNEDFLOAT)compare_type)r$   r   rQ   r   r7   rZ   is_unsignedr&   compareComparisonDirectionAttrrA   ComparisonTypeAttr)r2   r&  r  r  	elem_types        r:   compare_hlor    s    #AF++8I)R^,,  &/&;I

ooo		!%%i00)--o>>	
@ 
@ 
@ @r<   c                   t          j        |j                  }t           j                            |j                  rt          j        |          }t          j        |          }t          ||dd          }t          |||d          }t          t          j	        |          t          j	        |          |d          }	t          j
        ||	|          }
t          j
        |
||          S  | ||          S )z@Min/max that compares complex values lexicographically as pairs.EQr  )r$   r   rQ   r   r7   r   r&   realr  imagselect)r  cmpr2   r&  tensor_typerxryreal_eqreal_cmpimag_cmpwhichs              r:   _minmax_hlor    s    #AF+++^{788 		!B	!B"b$00G2r300H38A;;S'BBHJw(33E:eQ"""2a88Or<   LTGTc                   t          j        |j        t           j                  s|j        t	          j        t          j                  k    rqt          j        |j        t          j                  rd}n)t          j        |j        t          j                  rd}nd}t          |t          | d|          d|          }t          j        t          |          |          S )zVariant of convert that has HLO semantics.

  In particular, treat casts to boolean as x != 0, rather than truncating
  integer values (b/209440332).r  r  r  r   NE)r   rm  r   rn  rB   rU   inexactsignedintegerr  r|  r&   rw   r}   )r   r2   aval_inr  r  s        r:   convert_hlor    s    
 
HNFO
<
< Ln****
33  ll		7="*:	;	;  lllA~c1g66lKKA	_X..	2	22r<   )xc.OpSharding | sharding.SdyArrayShardingr  set[int] | Noner  allow_shardy_loweringc           	     @   t           j        j        r4|r2t          j                            | |j                              j        S |r3dd                    d t          |          D                       z   dz   }nd}t          |          }	t          |	t          j                  s
J |	            t          j        |          j        }
t          j        |
          rd }nt%          ||
          g}t'          | |	g|g|d||          }t)          ||           |j        S )Nzunspecified_dims=[,c                ,    g | ]}t          |          S rr   r  )r   rg   s     r:   r   z&_wrap_with_spmd_op.<locals>.<listcomp>  s    222AQ222r<   ]r,   rN   )r  r  backend_configr  result_shapesr  )r   rq  r  r#   r  ShardingConstraintOpbuildr  r  r  r}   r7   r$   r   r   r   rR   r   rR  r  set_sharding)r/  r   r2   r  r   r  r  r  r  r  rR  r  r  s                r:   _wrap_with_spmd_opr    sA    "( I-B I<,,Q0@0@AAHH  )CHH22 011222-4 -4 469:NN N))+	K	)	)66;666 **0)	I&& CMM1#yAABM4{mqc"0 !.#2	4 4 4"
 r8	r<   r    )r  SPMDFullToShardShapeSPMDShardToFullShapec                    t           j        j        rt          |          | j        d<   d S t          |          | j        d<   d S )Nr*  r+  )r   rq  r  rJ  r`  )r  r   s     r:   r  r  	  sE    "( A$5h$?$?BM.!!!%6x%@%@BM/"""r<   ir.Attributec                \   t           j        j        r | j                    S t	          | j                  dk    r,t          j                             | j	                              S t          j                            t          t          j                            |                               S )Nd   )r   rq  r  r  rV   tile_assignment_devicesr$   ra  rA   SerializeToStringreprr  r  
from_protor   s    r:   rJ  rJ  	  s     "( 	J8>
 8+,,s22]9x9;;<<<]tBN$=$=h$G$GHHIIIr<   r   r  c                   t          |          }t          |t          j                  s
J |            t	          j        |          j        }t	          j        |          rd }nt          | |          g}t          d|g|gd|t          t          |j                            g|j        d d d         g          }|j        S )NLayoutConstraintrN   )r  r  r  r  operand_layoutsr  )r}   r7   r$   r   r   r   rR   r   rR  r  r  r   r   major_to_minorr  )	r   r2   r  rt  r  r  rR  r  r  s	            r:   wrap_with_layout_opr  	  s    
  ))+	K	)	)66;666 **0)	I&& CMM1#yAABM%[MQC !.$(w|)<)<$=$=#> $*#82#>"?A A A" 
r<   c                F     t          j                    fd            }|S )a  Decorator that causes the contents of a lowering rule to be reused.

  The lowering will be emitted out-of-line in a separate function, together with
  a call to that function. If the same primitive is called with the same shapes
  and parameters, a new call to the original function will be added, without
  emitting a new function. We allow for different lowering for the same
  primitive for different platforms in the same module.
  c                   | j         J | j         t          | j                  t          | j                  t          |                                          f}	 | j        j                            |          }n# t          $ r  | g|R i |cY S w xY w|*t          t          fi ||           }|| j        j        |<   t          t          | j                  }t          | j                  |z   }t          |          }t          j        |t"          j                            |j        j                  t+          |                    }t-          |j        |          S r?   )r@  r   r  r  r  r  r  rA   r   r  r   r   r}   r  r  rE  r4  r$   r5  r/  r  r  r'  r  )	r   r   rB  r  r%   rZ  re  r7  r  s	           r:   cached_loweringz'cache_lowering.<locals>.cached_lowering:	  sd   =$$$cmcm 4 4  "C%:>>sCCdd % % % Qs$T$$$V$$$$$	%
 |'(<(<V(<(<cBBd;?c3C866L#$$t+D(660!377	HH0668 8D *$,EEEs   A; ;BB)r  wraps)r  r  s   ` r:   cache_loweringr  1	  s@     ?1F F F F F. 
r<   xla_computationxc.XlaComputationc                    t           j        j                            |           }t          j                            |          S r?   )r  _xlamlirxla_computation_to_mlir_moduler$   r  parse)r  
module_strs     r:   r  r  U	  s-    w|::?KK*		$	$$r<   
dst_modulerV  
src_module
dst_symtabr  c                x   | j         |j         k    sJ t          j        |j                  }|pt          j        | j                  }t	                      }i }|j        j        D ]f}|j        j        }||v p|dk    }	|	rL|dk    r|n|}
|
}d}||v s||v s||v r|
 d| }|dz  }||v ||v ||v |||<   |	                    |           gt          j
                            d          }|j        j        D ]@}|j        j        |v r&|                    |||j        j                            ||j        d<   A|                                D ]+\  }}|j        j        D ]}|                    |||           ,|j        j        D ]1}| j                            |           |                    |           2|d         S )a~  
  Args:
    dst_module: the module into which the contents of src_module should be
      moved. Nothing in dst_module will be renamed.
    sym_name: the desired name for the "main" function of src_module after
      merging. This is a hint: the true name may be different because of symbol
      uniquification, and the true name is returned by this function.
    src_module: the module whose contents are to be alpha-renamed, set to
      private visibility, and merged into dst_module. src_module must contain
      exactly one symbol named "main".
    dst_symtab: the symbol table of `dst_module`

      Functions in src_module will be renamed such that they do not collide with
      functions in dst_module.

      This function mutates `src_module`. On return, `src_module` is left in an
      undefined state.

  Returns:
    the name of src_module's main() function, after renaming.
  r  r   r   rN   r  r  )r  r$   r  r_  r  r  
operationsr/  r  rf  ra  rA   set_symbol_namer`  r  replace_all_symbol_usesr1  rG  )r  rV  r  r  
src_symtab
used_names	renamingsr  r/  should_rename	base_namenew_namerg   r  old_names                  r:   merge_mlir_modulesr  Z	  s5   2 
	z1	1	1	1	1~j233*AR^J,@AA*uu* )O&  b7=DJ&8$&.M "fnn(($ih
a
 ##x:'='=##%%!%%	Q ##x:'='=## !ionnX Mi(('O& . .b	w}	!!  Yrw}%=>>>&-BM"## &OO-- A Aho( A A((8R@@@@A O&  bO2b	6	r<   c                .     t           d fd            }|S )Nr   r  c                   | j         }|j        }t          |t          j                  r|j        }n|j        }t          d | j        | j	        z   D                       rt          d| j         d          t          |j                  dk    rt          d          t          j        |j        d         || j        | j	        fi |}t!          |          }t#          |j        dj         ||j                  }t+          t,          | j	                  }	t/          |	          }
j        rt2          j                            |
          n|
d         }t9          j        |gt2          j                            |          tA          |                    j!        j        sgS fd	tE          t          |
                    D             }tG          ||	          S )
Nc              3  j   K   | ].}t          |d           ot          j        |j                   V  /dS )rR   N)r   r   r   rR   r  s     r:   r   z:xla_fallback_lowering.<locals>.fallback.<locals>.<genexpr>	  s`       X X34 1g /%ag...X X X X X Xr<   zAShape polymorphism for xla_fallback_lowering is not implemented (z); b/261682623rN   z=fallback lowering not implemented for multi-platform loweringr   xla_fallback_)r  c                T    g | ]$}t          j        t          |                    %S rr   r&   get_tuple_elementrh   )r   rg   r7  s     r:   r   z;xla_fallback_lowering.<locals>.fallback.<locals>.<listcomp>	  s<     < < < )$<< < < <r<   )$r  r  r7   r   rp  unsafe_axis_envr  r  r  r  r  r@  rV   r  r   primitive_subcomputationr  r  rO  r/  r  r   r}   r  r4  r$   	TupleType	get_tuplerE  r4  r5  rA   r  r  r   r'  )r   r   rB  
module_ctxrZ  r  r  
xla_modulecallee_namerZ  re  output_typeflat_resultsr7  r  s                @r:   fallbackz'xla_fallback_lowering.<locals>.fallback	  s0   #J&H(N:;; %)hh$h
 X X9<9UX X X X X m
kcm
k
k
km m m :  1$$GI I I2Q4! !! !O 0@@J$64966
*, , ,K 66L(66+F2<))*;<<<1B11E  }!377DD0668 88> 	   Vm< < < <"3'8#9#9::< < <L *,EEEr<   r%  )r  )r  r  s   ` r:   r  r  	  s9    $F $F $F $F $F >$FJ 
/r<      c                4    t          d | D                       S )Nc              3  "   K   | ]
}|d k    V  dS r   rr   rI  s     r:   r   z!is_empty_shape.<locals>.<genexpr>	  s&      Q!Vr<   )r  rs  s    r:   is_empty_shaper  	  s    	Q		r<   channeltokenhlo.TokenTypeoperandxc.OpSharding | Nonec          
        t           j                            | t                    }t          j        |g||t
          j                            d                    }t
          j                            t          t
          j	                            t          |                    t
          j	                            t          |                                        |j        d<   |t          ||           |j        S NT)is_host_transfer)_xla_host_transfer_handler_name_xla_host_transfer_rendezvousr  )r&   ChannelHandlerA   SEND_TO_HOST_TYPESendOpr$   rI  r.  dictra  r  r`  r  r  )r  r  r  r   r/  r   channel_handlesend_ops           r:   send_to_hostr  	  s     $((2CDD.Jy%)+)>)>@ @ @'35;??
*,-*;*;CII*F*F(*(9(9#d))(D(DF F F4G 4G'/0 (###	r<   out_avaltuple[ir.Value, ir.Value]c          
     `   t           j                            | t                    }t          j        t          |          t           j                                        g||t          j                            d                    }t          j	                            t          t          j                            t          |                    t          j                            t          |                                        |j        d<   |t          ||           |j        \  }}||fS r  )r&   r  rA   RECV_FROM_HOST_TYPERecvOpr}   r   r$   rI  r.  r  ra  r  r`  r  r  )r  r  r  r/  r   r   recv_opr  s           r:   receive_from_hostr	  	  s    $((2EFF.J11))++-.3^)+)>)>@ @ @' 46;??
*,-*;*;CII*F*F(*(9(9#d))(D(DF F F4G 4G'/0 (###/-&%	r<   r  r  rU  r  operand_avalsSequence[core.ShapedArray]operand_shapesSequence[xc.Shape]result_avalsr  tuple[Sequence[ir.Value], Any]c	          	        |pt          j                    }|}
g }|s|
fd}
|j                                        }t	          j        dt          j                  }t          t          j	        dt          j                            }g |t          j        |          d         }t          |||||j        |	          }|                    |           n^t          ||          D ]M\  }}|j                                        }t          |||||j        |	          }|                    |           Ng }g }|D ]}}|j                                        }t!          |t          j                  sJ t#          ||||j        |	          \  }}|                    |           |                    |           ~|                     |
||||t&          j                  }|j                            |           ||fS )Nc                     ~              S r?   rr   )r   callback_without_argss    r:   _wrapped_callbackz4_emit_tpu_python_callback.<locals>._wrapped_callback
  s    
""$$$r<   rq   rN   r   r   )r&   r  r  r  r   r~   rB   float32ru   r   r   aval_to_xla_shapesr  r   r1  r>  r7   r	  ,make_python_callback_from_host_send_and_recvr   dumpsr  )r  r   callbackr  r  r
  r  r  r  r   r  send_channelssend_channeldummy_send_avaldummy_send_valr  operand_avalr  recv_channelsoutputsresult_avalr   ifrt_callbackr  s                          @r:   _emit_tpu_python_callbackr"  	  sF    
%3#%%%-	 $
 .% % % % % %1133L&tRZ88O !RZ!8!899NB~ B,_==a@BNuno!*X? ? ?E&&&&!$X}!=!= $ $"..00g7E7L#,xA A Ae7####-'! " "k ,,..Gk4#344444"7E;#+#4xI I IJE3NN3!!!!FF[&( (- &&}555	%r<   minor_to_majorSequence[int] | Nonec                    | t          j        dd          }nt          j        | d          }t          j                            |t          j                                                  S )Nr   rD   r   rQ   )rB   r   rT   r$   r@   rA   	IndexType)r#  rt  s     r:   _layout_to_mlir_layoutr)  -
  s\    Xd'***FFXnG444F		 	$	$V",2B2B2D2D	$	E	EEr<   c                D    t          j        |           g}d |D             S )Nc           	     Z    g | ](}t          t          |j        d z
  dd                    )S )rN   r  )r  r   r   r   r   s     r:   r   z,_aval_to_default_layouts.<locals>.<listcomp>8
  s2    	?	?	?$uTY]B++
,
,	?	?	?r<   )r   r   )r   avalss     r:   _aval_to_default_layoutsr.  5
  s)    d##
$%	?	?	?	?	??r<   )r   r  r  r  %Sequence[Sequence[int] | None] | None#tuple[Sequence[IrValues], Any, Any]c                  t          | j        j                  dk    rt          d          | j        j        d         dvrt	          d d          | j        j        }
t          j        d D                       }t          j        d |D                       }|'t          j        t          t          |                    }t          t          |          }|	't          j        t          t                              }	t          t          |	          }fd
}dk    rmt          j        d t          |          D                       \  }}t          |
| ||||||||
  
        \  }}t          |          fdD             }||d	fS t!          d D                       }|r|fd}t#          j        t&          j                  d         g|}t#          j        t&          j                  d         g|}|g|}t+                      g|}t          d	          g|}t          d	          g|}|
                    |||          \  }}| j                            |           t1          |          }|g|}|t          g           g|}t2          j                            |          }dv rdnd}t9          j        |g|t2          j                            |          t2          j                             |          tC          d          t2          j"                            g           t2          j                            tG          |                    |d	nt2          j"                            |          |d	nt2          j"                            |          	  	        |tI          |           fdtK          t          |                    D             }|r|^}}|||fS )z9Emits MLIR that calls back to a provided Python function.rN   z+multi-platform lowering for python_callbackr   >   rV  rY  rW  rX  z&`EmitPythonCallback` not supported on z	 backend.c                6    g | ]}t          j        |          S rr   r   r  )r   r   s     r:   r   z(emit_python_callback.<locals>.<listcomp>Q
  s#    KKK{sk**KKKr<   c                6    g | ]}t          j        |          S rr   r3  )r   rU  s     r:   r   z(emit_python_callback.<locals>.<listcomp>S
  s#    DDD7sg&&DDDr<   Nc            	     p    |  }t          |          t                    k    r=t          d                    t                    t          |                              t          d |D                       }t	          t          |                    D ]l\  }\  }}|j        |j        k    r"t          d| d|j         d|j                   |j        |j        k    r"t          d| d|j         d|j                   mdk    r)t          d t          |          D                       }|S |S )	NzDMismatched number of outputs from callback. Expected: {}, Actual: {}c              3  b   K   | ]*}t          j        t          j        |                    V  +d S r?   )r   r>  rB   rC   r  s     r:   r   zBemit_python_callback.<locals>._wrapped_callback.<locals>.<genexpr>f
  s5      MMqS+BJqMM::MMMMMMr<   z)Incorrect output shape for return value #z: Expected: z
, Actual: z)Incorrect output dtype for return value #rY  c              3  H   K   | ]\  }}t          |j                  |V  d S r?   r  rR   )r   out_valr   s      r:   r   zBemit_python_callback.<locals>._wrapped_callback.<locals>.<genexpr>v
  sK       !4 !4"g{ 122!4
!4 !4 !4 !4 !4 !4r<   )rV   RuntimeErrorformatr   r'  r>  rR   r   )	r   r  rg   r9  r  non_empty_out_valsr  r  r  s	         r:   r  z/emit_python_callback.<locals>._wrapped_callback_
  s   xH
8}}L))))%%+VC,=,=s8}}%M%MO O O MMHMMMMMH"+C,,G,G"H"H D DGX	(.	(	(C C C!C C3:=C CD D 	D 
(.	(	(C C C!C C3:=C CD D 	D 
)
 5
 ! !4 !4&)(L&A&A!4 !4 !4 4 4  or<   rY  c                D    g | ]\  }}t          |j                  ||fS rr   r8  )r   r   rR   s      r:   r   z(emit_python_callback.<locals>.<listcomp>
  sG     C+ C+ C+D%dj))C+	uC+ C+ C+r<   r   c                    g | ]R}t          |j                  r-t          t          j        |j        |j                             nt                    SS )r   )r  rR   ru   rB   r   r   r]  )r   r   non_empty_outputs_iters     r:   r   z(emit_python_callback.<locals>.<listcomp>
  sj     ) ) )  ++,,	OBH[.k6GHHHIII267M2N2N) ) )r<   c                ,    g | ]}t          |          S rr   r}   r,  s     r:   r   z(emit_python_callback.<locals>.<listcomp>
  s     "R"R"RT?4#8#8"R"R"Rr<   c                    | g | R S r?   rr   )r  r   callback_without_tokens     r:   r  z/emit_python_callback.<locals>._wrapped_callback
  s    4,,d3444r<   >   rW  rX  xla_python_gpu_callbackxla_python_cpu_callbackr   )call_target_namer  r  called_computationsr  r  r  c                T    g | ]$}t          j        t          |                    %S rr   r  )r   rg   r  s     r:   r   z(emit_python_callback.<locals>.<listcomp>
  s<       
 
FHQKK00  r<   )&rV   r  r  r  r   r  r   rH  r   r   r.  r)  unzip2r>  r"  r  r  r   r  r   ry  r  #get_emit_python_callback_descriptorr  ru   r$   r  r  r&   CustomCallOpra  rA   rI  rh   rL  r  r  r   ) r   r  r  r  r
  r  r  r   r  r  r  r  r  operand_mlir_layoutsresult_mlir_layoutsr  non_empty_result_avalsnon_empty_result_shapesnon_empty_outputsr  r  callback_descriptorr!  descriptor_operandcallback_operandsr  rF  r  rC  r?  r  r  s     `   `                      @@@@r:   emit_python_callbackrT  ;
  s    			%&&**
K
L
LL)!,(333
DDDDF F F&',KKlKKKM M-<DDmDDDF F. &$m446 6O3_EE%c*BL&Q&QRRN2NCC      > 6:k C+ C+|];;C+ C+ C+ 7, 7,33  9'%- 7	     u
 ""344) ) ) ) () ) )G E4!"R"R\"R"R"RSS,
 O.5 5 5 5 5 	t233A69GN 	t233A69FM !!HLL0<0L2488P;OP1$77N:MN112C2@2?A A %} &&}555"#677)5H5%2266N9MN&&|44+ $444 0/:S m}(()9::kooo661++,**2..]&&s+>'?'?@@$,\233#+\1225 5 5& """   S&&''  '  OEG	%	&&r<   closed_jaxprc                   t                               | j                  }|rt          d| j                   t	          || |g t          j                    dgt          | j        j	                  z  ||t                      	  	        }|j        S )zGHelper to generate pmap-style XLA computations for custom partitioners.r  F)r  r  rH  r  r  r  r  )r1   r  r   r   r  r   r  rV   r  r  r  rO  )rU  r/  r  r  r  r  lowering_results          r:   build_mlir_module_helperrX  
  s    
 *778LMM Q
O9MOO
P
PP)$%r!+--7S!3!:;;;9,..0 0 0/ 
	r<   rr   )	r  r  r  rG  r  operand_output_aliasesr  r  extra_attributesrF  r  Sequence[ir.Type]r  %str | bytes | dict[str, ir.Attribute]Sequence[ir.Value] | NonerG  r  rY  dict[int, int] | NoneSequence[Sequence[int]] | NonerZ  dict[str, ir.Attribute] | Noneir.Operationc                  t          |          }| t          j                            d          }nt	          |t
          t          f          r t          j                            |          }nVt	          |t                    r"t          j                            d          }d}nt          dt          |          z             t          t          j                            |           t          j	                            |          |t          |          t          j                            d |D                                 }|Bt          j                            fd|                                pdD                       |d	<   ||                    |           |t          j                            t          j        t          t#          t%          |          t%          |          t%          |          z                       t          j        
                    |d<   |	Ot%          |	          t%          |          k    sJ |	|f            t          |	          dgt%          |          z  z   }	t          |          t          |          z   }|	,t          j                            d |	D                       |d<   |
\|
J t%          |
          t%                    k    sJ |
f            t          j                            d |
D                       |d<   t(          j                            ||          }t	          |t                    r,t          j                            |          |j        j        d<   |S )a  Helper function for building an hlo.CustomCall.

  Args:
    call_target_name: the name of the custom call target
    result_types: the MLIR types of the results of the custom call
    operands: the MLIR IR values that are arguments to the custom call
    backend_config: an opaque string passed to the custom call kernel
    has_side_effect: if True, marks the custom call as effectful
    result_shapes: tensors that represent the result shapes, to be used when
      the results have dynamic shapes. If not-None, its length must match the
      number of the results.
    called_computations: the list of function names called by the custom call.
    api_version: the ABI contract version of the custom call
    operand_output_aliases: a dict mapping operand numbers to outputs they alias
    operand_layouts: a sequence of layouts (dimension orders) for each operand
    result_layouts: a sequence of layouts (dimension orders) for each result
    extra_attributes: additional IR attributes to apply to the custom_call.
  Nr,   rN   z,custom_call backend_config unexpected type: c                L    g | ]!}t           j                            |          "S rr   )r$   r5  rA   )r   r/  s     r:   r   zcustom_call.<locals>.<listcomp>  s)    
J
J
Jd2##D))
J
J
Jr<   )rF  r  r  r  rG  c                    g | ]=\  }}t           j                            t                    d k    r|gng |g           >S )rN   )output_tuple_indicesoperand_indexoperand_tuple_indices)r&   OutputOperandAliasrA   rV   )r   	input_idx
output_idxr  s      r:   r   zcustom_call.<locals>.<listcomp>  sj     	= 	= 	=  )Z 
   03</@/@1/D/D
||"! " !  	= 	= 	=r<   rr   output_operand_aliasesr   indices_of_shape_operandsr&  c           
         g | ]o}t           j                            t          j        t          j        |t          j                             t           j                                                  pS r   r'  r$   r@   rA   rB   
atleast_1drC   rD   r(  r   r"  s     r:   r   zcustom_call.<locals>.<listcomp>-  sq     6 6 6 *+ 	##M"*Qbh77788!!## 	$ 	% 	%6 6 6r<   r  c           
         g | ]o}t           j                            t          j        t          j        |t          j                             t           j                                                  pS rn  ro  rq  s     r:   r   zcustom_call.<locals>.<listcomp>6  sq     5 5 5 *+ 	##M"*Qbh77788!!## 	$ 	% 	%5 5 5r<   r  )r  r  r`  zmhlo.backend_config)r  r$   ra  rA   r7   r  r  r  r   rI  rh   rL  r  r,  r@   rB   rC   r   rV   rD   r&   rK  build_genericr.  r_  r`  )rF  r  r  r  r  r  rG  r  rY  r  r  rZ  backend_config_attrr`  r  s    `             r:   r  r  
  s   B (^^(-++B//.3,// [-++N;;.$'' [ -++B//KK
Cc.FYFYY
Z
ZZ}(()9::kooo66(;'',**
J
J6I
J
J
J   * '+-<+;+; 	= 	= 	= 	= %;$@$@$B$B$Hb	= 	= 	= 	, 	,J'( !&''' /1.E.I.I

4c(mmS]]S=O=O-OPPQQ	# 	# 	#/$ /$J*+ "!!S]]222_h4O222_--]9K9K0KKoH~~] 3 33H $&L$4$4 6 6 />6 6 6 % %J !
 %%%~#l"3"33336&333#%<#3#3 5 5 /=5 5 5 $ $J  %%lX1; & = ="%% 57[__6 6BL12	)r<   reducer_namereducer_body(Callable[[ir.Block], Sequence[ir.Value]]init_valuesinit_values_avalsr  c                   t          d |D                       }t          d ||||
g|	D                       rt          t          j        dt
          j                            d fd}t          j        t          t          ||	                    t          d                    }t          j                            ||z   |          }t          j                             j        j        j                  5  t)          j        ||          }ddd           n# 1 swxY w Y    j        j                            |           |                                }t          j        |          5  t          j         ||                     ddd           n# 1 swxY w Y   t5          d	t          t          t          |                    g ||t7           |          t7           |          t7           |
          t7           |          ||j        j        g
          }n!t          j        t          t          t          |                    ||t?          |          t?          |          t?          |
          t?          |          t          j                             t          j!        |	t
          j"                  tG          |	          dg                    } |j$        d         j%        j&        ||z    }t          j        |          5  t          j         ||                     ddd           n# 1 swxY w Y   |j'        S )z9Builds a ReduceWindowOp, with support for dynamic shapes.c                ,    g | ]}t          |          S rr   rA  r,  s     r:   r   z!reduce_window.<locals>.<listcomp>R  s     "W"W"WT?4#8#8"W"W"Wr<   c              3  @   K   | ]}t          j        |           V  d S r?   r^  r  s     r:   r   z reduce_window.<locals>.<genexpr>S  sL       	a 	a #A&&	& 	a 	a 	a 	a 	a 	ar<   )rN   r   	pad_lo_hi!tuple[core.DimSize, core.DimSize]c                N    t          |           }t          j        |          S r?   )rR  r&   rx   )r}  padsr   int2ds     r:   prep_one_padz#reduce_window.<locals>.prep_one_padX  s$    )#y99d[%%%r<   r   Nzstablehlo.dynamic_reduce_window)r  r  rG  r   r  )window_stridesbase_dilationswindow_dilationspadding)r}  r~  )(r  r  r}   r   r~   rB   rv   r&   r   r  r   rl   r$   rD  rA   r  at_block_beginr  rO  r  rE  rF  r  rG  rO  rU  r  rR  r/  r  ReduceWindowOpr  r@   rC   rD   rV   r  r  r1  r  )r   ru  rv  r  rx  ry  r  window_dimensionsr  r  base_dilationwindow_dilationscalar_typesr  	d_paddingreducer_typereducerr~  rwr  s   `                  @r:   reduce_windowr  D  s'    ""W"WEV"W"W"WXX, 	a 	a$o~}_W^_	a 	a 	a a a +) D,VRX>>??E& & & & & & & Sw%?%? @ @(1++NNI?&&\!<1 1L			)	)#*<*C*H	I	I @ @#L,??g@ @ @ @ @ @ @ @ @ @ @ @ @ @ @#**7333))++K		;	'	' - -	k,,{++,,,- - - - - - - - - - - - - - - 
'#C$C$CDD	$S*;<< 	%S.99 	%S-88	
 	%S/:: 	 $L./
 
 
BB 
	S),,--+)**&~66&}55(99'++BJw,I,I36w<<2C , E E
F 
F 
FB *bjm")L<,GIG		7	#	# ) )	k,,w''((() ) ) ) ) ) ) ) ) ) ) ) ) ) )	s6   =DD#&D#1FF"FMM	Mc                \   	 t           j                            t          |           dd          }n3# t          $ r&}t          dt          | d          z             |d}~ww xY wt                      }|5  t          j	        
                    |          cddd           S # 1 swxY w Y   dS )a  Refines the polymorphic shapes inside a module.

  Given a module with static input shapes, but using dynamic shapes due to
  shape polymorphism, runs shape refinement to resolve all the dynamic shapes.
  Then verifies that there are no more dynamic shapes in the module.
  T)enable_shape_assertionsvalidate_static_shapeszError refining shapes.  before_refine_polymorphic_shapesN)r"   r  refine_polymorphic_shapesr  	Exceptionr   rp  r  r$   r  r  )rO  refined_module_strr  r  s       r:   r  r    s   P&+EE  $! F # # 
 P P P
!F$FGG	HI INOPP
 ' / /9??-../ / / / / / / / / / / / / / / / / /s&   /2 
A"!AA"5B!!B%(B%)r2   r3   r4   r5   )r4   r=   )rF   rH   r4   rI   )rF   rH   r4   r^   )rm   rn   r4   ro   )r   r   r4   r   )r   r   r4   r   )r   r   r4   r   )r   r   r4   r   )r   r   r4   r   )r   rQ   r   r   )r   rQ   r4   r   r   )r2   r   r4   r3   )r   r   r4   r3   )r
  r  r  r  r4   r  )r   r  r
  r  r4   r5   )r   r  r  r  r4   r   )
r   r  r@  rA  rB  rC  rD  rE  r4   r   )rO  rP  rQ  r  r4   rR  )rO  rP  rQ  r  r4   r  )rq  r  r4   r  r?   )rO  rP  r4   r  )rO  rP  r4   r  )r4   r  )r  rA  r  r  r  rR  )rF   r  r4   r  )rF   r
  r4   r  )rF   r  r  r  r4   r  )r2   r   r4   rt   )rF   r  r   r!  r4   r"  )r   r   r   r(  r4   r   )r   r  rR   r/  r4   r0  )r   r  rR   r/  r4   r@  )r   r  rR   r/  r4   r   )rZ  r[  )r   r  r   r   r   r(  r4   rh  )rt  ru  r   r   r4   rR  )rq  r(  r4   rR  ).ri  r  r  r  r  r  r  r  r  r  r  r  rH  r  r  rH   r}  r  r~  r  r  r  r  r  r  r  r  r  r  r  r  rt   r  rt   r  r5   r  r  r  r  r  r  r  r  r4   rT  )*r   r  r/  r  r  r  r   r  rH  r  r  r5   r}  r  r~  r  r  r  r  r5   r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r4   r  )r2   r  r]  r  r  r   r4   r  )r   r  r4   r  )r   r  r4   r  )r   r  r  r  rH  r  r  r  rS  r  r   r3   r  r  r4   r  )r  r  r4   r  )r   r  r  r  r  r  r  r  r   r  r  r  r4   r  )r4   r"  )T)r
  r   r4  r5   r4   r   )NN)r+  rR  r  r  )r   r  r  r  r%  )r   r  r  r   r4   r  )
r   r  rP  r  rQ  r  rR  r/  r4   r  )r   r  r4   r  )r   r  rv  rt   )r   r  r   r   r4   r  )r  r  r  rR  )NFF)r/  r  r   r  r2   r  r  r   r   r  r  r  r  r5   r  r5   )r   r  )r   r  r4   r  )
r   r  r2   r  r  r   rt  r   r  r   )r  r  r4   rP  )
r  rP  rV  r  r  rP  r  r  r4   r  )r  rA  )rq  r/  r4   r5   )r  rt   r  r  r  r	   r   r   r/  r  r   r  r4   r  )r  rt   r  r  r  r   r/  r  r   r  r4   r  )r  r  r   r  r  rU  r  r  r
  r  r  r  r  r  r  r  r   r  r4   r  )r#  r$  )r   r  r  rU  r  r  r
  r  r  r  r  r5   r   r  r  r/  r  r/  r4   r0  )rU  r  r/  r  r  r  r  r  r  r  r4   rP  )rF  r  r  r[  r  r  r  r\  r  r5   r  r]  rG  r  r  rt   rY  r^  r  r_  r  r_  rZ  r`  r4   ra  )r   r  ru  r  rv  rw  r  r  rx  r  ry  r  r  r  )rO  rP  r4   rP  (   
__future__r   r  r  collections.abcr   r   r   r   r  r  r   rx  r  r7  r[  r  r  typingr	   r
   r   r   r   r   r  numpyrB   jax._srcr   r   r   r   r   r  r   r  r   r   r   r   r   r   r   r  jax._src.interpretersr   r  r   jax._src.layoutr   r   jax._src.shardingr    rj  jax._src.libr!   r  r"   jax._src.lib.mlirr#   r$   jax._src.lib.mlir.dialectsr%   rE  r&   r'   jax._src.state.typesr(   safe_mapr   
unsafe_mapsafe_zipr>  
unsafe_zipTypeVarr)   r   MYPYstring_flaggetenvrZ  r{  r1   r  r   r3   r;   rG   DenseI64ArrayAttrrA   r  r]   ra   rh   rl   r   r   r   r   r   r   r   rZ   r[   rU   int4int8int16rv   rD   uint4get_unsignedrX   uint16uint32uint64float8_e4m3b11fnuzFloat8E4M3B11FNUZTypefloat8_e4m3fnFloat8E4M3FNTypefloat8_e4m3fnuzFloat8E4M3FNUZTypefloat8_e5m2Float8E5M2Typefloat8_e5m2fnuzFloat8E5M2FNUZTypebfloat16BF16Typefloat16F16Typer  r   float64r   	complex64
complex128r   int2uint2r   r   r   r   r}   r~   ConcreteArrayr*  r+  r   r   r   r   r   ru   r   r   maMaskedArrayr  ndarraylonglong_scalar_typer  python_scalar_dtypesr  ptyper	  r  r  r  r?  rN  DialectRegistryr  register_dialectsr  r^  rl  rp  rb  rf  r  r  rp  ReplicaAxisContextShardingContextr  r  	dataclassr  r  r  r  r  r  r  r  r  r  r  objectr  r	  r  r  r  r'  compiler  r.  r3  rK  rP  rR  rT  r  rg  rk  rz  r|  r  r  r   r  r  r  r  rB  r9  r  rT  r  r  r	  r<  r*  r/  r8  r:  call_pclosed_call_prA  r  r   rW  rx   r_  rg  rm  rp  ry  r|  r  add_jaxvals_pstop_gradient_pr  r  minimummin_hlomaximummax_hlor  r  r4  wrap_with_full_to_shard_opwrap_with_shard_to_full_opr  rJ  r  r  r  r  r  DEVICE_TO_DEVICE_TYPEr  r  r  r  r	  r"  r)  r.  rT  rX  r  r  r  rr   r<   r:   <module>r     s    # " " " " " "         B B B B B B B B B B B B               				      				 				   F F F F F F F F F F F F F F                              + + + + + + & & & & & &                         # # # # # # % % % % % %       % % % % % % 4 4 4 4 4 4 % % % % % % 9 9 9 9 9 9 9 9 3 3 3 3 3 3 ) ) ) ) ) ) & & & & & & & & & & & &             ; ; ; ; ; ; * * * * * * 3 3 3 3 3 3 , , , , , ,-Z-ZFN3 $&$ibi 0"55
=> > > $66#5#")-v66#$$ $$ $$   0;/L  L L L L
 53//0- - - -
? ? ? ? &*? ? ? ?' ' ' ' O N N N N N, , , ,(   rw|,,
-7 7 7 7=
"(6=772>#>BB=
"(28ggbn91=== "(6;!<a@@= "(27WWR^8!<<	=
 "(28ggbn92>>= "(28ggbn92>>= "(28ggbn92>>= "(6<''"."=qAA= "(28ggbn91=== "(29wwr~:B??= "(29wwr~:B??= "(29wwr~:B??= "(6$%%r'?'C= "(6  ""5"9= "(6!""B$9$==  "(6 1 5!=" "(6!""B$9$=#=$ "(6?R[_
"(2:

"(2:

"(2:

"(2<FF
"(2=GG/= =     6 
;		!	!	!-4Wn!1. .HBHV[))* /6gn!1/ /HBHV\**+
   K K K KF F F F
 NP  O O O OP P P P &5 ! "'6 # $'D'D # $&= " #5 5 5 5= = = = =h = = = 46  5 5 5 5* * * *# # # #@ @ @ @&   W W W  "%+-K L L L& & & &B  "*&? @ @ @Wbh"(Xry")RYZRZ\2=Xr{FO	= E EL
 L*CDDDD5 5 5 /5577 K KLE5E77+A5#I#IJJJJ    $*&= > > >	 	 	 	   ' ' ' 'R   , 'B&((  =))*;<<< #9?$$ ! ! ! !F7 7 7 7& & & &          * "%"$       > d###4 4 4 4 4 4 4 $#4* 
! 
! 
! 
! 
! 
! 
! 
! ^B ^B ^B ^B ^B ^B ^B ^BB        D  ; ; ; ; ;X ; ; ; ; ,13
 3 3 3 3 K K K K6{6t<<  .2        &(( E E E    &((
 
 
 
1 1 1 1     RZ
++ 	7 	7 	7 	7D D D D$G G G GG G G G6 6 6 6
+ + + + +Z + + + :99 ) ) ) )= = = =,+ + + +   " .27;:>IMJN-1)-!%:>>B;?/\. \. \. \. \. \.~(0 (0 (0T 	]
*  *  *  *  *  *  *  * f -17;:>%)8<.2-1)-487;JNMQ>B+\ \ \ \ \ \~
   
. 
. 
. 
.# # # #LF7 F7 F7 F7R   F F F FP4 4 4 4+ + + + +^ =A   &= = = =* !t	     6 -1       $+ww'9LLL M M M  $$',4888: : :5 5 5U U U   <   $6 6 6 61 1 1 16!M !M !M !MF> > > >"G G G G"	D 	D 	D 	DD D D D
  
  '')= > > >  ')+=+= > > >@ @ @ @ @   '+s{D
1
1
'+s{D
1
13 3 3 3, <@/45:         F   2J6:< < < $W%79OPP $W%79OPP A A A AJ J J J   4! ! !H% % % % <@D D D D DN' ' ' 'T           
 37     " 8<     < &*6 6 6 6 6 6pF F F F@ @ @ &*=A<@G' G' G' G' G' G'R       * =?!/3)+486:597;m m m m m m`; ; ; ;|/ / / / / /r<   