
    ukiW                   :#   U d dl m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mZ d dlmZ d dlmZ d d	lmZ d d
lmZ d dlmZ  d dlm!Z! d dlm"Z" d dlm#Z# d dlm$Z$ d dlm%Z& d dlm'Z' d dlm(Z( d dlm)Z) d dlm*Z* d dlm+Z, d dl-m.Z/ d dl0m1Z1m2Z2 d dl3m4Z4 d dl3m5Z5 d dl3m6Z7 d dl8m9Z9m:Z:m;Z; d dl<m=Z>m?Z? d dl@mAZA d dlBmCZC d dlDmEZF d d lGmHZHmIZImJZJmKZKmLZL d d!lMmNZN d d"lOmPZP d d#lQmRZR d dlSZTe*j                  eVcZVZWe*j                  eYcZYZZ ej                  d$      Z\eZ]d%Z^ee:j                  e_e:j                  d&f   f   Z`dd'Zadd(Zbe:j                  j                  Zedd)Zfd* Zgd+ Zh	 	 dd,Zid- Zjee:j                  e_e:j                  d&f   f   Zldd.Zmi  eTj                  ej                         e
e:j                  j                  d/       eTj                  eTj                         e
e:j                  j                  d/       eTj                  ej                         e
e:j                  j                  d0       eTj                  eTj                         e
e:j                  j                  d1       eTj                  eTj                         e
e:j                  j                  d2       eTj                  eTj                         e
e:j                  j                  d3       eTj                  eTj                         e
e:j                  j                  d4       eTj                  ej                         e
e:j                  j                  d0       eTj                  eTj                         e
e:j                  j                  d1       eTj                  eTj                         e
e:j                  j                  d2       eTj                  eTj                         e
e:j                  j                  d3       eTj                  eTj                         e
e:j                  j                  d4       eTj                  ej                        e:j                  j                   eTj                  ej                         e:j                  j                   eTj                  ej                        e:j                  j                   eTj                  ej                        e:j
                  j                   eTj                  ej                        e:j                  j                   eTj                  ej                        e:j                  j                   eTj                  eTj                        e:j                  j                   eTj                  eTj                        e:j                  j                   eTj                  eTj                        e:j                  j                   eTj                  eTj                         d5  eTj                  eTj"                        d6  eTj                  ej$                         e
e:j                  j                  d7       eTj                  ej&                         e
e:j                  j                  d7       eTj                  ej(                        e:j*                  j                   eTj                  ej,                        e:j.                  j                   eTj                  ej0                        e:j2                  j                   eTj                  ej4                        e:j6                  j                  iZd8ed9<   dd:Zdd;Zdd<Zi Zd=ed><   dd?ZeeejF                  <   d@ eejH                  <   ddAZ G dB dCe      Zi ZdDedE<   ddFZddGZdddH	 	 	 	 	 	 	 ddIZddJZdK Z eeTjZ                  j\                  e       dL Z eej`                  e       	 	 	 	 ddMZ eeTjd                  e        ee#jf                  e       eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                   eTj"                  eTj                  eTjh                  ej                  fD ]  Z eee        ddNZejn                  D ]  Z eee        ddOZ eejt                  e       eege:jv                  f   Zi ZdPedQ<   ddRZddSZddTZddUZddVZ eeTjd                  e«        ee"j                  dW        eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                  eTj                   eTj"                  eTj                  eTjh                  ej                  fD ]  Z eee«        ddXZ eeTj                  eī        eeTj                  eī       dY Zej                  j                         D ]  \  ZZn ee e
een               eee:j                  j                          eee:j                  j                         ddZZ eee̫       dd[Z eeeΫ        ee_eΫ        ee:jv                  d\         ee:j                  d]        dd^Zdd`ZddbZ	 	 	 	 	 	 	 	 	 	 ddcZ e:j                         Z e5j                  eի        ej                         ZdddZddeZddfZdd dgZܐddhZ e:j                         Z G di dje:j                        ZddkZee(j                  e(j                  e(j                  f   Z G dl dm      Z ej                  dno       G dp dq             ZddrZej                   G ds d_             Z ej                  dno       G dt du             Z ej                  dno       G dv dw             Zej                   G dx da             Zej                   G dy dz             Ze^s G d{ d|e      ZneZ ej                  dno       G d} d~             Zi Zded<   ded<    ej                  eͫ      Z	 	 d	 	 	 	 	 	 	 ddZddZddZ e       ZddZd	dZ	 	 	 	 d
dZ ej                  d      ZddZ	 	 	 	 ddZ 	 	 	 	 ddZ	 	 	 	 ddZ	 	 	 	 	 	 ddZ	 	 	 	 ddZ G d de      Zg dZddZ	 	 	 	 	 	 	 	 ddZ	 	 	 	 ddZ	ddZ
d Zd Z G d de      ZddZddZd/Zdddddddd/d/dnddd	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddZd Ze:j                  Ze?j&                  j                  Ze?j*                  Z G d d      Zd%ddddndddddddddd	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddZ	 	 	 	 	 	 	 	 ddZddZ e       Zded<   	 	 	 	 	 	 	 	 	 	 	 	 	 	 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Z$d#d$dZ%	 d%	 	 	 	 	 d&dZ&	 	 d'dZ'	 	 d(	 	 	 	 	 d)dZ(dddd	 	 	 	 	 	 	 d*dZ)dd	 d+dZ* eejV                   e
e*d              eejX                   e
e*d      d%       d,dZ-d-dZ.d-dZ/	 	 d.dZ0	 	 	 	 	 	 	 	 d/dZ1d.dZ2	 	 d0dÄZ3	 	 d0dĄZ4	 	 d0dńZ5	 	 d0dƄZ6d1dǄZ7d2dȄZ8dɄ Z9 eejt                  e9        eejv                  dʄ        dd3d˄Z<d̄ Z= e
e=e?j|                  dͫ      Z? e
e=e?j                  dΫ      ZAd4dτZB	 	 	 d5	 	 	 	 	 	 	 	 	 	 	 	 	 d6dЄZC e
eCddnҫ      ZD e
eCdӫ      ZE e
eCdԫ      ZFddՄZGd7dքZH	 	 	 	 d8dׄZI	 	 	 	 	 	 	 	 d9d؄ZJ	 d	 	 	 	 	 	 	 d:dلZKd/ZLd7ZMdڐZN	 	 	 	 	 	 	 	 	 	 	 	 d;dۄZOdd%ddd7dddddޜ		 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d<d߄ZP	 	 	 	 	 	 	 	 	 	 	 	 	 	 d=dZQd>dZRy(?      )annotationsN)CallableIterableIteratorSequence)partial)Any
NamedTupleProtocolUnioncast)ad_util)api_util)config)core)dtypes)effects)frozen_dict)hashable_array)literals)
jaxpr_util)linear_util)path)sharding_impls)source_info_util)util)
xla_bridge)partial_eval)
AutoLayoutLayout)_jax)jax_mlir_ext)
xla_client)dialectsirpassmanager)funchlo)AxisType)PartitionSpec)Sharding)AUTONamedShardingSdyArraySdyArrayList"modify_sdy_sharding_wrt_axis_types)AbstractRef)	ArrayLike)foreachTF.c                8    t        | t        j                         S N)
isinstancer%   BlockArgumentxs    U/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/_src/interpreters/mlir.py_is_not_block_argumentr<   T   s    2++,	,,    c                    t         j                  j                  t        j                  | t        j
                              S r6   )r%   DenseIntElementsAttrgetnpasarrayint64)xss    r;   dense_int_elementsrE   W   s(    		 	 	$	$RZZBHH%=	>>r=   c                   t        j                  t        j                  | t         j                        d      }t        j
                  j                  |t        j                  j                  d      t        |       g      S )Nlittlebitorder   typeshape)
rA   packbitsarraybool_r%   DenseElementsAttrr@   IntegerTypeget_signlesslen)rD   as     r;   dense_bool_elementsrV   \   s[    kk"((2rxx(8<!				!	!bnn))!,SWI 
" 
? ?r=   c                |    t         j                  j                  t         j                  j	                  d      |       S )N    r%   IntegerAttrr@   rR   rS   is    r;   i32_attrr]   a   '    **2>>+F+Fr+JANNr=   c                |    t         j                  j                  t         j                  j	                  d      |       S )N@   rY   r[   s    r;   i64_attrra   b   r^   r=   c           	        t        t        j                  dt        j                              t        t        j                  dt        j                              fd}t        ||       }|sEt        t        j                  t        t        j                  g t        j                                    S t        |      dk(  r|d   S t        j                  |t        d            S )N)rJ    c                    t        |       t        u r.t        t        j                  | gt        j
                              S | j                   k7  rt        j                  |       } t        j                  |       S r6   )	rL   intir_constantrA   rO   int32r(   convertreshape)di32_typeint1ds    r;   	lower_dimzshape_tensor.<locals>.lower_dimh   sW    Aw#~1#rxx011	
8	KK!$[[""r=   rJ   r   )aval_to_ir_typer   ShapedArrayrA   rg   map	type_castr%   RankedTensorTyperf   rO   rT   r(   concatenatera   )sizesrm   dsrk   rl   s      @@r;   shape_tensorrv   d   s    
$**4:
;%T--b"((;<(# 9e"	R((+bhhr2886L*MNN
2w!|a5L??2x{++r=   c                r     | j                   di |} ||g| }| j                  |j                         |S )zSide-effects on `ctx`rc   )replaceset_tokens_out
tokens_out)ctxlowering_funargsctx_override_kwargsctx_newouts         r;   delegate_loweringr   x   s>    CKK.-.'W$t$#W''(	*r=   c                    t        | t        j                        ryt        | t              xr" t	        |       dk7  xr t        d | D              S )z8Returns true if `x` is an ir.Value or tuple of ir.ValuesTrJ   c              3  P   K   | ]  }t        |t        j                           y wr6   )r7   r%   Value.0vs     r;   	<genexpr>z _is_ir_values.<locals>.<genexpr>   s     5a*Q)5   $&)r7   r%   r   tuplerT   allr9   s    r;   _is_ir_valuesr      sD    288
Q
 63q6Q; 651557r=   rJ            rX   r`   c                 x    t         j                  j                  t         j                  j                               S r6   )r%   ComplexTyper@   F32Typerc   r=   r;   <lambda>r      s     ".."4"4RZZ^^5E"F r=   c                 x    t         j                  j                  t         j                  j                               S r6   )r%   r   r@   F64Typerc   r=   r;   r   r      s     2>>#5#5bjjnn6F#G r=      z%dict[np.dtype, Callable[[], ir.Type]]_dtype_to_ir_typec                |   t        | t        j                        r#t        j                  t        j
                        } t        | t        j                  t        j                  f      sJ t        |              t        j                  |       } 	 t        |    } |       S # t        $ r}t        d|        |d }~ww xY w)Nz'No dtype_to_ir_type handler for dtype: )r7   r   bintrA   dtyperg   genericrL   r   KeyError	TypeError)r   ir_type_factoryerrs      r;   dtype_to_ir_typer      s    tyy! HHRXXE	EBHHbjj1	2?DK?	2
((5/%D'.O 
	 
 D

1%9;@CDDs   	B 	B;'B66B;c                    t        j                  |       } t        j                  | j                        st	        |       S t
        j                  j                  | j                  t        | j                              S r6   )
r   physical_avalis_constant_shaperM   _dynamic_array_ir_typesr%   rr   r@   r   r   avals    r;   _array_ir_typesr      sV    			D	!$			

	+"4((				 	 -=djj-I	JJr=   c                   t         j                  j                         }| j                  D cg c]  }t	        |      t
        u r|n| }}t         j                  j                  |t        | j                              S c c}w r6   )
r%   
ShapedTypeget_dynamic_sizerM   rL   re   rr   r@   r   r   )r   dyn_sizerj   rM   s       r;   r   r      sc    ]]++-(6:jj
AQ31H,
A%
A				 	 (8(D	EE Bs   A>z8dict[type[core.AbstractValue], Callable[[Any], IrTypes]]ir_type_handlersc                    	 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   rL   r   r   )r   r   s     r;   rn   rn      sG    
PDJ'--	 P
8dE
FCOPs    	?:?c                >    t         j                  j                         S r6   )r(   	TokenTyper@   )_s    r;   r   r      s    1B1B1D r=   c                V    t        |       }t        |t        j                        r|fS |S r6   )rn   r7   r%   Type)r   typs     r;   aval_to_ir_typesr      s&    #c277+#44r=   c                      e Zd ZddZy)ConstantHandlerc                     y)zlBuilds an IR representation for a constant `val`.

    A JAX value is represented by zero or more IR values.Nrc   )selfvalr   s      r;   __call__zConstantHandler.__call__       r=   N)r   r	   r   core.AbstractValue | NonereturnIrValues__name__
__module____qualname__r   rc   r=   r;   r   r      s    =r=   r   zdict[type, ConstantHandler]_constant_handlersc                    |t         | <   y r6   r   type_handler_funs     r;   register_constant_handlerr      s    )Ur=   c                    t         |    S r6   r   r   s    r;   get_constant_handlerr      s    	E	""r=   const_loweringr   c                  |5t        j                  |       r |j                  t        |       |f      x}|S t	        |       j
                  D ]C  }t        j                  |      }|s || |      }t        |      sJ t	        |       |f       |c S  t        | d      rt        | j                               S t        dt	        |              )a)  Translate a Python `val` to an IR constant.

  See https://docs.jax.dev/en/latest/internals/constants.html.
  Args:
    val: a Python value to be translated to a constant.
    const_lowering: an optional dictionary with known lowering for some
      constants, indexed by `id`. This is used, e.g., when we pass constants
      as MLIR function arguments.
    aval: the abstract value of `val`, if known. Required where ambiguous, e.g.
      for Python scalars.

  Returns:
    A representation of the constant as an IR value or sequence of IR values.
  __jax_array__zNo constant handler for type: )rA   rM   r@   idrL   __mro__r   r   hasattrrf   r   r   )r   r   r   c_valthandlerr   s          r;   rf   rf      s    & 	xx}>#5#5r#wo#FF%Sl9 a $$Q'GCc31$s)S!11j S/"s((*++249+>??r=   c                D   t        | j                        }| j                  }| j                  t        j                  k(  rt        j
                  | d      } t        j                  |       } t        j                  j                  | ||      }t        j                  |      S NrG   rH   rK   )r   r   rM   rA   rP   rN   ascontiguousarrayr%   rQ   r@   r(   constant)r:   element_typerM   attrs       r;   _numpy_array_constantr   
  su    !!''*,
''%WW
A)A1!				!	!!,e	!	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r=   c                     t        d      )Nz\A ShapeDtypeStruct does not have a value and cannot be used as a constant in a JAX function.)r   r   s     r;   $_shape_dtype_struct_constant_handlerr     s     : 	; ;r=   c                   | j                   t        j                  k(  r8t        t	        j
                  | j                  t        j                              S d| j                  v r	| j                  dkD  rt	        j                  t	        j                  d| j                              \  t	        j                  t	        j                  d| j                              \  }| t        fdt        | j                        D                 }t!        j"                  t$        j&                  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  @   K   | ]  }|v rd n
t        d        yw)r   N)slice)r   axzero_stride_axess     r;   r   z,_ndarray_constant_handler.<locals>.<genexpr>8  s*      9"$ $&)9#9auT{J 9   )r   r   float0r   rA   zerosrM   rP   stridessizewhereequal	not_equalr   rangendimr(   broadcast_in_dimr%   rr   r@   r   dense_int_array)r   r   
other_axescollapsed_valr   r   s        @r;   _ndarray_constant_handlerr   !  s
   $ 	YY&-- #))288!DEECKKCHHqL!S[[!9:((2<<3;;78KJ 9(-chh9 9 :M



II'(;(;<	>m,
#	%C
 J %%r=   c                    t        |t        j                        sJ |       |j                  dk(  sJ |       t	        t        j                  | |j                              S Nrc   )r7   r   ro   rM   r   rA   rO   r   r   r   s     r;   _python_scalar_handlerr   M  sN    	D$**	+1T1	+	r	4		rxxTZZ8	99r=   c                *    t        j                         S r6   )r(   create_tokenr   s     r;   _token_constant_handlerr  U  s    				r=   z!dict[type[Any], AttributeHandler]_attribute_handlersc                    |t         | <   y r6   r  r   s     r;   register_attribute_handlerr  ^  s    *er=   c                    t         |    S r6   r  r   s    r;   get_attribute_handlerr  a  s    	U	##r=   c                B   t        | j                        }t        |t        j                        r t        j
                  j                  ||       S t        |t        j                        r t        j                  j                  ||       S t        dt        |              )Nz#Unsupported scalar attribute type: )r   r   r7   r%   rR   rZ   r@   	FloatType	FloatAttrr   rL   )r   	mlir_types     r;   _numpy_scalar_attributer  d  sq    syy))	2>>*>>i--)R\\*<<Is++
9$s)E
FFr=   c                   t        | j                        }| j                  }| j                  t        j                  k(  rt        j
                  | d      } t        j                  |       } t        j                  j                  | ||      S r   )
r   r   rM   rA   rP   rN   r   r%   rQ   r@   )r:   r   rM   s      r;   _numpy_array_attributer  m  sg    !!''*,
''%WW
A)A1!				!	!!,e	!	DDr=   c                j   d| j                   v r| j                  dkD  rt        d      | j                  t        j
                  k(  r/t        j                  | j                  t        j                        } t	        j                  |       st        j                  |       rt        |       S t        |       S )Nr   zCNumPy arrays with zero strides are not supported as MLIR attributesr   )r   r   r   r   r   r   rA   r   rM   rP   is_python_scalarisscalarr  r  )r   s    r;   _numpy_array_attribute_handlerr  u  s    #++#((Q,
MO OYY&--
((399BHH
-CS!R[[%5"3''!#&&r=   c                ,    t        | j                        S r6   )r  r   r9   s    r;   r   r     s    %CAEE%J r=   c                R    t         j                  j                  t        |             S r6   )r%   TypeAttrr@   r   r   s    r;   _dtype_attribute_handlerr    s    	)%0	11r=   c                @    t        t        j                  ||             S r6   )r  rA   rO   )r   r   s     r;    _python_scalar_attribute_handlerr    s    	 #u!5	66r=   c           
         t         j                  j                  | j                         D ci c]  \  }}|t	        |       c}}      S c c}}w r6   )r%   DictAttrr@   itemsir_attribute)r   kr   s      r;   _dict_attribute_handlerr    s6    	EA!\!_,E	FFEs   A
c                x    t         j                  j                  | D cg c]  }t        |       c}      S c c}w r6   )r%   	ArrayAttrr@   r  )r   r   s     r;   _sequence_attribute_handlerr!    s)    			C8q<?8	998s   7c                    | S r6   rc   r9   s    r;   r   r     s    1 r=   c                    | S r6   rc   r9   s    r;   r   r     s    a r=   c                L   t        |       j                  D ]Q  }t        j                  |      }|s ||       }t	        |t
        j                        sJ t        |       |f       |c S  t        | d      rt        | j                               S t        dt        |              )z,Convert a Python value to an MLIR attribute.r   z'No attribute handler defined for type: )rL   r   r  r@   r7   r%   	Attributer   r  r   r   )r   r   r   r   s       r;   r  r    s    9 a!%%a(GCLcR\\*<T#Y,<<*j S/"))+,,;DI;GHHr=   TracebackCachesc                    |j                   j                  | d       }||S t        j                  j                  }|rt        j                  |d|       } | |j                   | <   | S )N )canonical_name_cacher@   r   &hlo_source_file_canonicalization_regexvalueresub)	file_namecachescanonical_file_namepatterns       r;   get_canonical_source_filer2    sc    3377	4H$99??'wI.I+4&i(	r=   ModuleContextc                L    | j                   j                  j                  |      S )z8Converts a full traceback to a callsite() MLIR location.)traceback_cachestraceback_to_location_cacher@   )r{   tbs     r;   _traceback_to_locationr8    s    				9	9	=	=b	AAr=   c                   t         j                  j                  r.|t        j                  j                         }nt        | |      }nt        j                  |      }|t        j                  j                         }nSt        j                  j                  t        |j                  | j                        |j                  |j                        }|8|j                  r*t        j                  j!                  t#        |      |      }|S |j                  r| d|j                    n|j                   }t        j                  j!                  ||      }t        j                  j!                  |j                    d|      }|S )N)childLoc/:)r   $include_full_tracebacks_in_locationsr+  r%   Locationunknownr8  r   
user_framefiler2  r.  r5  
start_linestart_columnstacknamestr)r{   	primitive
name_stack	tracebacklocframeeqn_strs          r;   source_info_to_locationrM    s>    0066KK!c"3	2c''	2E}KK!cKK6u7:7K7KM"--u/A/ACc KKS_s;c 
*	 -7,<,<:,a	'()..  ++

7S

1C
++

inn-Q/#

>C	*r=   c                   t        j                  t        j                  j                        x}syt        j
                  j                  j                  d      }d|vryt        t              }| j                  j                  d   }t        j                  |      j                  }d|ddt        |       d| d}||z  }|j                  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,	stablehlosym_namejax_ir04dr   z.mlir)r   make_jax_dump_dirr   jax_dump_ir_tor+  jax_dump_ir_modessplitnext_ir_dump_counter	operation
attributesr%   
StringAttr_make_string_safe_for_filename
write_textmodule_to_string)	module
stage_nameout_dirmodesr   rQ  module_namerE  	full_paths	            r;   dump_module_to_filerf    s     ++F,A,A,G,GH
H'
H

"
"
(
(
.
.s
3%"((4(h'--+"S9+FGqTY	Z$n)'/0	+r=   c                ,    t        | |      }|rd| dS y)NzThe module was dumped to .z)Define JAX_DUMP_IR_TO to dump the module.)rf  )r`  ra  	dumped_tos      r;   dump_module_messagerj    s#    !&*5)&yk336r=   c                0    t        j                  dd|       S )Nz
[^\w.)( -]r(  )r,  r-  ss    r;   r]  r]    s    	r1	%%r=   c                    t        j                         }|5t        t        j                  j
                        j                         }|dv}| j                  j                  ||       |j                         S )N)false0)rA  enable_debug_info)
ioStringIOrF  r   jax_include_debug_info_in_dumpsr+  lowerrZ  printgetvalue)r`  rq  outputenable_debug_flags       r;   r_  r_    se    ;;=&FBBHHIOOQ)?f8IJ		r=   c                    t        j                         }| j                  j                  |       |j	                         S )N)rA  )rr  BytesIOrZ  write_bytecoderw  )r`  rx  s     r;   module_to_bytecoder}    s1    ::<&!!v!.		r=   c                       e Zd Z fdZ xZS )JaxIrContextc                @    t        t        j                  |   |i | y r6   )superr%   Context__init__)r   r}   r   	__class__s      r;   r  zJaxIrContext.__init__#  s    
 
"**d$d5f5r=   )r   r   r   r  __classcell__)r  s   @r;   r  r  "  s    6 6r=   r  c                    t               } | j                  t               | j                          | j	                  t
               t        j                  j                  |        t        j                  j                  |        t        j                  j                  |        t        j                  j                  |        t        j                  j                  |        | S )z,Creates an MLIR context suitable for JAX IR.)r  append_dialect_registryupstream_dialectsload_all_available_dialectsset_thread_poolglobal_thread_poolr$   sdyregister_dialectmpmdmhloregister_mhlo_dialectchlor(   )contexts    r;   make_ir_contextr  *  s    N'	!!"34	%%'	,-
,,(
--  )
--%%g.
--  )
,,( 
.r=   c                  <    e Zd ZU ded<   ded<   ded<   	 	 	 	 ddZy)	ShapePolyLoweringStatetuple[str, ...]dim_varsbooluses_dim_varshas_platform_index_argumentc                    |$t        |      dkD  rdt        |      z   }d| _        nd| _        t        |      dkD  | _        || _        y )NrJ   )_platform_indexTFr   )rT   r   r  r  r  )r   r  lowering_platformss      r;   r  zShapePolyLoweringState.__init__Z  sM     %#.@*AA*E%h7h)-d&).d&h-!+DDMr=   N)r  r  r  ztuple[str, ...] | Noner   r   r   __annotations__r  rc   r=   r;   r  r  F  s/     
   $#	(	#9	r=   r  T)frozenc                  |    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<   ej                  j                  Zded	<   y)
LoweringParametersNz1tuple[tuple[core.Primitive, LoweringRule]] | Noneoverride_lowering_rulesFr  global_constant_computation
for_export#export_ignore_forward_compatibilityhoist_constants_as_args)r   r   r   r  r  r  r  r  r   use_simplified_jaxpr_constantsr+  r  rc   r=   r;   r  r  e  sQ    
 PTLS ',t+ *d /4%t3 #)"G"G"M"M4Mr=   r  c                    t        j                  | j                        syt        j                  j
                  }|r!t        j                  |d| j                        S | j                  S )zxReturns the canonicalized filename of a code object.

  Returns None if the filename should be omitted in tracebacks.
  Nr(  )r   is_user_filenameco_filenamer   r*  r+  r,  r-  )coder1  s     r;   _code_to_filenamer    sR    
 
	*	*4+;+;	<99??'29T--	.Ot?O?OOr=   c                  (    e Zd ZU ded<   ded<   d Zy)r&  r	   r6  zdict[str, str]r)  c                    t         j                  j                  }|dk\  r|nd}t        j                  t
        |      | _        i | _        y )Nr   i  )code_to_filenameframe_limit)r   traceback_in_locations_limitr+  r"   TracebackToLocationCacher  r6  r)  )r   r  s     r;   r  zTracebackCaches.__init__  sB    55;;K!,!1+tK'3'L'L*(ED$ "Dr=   Nr  rc   r=   r;   r&  r&    s    ""&&#r=   c                  J    e Zd ZU ded<   ded<   ded<   ded<   d	ed
<   ded<   y)LoweringCacheKeycore.PrimitiverG  core.JaxprEqnContexteqn_ctxztuple[core.AbstractValue, ...]avals_ineffects_lib.Effectsr   z frozen_dict.FrozenDict[str, Any]paramsr  	platformsNr   r   r   r  rc   r=   r;   r  r    s&    **
**r=   r  c                  @    e Zd ZU ded<   ded<   ded<   ded<   d	ed
<   y)LoweringCacheValuefunc_dialect.FuncOpr'   Sequence[IrTypes]output_typeszSequence[ArrayLike]
const_argsSequence[core.AbstractValue]const_arg_avalsr  inlineNr  rc   r=   r;   r  r    s    !!!!//
,r=   r  c            
      D   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<   ded<   ded <   ed+d!       Zd"d"d"d"d"d"d"d"d#d$		 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d,d%Zd-d.d&Zd/d'Z	d0d(Z
d1d)Zd* Zy")2r3  z2Module-wide context information for MLIR lowering.
ir.Contextr  	ir.Moduler`  zir.InsertionPointipzir.SymbolTablesymbol_tableSequence[str]r  xc.Client | NonebackendAxisContextaxis_context	list[Any]
keepalivesIterator[int]channel_iteratorhost_callbacksr  shape_poly_stater  all_default_mem_kindz*dict[LoweringCacheKey, LoweringCacheValue]lowering_cachezdict[Any, func_dialect.FuncOp]cached_primitive_loweringsr&  r5  r  lowering_parametersc                .    | j                   j                  S r6   )r  axis_envr   s    r;   r  zModuleContext.axis_env  s    %%%r=   NT)	r  r`  r  r  r  r  r5  r  r  c                  |xs
 t               | _        |	xsG t        j                  j	                  t        j
                  j                  | j                              | _        |
xs) t        j                  | j                  j                        | _
        |xs) t        j                  | j                  j                        | _        || _        || _        || _        |i n|| _        |i n|| _        | j                  5  |
t'               n|| _        d d d        || _        || _        || _        |xs t1        dt3        |            | _        || _        || _        y # 1 sw Y   KxY w)N)rJ  rc   )r  r  r%   Modulecreater>  r?  r`  InsertionPointbodyr  SymbolTablerZ  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  r  r  r5  r  r  s                    r;   r  zModuleContext.__init__  s5   ( /o/DLSBII,,1D1DT\\1R,SDK7B%%dkk&6&67DG$Mt{{7L7L(MDDLDN$D!/!72^D-G-Or,F 	#	 54D4L0#3 5 -D DO(DF0U95EF 	 4D2D5 5s   ?EE&c                   t        | j                        dkD  r|ry t        d      | j                  t	        j
                  | j                  j                        | j                  d   k7  rH|ry t        dt	        j
                  | j                  j                         d| j                  d          | j                  S t	        j                  | j                  d         S )NrJ   zaccessing .backend in multi-lowering setting. This can occur when lowering a primitive that has not been adapted to multi-platform loweringr   z'the platform for the specified backend z) is different from the lowering platform )	rT   r  NotImplementedErrorr  xbcanonicalize_platformplatformr   get_backend)r   optionals     r;   r  zModuleContext.get_backend  s    
4>>Q		  ||		!	!$,,"7"7	8DNN1<M	M
3%%dll&;&;<= >((,q(9':<= 	= \\>>$..+,,r=   c                P    t        | j                        }|dk\  rt        d      |S )Ni   z`Host callback lowering created too many channels. PjRt does not support more than 65535 channels)rX  r  RuntimeError)r   channels     r;   new_channelzModuleContext.new_channel  s3    4(()G7./ / Nr=   c                :    | j                   j                  |       y r6   )r  append)r   host_callbacks     r;   add_host_callbackzModuleContext.add_host_callback  s    }-r=   c                :    | j                   j                  |       y r6   )r  r  )r   	keepalives     r;   add_keepalivezModuleContext.add_keepalive  s    OO9%r=   c                .    t        j                  | fi |S r6   dataclassesrx   r   kws     r;   rx   zModuleContext.replace      +"5"5d"Ab"AAr=   )r   zsharding_impls.AxisEnv)r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  zir.Context | Noner`  zir.Module | Noner  zir.InsertionPoint | Noner  ir.SymbolTable | Noner  z"None | dict[LoweringCacheKey, Any]r  z%None | dict[Any, func_dialect.FuncOp]r5  zNone | TracebackCachesr  r  )F)r  r  r   r  )r   re   )r  r	   r   None)r  r	   r   r  )r   r   r   __doc__r  propertyr  r  r  r  r  r  rx   rc   r=   r;   r3  r3    sQ   :
 !!** =<<< $#))& & $(!%%),0;?JN15#'%'3 '3  	'3
  '3 '3 &'3  '3 .'3 !'3 '3 #'3 *'3 9'3 #H'3  /!'3$ !%'3R-&.& B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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y)#LoweringRuleContextz/Per-rule context information for MLIR lowering.r3  module_contextsource_info_util.NameStackrH  xc.Traceback | NonerI  core.Primitive | NonerG  r  r  r	   	avals_outTokenSet	tokens_inzTokenSet | Nonerz   .dict[tuple[int, core.AbstractValue], IrValues]r   Nzdict[core.Var, ir.Value] | Noneaxis_size_envrc   Sequence[ir.Value]dim_var_valuescore.JaxprEqnContext | Nonejaxpr_eqn_ctxSequence[str] | Noner  c                8    | j                   J d       || _         y )Nz"Should only set `tokens_out` once.)rz   )r   rz   s     r;   ry   z"LoweringRuleContext.set_tokens_out=  s    ??"H$HH" DOr=   c                .    t        j                  | fi |S r6   r  r  s     r;   rx   zLoweringRuleContext.replaceA  r  r=   c                    | j                   j                  }| j                  xs | j                   j                  }t        d |D              }|j                  xs |xr |j
                   S )zOReturns true if the lowering parameters are in forward compatibility mode.
    c              3  @   K   | ]  }|t         j                  v   y wr6   )r  'FORCE_FORWARD_COMPAT_LOWERING_PLATFORMSr   ps     r;   r   z8LoweringRuleContext.is_forward_compat.<locals>.<genexpr>K  s       <=R777   )r  r  r  anyr  r  )r   r  check_platformsforce_forward_compats       r;   is_forward_compatz%LoweringRuleContext.is_forward_compatC  sx     --AA 	7$--77   AP 
 	&&>*>F!EE
EFr=   )rz   r  )r   r  )r   r   r   r  r  r
  r  r  r  ry   rx   r  rc   r=   r;   r  r    s    7 )(  ""((.
 A@37-07 (*.$)/3-,3 %))!(! BFr=   r  c                      e Zd Z	 	 	 	 ddZy)LoweringRulec                     y)z.Converts a JAX primitive invocation into MLIR.Nrc   )r   r{   r}   r  s       r;   r   zLoweringRule.__call__V  r   r=   N)r{   r  r}   zir.Value | Sequence[ir.Value]r   z'Sequence[ir.Value | Sequence[ir.Value]]r   rc   r=   r;   r  r  U  s    ;5;A;r=   r  c                  "    e Zd ZU ded<   ded<   y)LoweringRuleEntryr  ruler  r  Nr  rc   r=   r;   r   r   ]  s    
,r=   r   z'dict[core.Primitive, LoweringRuleEntry]
_loweringsz2dict[str, dict[core.Primitive, LoweringRuleEntry]]_platform_specific_loweringsc           
        t        |t              rJ |st        j                  |        |t        ||      t        | <   yt        j                  |      sAt        t        j                               }t        d|  d| ddj                  |       d      t        j                  |      D ]  }t        ||      t        |   | <    y)a  Registers a lowering rule for a primitive.

  Args:
    prim: The primitive to register the rule for.
    rule: The lowering rule to register.
    platform: The platform to register the rule for. If None, this is a common
      rule applicable to all platforms. Platform-specific rules take precedence
      over common rules.
    inline: Whether to emit the lowering inline. If False, the lowering will be
      emitted in a separate function, called by similar instances of the
      lowering.
    uncacheable: Whether this primitive's lowering can be cached. This is a
      temporary flag that will be removed after primitives that have problems
      with caching are fixed.
  Nz0Registering an MLIR lowering rule for primitive z for an unknown platform z. Known platforms are: , rh  )r7   r   _uncacheable_primitivesaddr"  r  is_known_platformsortedknown_platformsr  joinexpand_platform_aliasr#  )primr!  r  r  	cacheabler*  r  s          r;   register_loweringr/  f  s    $ /00	0	%(v6Jt)r1134o<TF%hZ 0ii(),- - %%h/ N.?f.M"1%d+Nr=   c                    g }| D ]?  }t        |t        j                        r|j                  |       /|j	                  |       A |S )z@Concatenates/flattens a list of ir.Values or ir.Value sequences.)r7   r%   r   r  extendrD   r   r:   s      r;   flatten_ir_valuesr3    sA    
# a!RXX	jjm	jjm	
 
*r=   c                    g }| D ]?  }t        |t        j                        r|j                  |       /|j	                  |       A |S )z>Concatenates/flattens a list of ir.Types or ir.Type sequences.)r7   r%   r   r  r1  r2  s      r;   flatten_ir_typesr5    sA    
# a!RWW	jjm	jjm	
 
*r=   c           	         t        |       |D cg c]/  }|dk(  rt              nt        fdt        |      D              1 }}t        t              t        u sJ |S c c}w )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.rJ   c              3  4   K   | ]  }t                y wr6   rX  r   r   xs_iters     r;   r   z%unflatten_ir_types.<locals>.<genexpr>  s      4$ 59M 4$   )iterrX  r   r   _unflatten_done)rD   nsnunflattenedr:  s       @r;   unflatten_ir_typesrA    st    
 H' .01() #$q&ge 4$(4$ /$ $ 1+ 1	g	'?	::	:	1s   4A!c                P    t        | t        j                        rdS t        |       S )NrJ   )r7   r%   r   rT   r9   s    r;   len_ir_typesrC    s    BGG$0#a&0r=   c                   t        |       |D cg c]M  }t        |t        j                        rt	              n%t        fdt        t        |            D              O }}t	        t              t        u sJ |S c c}w )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              3  4   K   | ]  }t                y wr6   r8  r9  s     r;   r   z1unflatten_ir_values_like_types.<locals>.<genexpr>  s     <tG}<r;  )	r<  r7   r%   r   rX  r   r   rT   r=  )rD   ysyr@  r:  s       @r;   unflatten_ir_values_like_typesrH    s{     H'  #-Q"8g<eCFm<<= +  
g	'?	::	:		s   AA?z[^\w.-]c                .    t         j                  d|       S )z3Ensure a name is usable as module or function name.r   )_module_name_regexr-  rE  s    r;   sanitize_namerL    s    			T	**r=   c                    || S t        |t              r| S t        | t        j                        r| S t        | t        j                        st
        | j                  |j                  | j                        d      S )z5Returns the new aval sharded based on sharding proto.Nsharding)	r7   r,   r   AbstractTokenro   r  updateshard_shaperM   )r   rO  s     r;   sharded_avalrS    sk     K$Kd(()K	D$**	+
	X))$**5	EEr=   c           	        | j                  dt        j                         gt        | j                  j
                  j                        z  d       }  t        t        t        j                  || j                  j
                  j                        d      | g| j                   }t        d t        |t        |            D              S )Neval_dynamic_shape)rG  r  rz   T)multiple_resultsc              3  z   K   | ]3  \  }}t        j                  |      rt        j                  |      n| 5 y wr6   )r   is_constant_dimoperatorindex)r   rj   d_irs      r;   r   z%eval_dynamic_shape.<locals>.<genexpr>  s9      B1d %)$8$8$;x~~a E Bs   9;)rx   r   dim_value_avalrT   r  r  r  	lower_funr   evaluate_shaper  r   zipr3  )r{   rM   ress      r;   rU  rU    s    $##%&S-?-?-P-P-Y-Y)ZZ 	 	#
		d!!5#*<*<*M*M*V*VW	 	7#&#5#5	7# 
 B!%):3)?@B 
B Br=   c                H    ddt        fdt        | |      D              S )z-Evaluates the dynamic shapes as int32 values.c                (   t        |       t        u r.t        t        j                  | t        j
                              S t        t        j                  dt        j
                              }| j                   |k7  rt        j                  ||       S | S )Nr   rc   )rL   re   rf   rA   rO   rg   rn   r   ro   r(   rh   rj   rk   s     r;   convert_dimz/eval_dynamic_shape_as_vals.<locals>.convert_dim  sd    Aw#~!288455 !1!1"bhh!?@h	
8	{{8Q''r=   c              3  .   K   | ]  } |        y wr6   rc   r   r   rd  s     r;   r   z-eval_dynamic_shape_as_vals.<locals>.<genexpr>       F!{1~F   )rj   int | Valuer   rU  r{   rM   rd  s     @r;   eval_dynamic_shape_as_valsrl    s#     
F'9#u'EF	FFr=   c                H    ddt        fdt        | |      D              S )z7Evaluates the dynamic shapes as int or ir.int32 values.c                    t        |       t        u r| S t        t        j                  dt
        j                              }| j                   |k7  rt        j                  ||       S | S r   )	rL   re   rn   r   ro   rA   rg   r(   rh   rc  s     r;   rd  z0eval_dynamic_shape_as_ivals.<locals>.convert_dim  sO    Aw#~h !1!1"bhh!?@h	
8	{{8Q''r=   c              3  .   K   | ]  } |        y wr6   rc   rf  s     r;   r   z.eval_dynamic_shape_as_ivals.<locals>.<genexpr>  rg  rh  )rj   ri  r   zint | ir.Valuerj  rk  s     @r;   eval_dynamic_shape_as_ivalsrp    s#     
F'9#u'EF	FFr=   c                ,    t        t        | |            S )z4Evaluates the dynamic shapes as one 1d int32 tensor.)rv   rU  )r{   rM   s     r;   eval_dynamic_shape_as_tensorrr    s     
(e4	55r=   c                  6    e Zd ZU ded<   ded<   ded<   ded<   y	)
LoweringResultr  r`  z
Any | Noner  r  r  r  r  Nr  rc   r=   r;   rt  rt    s    
**r=   rt  )cpucudarocmtpuneuronc                  	 | j                   j                  }|j                   j                  }t        |t        j                        r(|j
                  |j
                  k(  r||j                  }}n*|t        j                  |j                  |      |      d   }}|j                  | j                  D ci c]  }|t        j                   c}      }t        j                  |||j                        }|j                   j                  	t        	fd|j                  D              rt        d|j                   d	 d      |S c c}w )Nr   )memory_kindc              3     K   | ]6  }|2|t         j                  ur t        |t              r|n|fD ]  }|v  
 8 y wr6   )r*   UNCONSTRAINEDr7   r   )r   rm  r  manual_axess      r;   r   z"add_manual_axes.<locals>.<genexpr>  sO      	9aMa}'B'BB"1e,11$	9 
k	 	9	 	9s   <?zpspec z contains a manual axes z of mesh which is not allowed. If you are using a with_sharding_constraint under a shard_map, only use the mesh axis in PartitionSpec which are not manual.)meshabstract_meshr7   r   r-   rM   specparse_flatten_op_sharding_to_xla_hlo_shardingupdate_axis_typesr~  r)   Manualr{  r  r   )
axis_ctxrO  r   r  sharding_meshout_meshr  rU   r   r~  s
            @r;   add_manual_axesr    s(   		$	$$----->778TZZ'"HMMdH>CC%%d+T3346dH ''#+#7#78aq(//8:($$Xt191E1E	G#$$+ 	9388 	9 9 

2;- @< 	<= =
 
* 9s   3Ec                X   |y t        ||      ry t        |t              r6t        j                  j
                  r|j                  |j                        S y t        |t              sJ t        |t              rt        | |j                  |      S t        |t        j                        sJ t        j                  |j                   t        j"                        r+t%        j&                  ||      }t        j(                  |      }| j*                  }t        |t$        j,                        r#|j.                  rt1        |||j                        }t        j                  j
                  r|j                  |j                        S |j3                  |j                        j5                         S r6   )all_unconstrainedr7   r,   r   use_shardy_partitionerr+  _to_sdy_shardingr   	JShardingr1   _to_physical_op_sharding
inner_avalr   ro   r   
issubdtyper   extendedr   physical_shardingr   r  SPMDAxisContextr~  r  r  to_proto)r{   r   rO  r  s       r;   r  r  *  s@    x&$$$**&&tyy11	Hi	((	(k"#C(CC	D$**	++	+tzz6??3//h?Hd#D(>99:x499=H""(($$TYY//		&	&tyy	1	:	:	<<r=   c                    | y t        | t              ry|t        j                  u ry t	        | j                  |j                              S )Nauto)r7   r   r   abstract_tokenrF  _to_xla_layoutr   )layoutr   s     r;   r  r  F  sB    ^
#	T   	V""4::.	//r=   c                f    | y t        | t              ry t        | t              sJ | j                  S r6   )r7   r,   r  r{  rl  s    r;   _get_mem_kindr  Q  s1    Y4	Ay	!!	!	
r=   c                ^    t        | t              xr t        j                  | j                  v S r6   )r7   r-   r*   r}  r  rl  s    r;   contains_unconstrainedr  Z  s(    
Q
& 4))QVV35r=   c                    t        | t              rO|j                  dk(  ry|j                  t        | j                        k7  ryt        d | j                  D              S y)Nr   Fc              3  @   K   | ]  }|t         j                  u   y wr6   )r*   r}  r  s     r;   r   z$all_unconstrained.<locals>.<genexpr>e  s     @AqM///@r  )r7   r-   r   rT   r  r   )rm  r   s     r;   r  r  _  sI    =!yyA~yyCK@@@@	r=   c                  "    e Zd ZU ded<   ded<   y)UnconstrainedVariantsr  r  r  Nr  rc   r=   r;   r  r  h  s    r=   r  c                F    t        |       }t        |t        | |            S )N)r  r  )r  r  r  )rm  r   uss      r;   _get_unconstrained_variantsr  l  s(    a "	3DQ3M
O Or=   c           
     t   t         j                  j                  ryt         j                  j                  x}dk(  ryt	        j
                  t        d | j                              \  }}t        |      x}|k  rydt        j                  |       d}t         j                  j                  x}s|dz  }t        j                  |       y|dt        |t        | j                               dz  }	 t!        || j"                  j$                  | j                        }t'        j(                  d	|t+        j,                  d
            D ]  \  }}	}
|dt/        |
       d|	j0                  j3                          dt        j                  |       dz  }t5        j6                  | j"                  |	      D ]=  }t9        j:                  |j<                  |      }|d|j?                  dd      z   dz   z  }?  t        j                  |       y# t@        $ r%}t        j                  |d| z          Y d}~yd}~ww xY w)zbCheck if a JAXPR contains an excessive amount of constants, if so, report where they were capturedNc                    t        | dd      S )Nnbytesr   )getattr)cs    r;   r   z'check_jaxpr_constants.<locals>.<lambda>{  s    GAx+ r=   z;A large amount of constants were captured during lowering (zg total). If this is intentional, disable this warning by setting JAX_CAPTURED_CONSTANTS_WARN_BYTES=-1. zjTo obtain a report of where these constants were encountered, set JAX_CAPTURED_CONSTANTS_REPORT_FRAMES=-1.zbThe subsequent report may be disabled by setting JAX_CAPTURED_CONSTANTS_REPORT_FRAMES=0.

Largest z allocation(s):
   r   )keyz  Constant r%  z captured at:
z    
z
    z

z+ Exception raised while generating report: )!r   r  r+  captured_constants_warn_bytes	itertoolsteerp   constssumr   pprint_bytes captured_constants_report_frameswarningswarnminrT   r_  jaxpr	constvarsheapqnlargestrY  
itemgetterrL   r   	str_shortr   eqns_using_varr   	summarizesource_inforx   	Exception)closed_jaxpr	threshold
total_iternbytes_itertotal_bytesmessage
num_framesnbytes_var_constr  varconsteqncall_frame_source_infoexcs                 r;   check_jaxpr_constantsr  r  s.   **00
77===i"D
 &MM	
+\-@-@A*k _$k	1
[)
* +PP 
 ??EE
E*
E	7G MM'
	Z\%8%8!9:;;LN'Q;(:(:(D(DlFYFYZ#nnQ0@hFYFYZ[F\] ]U;tE{m2chh.@.@.B-C2dFWFWX^F_E``oppg**<+=+=sC ]#!1!;!;COOZ!X84<<T?SSV\\\]] MM'	 QMM'I#OOPPQs   1DH	 		H7H22H7)replicated_argsarg_shardingsresult_shardings
in_layoutsout_layouts	arg_namesresult_namesnum_replicasnum_partitionsr  input_output_aliasespropagated_out_mem_kindsc               P  -. t        j                  d       t        t        t        j
                  |            }|
|nt        t        ||
      }||j                  nt        t        |j                  |      }|rd}d}n(|
t        t        |
      nd}|t        t        |      nd}d}|D cg c]  }|t        v r| }}|rt        |      t        |      k7  r|st        |      rt        d| d| d      t        |||||||||dkD  r|nd	      \  }}}|dkD  rN|t        d |D              r:|dgt        |      z  }t        t        |            D ]  }||   s	d	||<   d||<    t        |      rbt        ||      D cg c]  \  }}|s	t!        |       } }}d
}!|sd| d|! }!| r*t#        j$                  ddj'                  |        d|!        ~t(        j*                  j-                  |j.                        }"|"rt1        d|j.                         t3        j4                  t6        dz         }#g }$g }%|D &cg c]7  }&t9        |&d      s|&j:                  D ]  }t=        j>                  |      s| 9 }'}&}t        tA        tC        jD                  d |'tG                                 }(tI        ||||$|#|%|tK        |(|      |	      })|)jL                  5  tN        jP                  jS                  |)jL                        5  |)jT                  jV                  jX                  }*tN        jZ                  j]                  t_        |       ja                  d            |*d<   tc        |      |*d<   tc        |      |*d<   te        |)| ||f|d	|	||
||||||||||d ddd       ddd       	 |)jT                  jV                  jg                         s"t1        dti        |)jT                  d      z         	 |)jL                  5  tn        jp                  js                  d      },|,ju                  |)jT                  jV                         tv        jx                  jz                  rDtn        jp                  js                  d      },|,ju                  |)jT                  jV                         ddd       t        j                  d||)jT                         t}        |)jT                  |)j~                  |)j                  |)j                        S c c}w c c}}w c c}}&w # 1 sw Y   yxY w# 1 sw Y   ~xY w# tN        jj                  $ r\}+dg.-.fd-|+jl                  D ]
  } -|        t1        dj'                  .      dz   ti        |)jT                  d      z         |+d}+~+ww xY w# 1 sw Y   xY w) zLowers a top-level jaxpr to an MLIR module.

  Handles the quirks of the argument/return value passing conventions of the
  runtime.
  The inputs already account for the constant arguments.
  See https://docs.jax.dev/en/latest/internals/constants.html
  lower_jaxpr_to_moduleNzeIn multi-platform lowering either all or no lowering platforms should support donation. Lowering for z of which only z support donationrJ   c              3  b   K   | ]'  }|d u xs t        |t              xs t        |       ) y wr6   )r7   r,   r  r   rm  s     r;   r   z(lower_jaxpr_to_module.<locals>.<genexpr>  s6      ( $YJ*Q-J1G1JJ (s   -/FTzNSee an explanation at https://docs.jax.dev/en/latest/faq.html#buffer-donation.z Donation is not implemented for z.
z&Some donated buffers were not usable: r%  !Cannot lower jaxpr with effects: rM   c                @    | j                  |j                               S r6   )union	_get_vars)accnews     r;   r   z'lower_jaxpr_to_module.<locals>.<lambda>	  s    CIIcmmo<V r=   )	r  r  r  r  r  r  r  r  r  r   rQ  zmhlo.num_replicaszmhlo.num_partitions)num_const_argsmain_functionr  in_avalsr  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                    j                  d| j                          j                  d| j                          | j                  D ]
  } |        y )N	z		at )r  r  locationnotes)rj   r?  emit_diagnostic_info	msg_liness     r;   r  z3lower_jaxpr_to_module.<locals>.emit_diagnostic_info6  sP    AII;'(-.ww  !Q r=   r  zbuiltin.module(symbol-dce)z'builtin.module(sdy-lift-inlined-meshes)zmlir.collect_lowered_jaxprs)Br   
test_eventr   rp   r  r  rS  	out_avalsr  _platforms_with_donationrT   r  r  _set_up_aliasesr   r_  rF  r  r  r+  effects_liblowerable_effectsfilter_not_inr   r   r  countCOLLECTIVE_CHANNEL_IDr   rM   r   rX  r)  	functoolsreducesetr3  r  r  r%   r>  r?  r`  rZ  r[  r\  r@   rL  rstripr]   lower_jaxpr_to_funverifyrj  	MLIRErrorerror_diagnosticsr&   PassManagerparserunr   r  r+  rt  r  r  r  )/rd  r  r  r  ordered_effectsr  r  r  donated_argsr  r  r  r  r  r  r  r  r  r  r  r  r  sharded_in_avalssharded_out_avalsr  r  r  r  platforms_with_donationinput_idrU   rj   unused_donationsmsgunlowerable_effectschannel_iterr  r  r   all_dim_polyr  r{   attrsepipeliner  r  s/                                                @@r;   r  r    s   B //)*C00)<=)"/"7h,-@ *:*Bu<:JK  )4 M=9:>  .9 }.>??C  (1 ?1 !%= =  ? ?
"#s9~5C-	11: <'((9	;< < <K.0A<-z;*Q.D<:8,(8 			!	 (&( 
(		!!7S%66C-. )(!'+
8
$#(,x
 ) 	+./?+NT41aRSATT
ZC".ykSEBcmm 		"234Cu> ? #55CCmm
8H
II !6!:;,* . 0 FGD'4J::FT-A-A!-D  F! F, F6)**+V+7@ A B( 	g )!+'3%3*='=h	'R+?	A# {{ ;BKK''4 ; JJ  ++E))k"))#.0E*!),!7E
#+N#;E
 [%;%'#)1)!)/"!9!;; ;4M::&&(
5
cjj.
9:; ; )  {{ 	)&&,,$&HLL%%&$$**((..
35hll3::''(	) ///

C	

CNNC4F4F,,
. .e?0 U*F; ; ; ;> 
 
M;<I 
    1
TYYy)D0(^DE FKLM
M	) 	)sj   .S?
TT T
*T
*TBTT*AT* =B#VT	TT'*V=AVVV%c	                t   | d gt        |      z  } nt        |       } d }	t        |	|      }t        |	|      }|&|$t        d |D              st        d |D              rd gt        |      z  }d gt        |      z  }t	        j
                  t        j                        }
t        t        ||||             D ]&  \  }\  }}}}|s||
||f   j                  |       ( d }t        |      }|d uxr |d u}t        t        ||            D ]  \  }\  }}||f}|
j                  |d      s!|
|   j                         }d||<   |r>t        ||   t              r+t        ||   t              st        d||    d||    d      |r>t        ||   t              s+t        ||   t              rt        d||    d||    d	      ||||   ||   k(  r*|"||   x}!t        |t              st!        |      s|| |<   |dgt        |      z  }d
||<   
 | D ch c]  }||	 }}t	        j
                  t        j                        }t        t        ||            D ]@  \  }\  }}||vs|t"        j$                  us!||j&                  |f   j                  |       B t)        t        |            D ]s  }||   s	||   t"        j$                  us||   j&                  ||   f}|j                  |d      sF||   j                          d||<   |dgt        |      z  }d
||<   u | ||fS c c}w )Nc                ~    | t         j                  u r| S t        j                  | j                  | j                        S r6   )r   r  ro   rM   r   )rU   s    r;   r   z!_set_up_aliases.<locals>.<lambda>X  s0    1(;(;#;a "..qww@ r=   c              3  $   K   | ]  }|d u  
 y wr6   rc   )r   rU   s     r;   r   z"_set_up_aliases.<locals>.<genexpr>d  s     	.!t)	.   c              3  $   K   | ]  }|d u  
 y wr6   rc   r   rs     r;   r   z"_set_up_aliases.<locals>.<genexpr>e  s     	1!t)	1r  rc   FzInput layout being donated was z while output layout was z. Did you mean to set the **output layout** to **Layout.AUTO**?
This will allow for the input and output layout to be chosen by XLA and not the layout of the output which might not be optimal.z. Did you mean to set the **input layout** to **Layout.AUTO**?
This will allow for the input and output layout to be chosen by XLA and not the layout of the input which might not be optimal.T)rT   listrp   r  collectionsdefaultdictdeque	enumerater_  r  r@   popleftr7   r   r   r,   r  r   r  r   r   )r  r  r  r
  r  r  r  r  r  strip_metadata	donationsr\   r   amdonatedaliasedr  out_donated_argsin_out_layout_not_nonermr  r  rm  aliased_output_idsresults_not_matched	input_idxs                             r;   r  r  O  s    ! 6CM1 45B.*(.),) "5"=		.-	..		10	11vH-&3y>1%%k&7&78))2	($l4HI*K &%a	%$GW7?r
""1%&
 ,'%T1Mk6M Y0C!DE #*ma$ *C}}S"3'')h#(x 
 
Z):
6Q4-j.B-C D""-a.!1 2HHI 	I !H-z:
[^Z
0-j.B-C D""-a.!1 2??@ 	@ 
 3
X
+a.
0&$Q''q0D!%;A%>)*X& ##Ws8}4
%)"G#*J $8Ia1=II#//0A0AB Y0C!DE 5ma$""t43F3F'F499b/*11!45 -./ +i 	#Yt':'::i %%'7	'BCc		 	 b	)C ((*&+###Ws8}4
&*#+ 
/1A	AA7 Js   L5%L5c                  b    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y)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                8    t        j                  |i || _        y r6   )r  OrderedDictr0  )r   r}   r   s      r;   r  zTokenSet.__init__  s    **D;F;DLr=   c                ,    t        | j                        S r6   )rT   r0  r  s    r;   __len__zTokenSet.__len__  s    t||r=   c                     | j                   |   S r6   r0  )r   effects     r;   r@   zTokenSet.get  s    <<r=   c                d    |D cg c]  }t                }}t        t        ||            S c c}w )z?Creates a `TokenSet` corresponding to a list of `core.Effect`s.)r   r  r_  )clsr   r   tokenss       r;   r  zTokenSet.create  s0     '..ln.F.C()) /s   -c                H    t        | j                  j                               S r6   )r   r0  r  r  s    r;   r  zTokenSet.items  s    ##%&&r=   c                H    t        | j                  j                               S r6   )r   r0  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  @   K   | ]  }|j                   |   f  y wr6   r6  )r   effr   s     r;   r   z"TokenSet.subset.<locals>.<genexpr>  s     @S$,,s+,@r   )r  )r   r   s   ` r;   subsetzTokenSet.subset  s    @@@@r=   c                    g }| j                         D ]Q  }||j                  v r!|j                  ||j                  |   f       2|j                  || j                  |   f       S t        |      S )zRReturns a new `TokenSet` with tokens replaced with ones from the input `TokenSet`.)r   r0  r  r  )r   r:  
new_tokensr@  s       r;   update_tokenszTokenSet.update_tokens  sm    J||~ 4		3s 3453S 123	4
 Jr=   N)r7  zcore.Effectr   Token)r   Sequence[core.Effect]r   r  )r   z#Sequence[tuple[core.Effect, Token]])r   zset[core.Effect])r:  r  r   r  )r   r   r   r  r  r  r4  r@   classmethodr  r  r   rA  rD  rc   r=   r;   r  r    sG     76<  * *
'$A r=   r  )r  r  r  r  use_sharding_annotationsr  r  r  r  r  r  r  r  r  c               #   t        j                  d|       t        j                  j                  st        |       t        | j                  j                        }t        j                  dt        j                               g|z  }t        t        |      }|t        |j                        z   }|t        |      k(  s	J ||f       ||t        |      k(  s	J ||f       ||t        |      k(  s	J ||f       ||t        |      k(  s	J ||f       ||t        |      k(  s	J ||f       ||t        |      k(  s	J ||f       t        t        |      }t        t        |j                         }t        |      }|D cg c]  }t#                }}t        j$                  g|z  }||z   t'        |      z   }g |||}t        j$                  g|z  |j                   z   }g ||}|%dg||z   z  } g | |}|D !cg c]  }!|!dn|!|z    }}!|dg||z   z  }"g |"|}|	dg|z  }#g |#|	}	|dg||z   z  }$g |$|}|dg||z   z  }%g |%|}|dg|z  }&g |&|}|dg||z   z  }'g |'|}|dg|z  }(g |(|}|g dg||z   z  |}t)        |      })t)        |      }*t*        j,                  j/                  |)|*      }+|rdn|},t1        j2                  |,|+| j4                        }-t*        j6                  j/                  |rdnd      |-j8                  d	<   | j:                  j=                  |-       d}.|Nt        j>                  tA        |||      D !/0cg c]   \  }!}/}0tC        | |!|/      gtE        |0      z  " c}0}/}!      }.d}1|?t        j>                  tA        ||      D 20cg c]  \  }2}0|2gtE        |0      z   c}0}2      }1d}3|Mt        j>                  tA        |||      D 4!0cg c]  \  }4}!}0tG        |4|!      gtE        |0      z  ! c}0}!}4      }3d}5|?t        j>                  tA        ||      D 60cg c]  \  }6}0|6gtE        |0      z   c}0}6      }5d}7d}8|	t        j>                  tA        ||	|      D !/0cg c]   \  }!}/}0tC        | |!|/      gtE        |0      z  " c}0}/}!      }7t        j>                  tA        ||	|      D !/0cg c]  \  }!}/}0tI        |/|!      gtE        |0      z  ! c}0}/}!      }8d}9d}:||d
t        |      z  }g g }<};tA        |||      D ]e  \  }=}2}0|=!|2|;jK                  |=gtE        |0      z         n|;jK                  |2gtE        |0      z         |<jK                  |2gtE        |0      z         g t        j>                  |;      }9t        j>                  |<      }:d}>|Mt        j>                  tA        |||      D 4!0cg c]  \  }4}!}0tG        |4|!      gtE        |0      z  ! c}0}!}4      }>||.|1|3||5||dkD  s|dkD  s|dkD  rtM        t        |)            D cg c]  }i  }?}|wtA        ||      D @0cg c]  \  }@}0|@gtE        |0      z   }A}@}0tA        |?t        j>                  |A            D ]*  \  }B}@|@s	t*        jN                  j/                  d      Bd<   , |
rP|.NtA        |?|.      D ]?  \  }B}C|C	t        jP                  j                  rtS        C      Bd<   2tS        C      Bd<   A |19tA        |?|1      D ]*  \  }B}D|D	t*        j6                  j/                  D      Bd<   , |39tA        |?|3      D ]*  \  }B}E|E	t*        j6                  j/                  E      Bd<   , |59tA        |?|5      D ]*  \  }B}6|6s	t*        jN                  j/                  d      Bd<   , |t        jT                  t'        tM        t        |*                  t        tD        |            }Fg }GtA        ||      D ]:  \  }H}I|IGjW                  dgtE        H      z         'GjW                  FI          < tA        |?G      D ]  \  }B}I|I	tY        I      Bd<    |dkD  rNtA        | j                  j                  |?d|       D ]'  \  }J}Bt*        j6                  j/                  |J      |Bd<   ) n?| jZ                  j\                  r)|?D ]$  }Bt*        j6                  j/                  d      |Bd<   & |dkD  r1|?|||z    }K|KD ]$  }Bt*        jN                  j/                  d      |Bd<   & |dkD  r7|?||z   ||z   |z    }L|LD ]$  }Bt*        jN                  j/                  d      |Bd<   & t*        j^                  j/                  |?D Bcg c]!  }Bt*        j`                  j/                  |B      # c}B      |-_1        tM        t        |*            D cg c]  }i  }M}|dkD  r.Md| }N|ND ]$  }Bt*        jN                  j/                  d      |Bd<   & |rRM|d }Ot        |O      t        |      k(  r6tA        O|      D ]'  \  }B}Pt*        j6                  j/                  |P      |Bd<   ) |
r_|7]tA        M|7|8      D ]M  \  }B}C}Q|C
Qjd                  rt        jP                  j                  rtS        C      Bd<   @tS        C      Bd<   O |99tA        M|9      D ]*  \  }B}R|R	t*        j6                  j/                  R      Bd<   , |>9tA        M|>      D ]*  \  }B}E|E	t*        j6                  j/                  E      Bd<   , t*        j^                  j/                  MD Bcg c]!  }Bt*        j`                  j/                  |B      # c}B      |-_3        |rt*        jh                  jk                         g||z   z  }S|D ]P  }TSjK                  |Trt*        jh                  jm                  T      nt*        jh                  jk                                R |-jo                  S      }Un7t*        jh                  jk                         5  |-jo                         }Uddd       |rtq        jr                  |      ntq        jr                         }Vt+        jt                  U      5  Ujv                  }Wt        jx                  |W|||g      \  }X}}Y}t        jz                  |j|                        }Z|dk(  rZD [\cg c]  \  }[}\t        |[|\       }Y}[}\tA        ZY      D [\]ci c]  \  \  }[}\}]t        |[      |\f|] }^}\}[}]t        | Vddg dt        j                  g       ddX|^      }_|
s4|.2tA        W|.|      D !/`cg c]  \  }!}/}`|/|!nt        _|!`|/       }W}/}!}`|.|rtA        W|.||      D a/!bcg c]d  \  }a}/}!}b|!t        j$                  urIt        j                  |!j                  t        j                        r|/t        b|!      rt        _a|!      naf }W}!}/}a}bt        jx                  t        W|      |||g      \  }}c}}dt        tA        ||c            }e|d}ftA        |j                  |j|                  j                        D [gci c]&  \  }[}gt        |[      t        |[|gj                        ( }h}[}g|j                  D [cg c]  }[ht        |[          }i}[t        | |j|                  Ve|igfX^d\  }j}kg }l|D ]"  }mljK                  kj/                  |m             $ ljW                  j       t        |l      }n|
s4|72tA        n|7|      D a/ocg c]  \  }a}/}o|/ant        _ao|/       }n}/}a}o|7g }ptA        n|7||8      D ]  \  }a}/}o}Q|/Qjd                  rQj                  st        jP                  j                  r#t        |/oj                  j                        }/d}qn@oj                  j                  j                  rt        tM        oj                              nd}qpjK                  t        _ao|/q             pjK                  a        p}n|:3|r1tA        n|:|      D a2ocg c]  \  }a}2}o|2ant        a|2o       }n}2}a}o|7|rtA        n|7||	      D a/!bcg c]d  \  }a}/}!}b|!t        j$                  urIt        j                  |!j                  t        j                        r|/t        b|!      rt        _a|!      naf }n}!}/}a}bt1        j                  n       ddd       |-S c c}w c c}!w c c}0}/}!w c c}0}2w c c}0}!}4w c c}0}6w c c}0}/}!w c c}0}/}!w c c}0}!}4w c c}w c c}0}@w c c}Bw c c}w c c}Bw # 1 sw Y   8xY wc c}\}[w c c}]}\}[w c c}`}/}!w c c}b}!}/}aw c c}g}[w c c}[w c c}o}/}aw c c}o}2}aw c c}b}!}/}aw # 1 sw Y   |-S xY w)a	  Lowers jaxpr and its callees to an IR function.

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

  Note: this function does *not* take a name stack. Name stacks do not cross
  the boundaries of HLO functions.

  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.
    num_const_args: how many constant arguments is this function going to have.
      See https://docs.jax.dev/en/latest/internals/constants.html
    main_function: if true, this is the main function in the module. This has
      several effects:
      * the function's visibility is set to "public".
      * the function's symbol name will be "main"
      * the function's name will be used as the root name stack entry.
    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.
  Returns:
    MLIR func op
  r  rc   NFmainr  publicprivatesym_visibilityr6   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.tokenz	jax.constzjax.result_infor   )r  rH  rI  rG  r  r  r  rz   r
  r  r   r  r   unspecified_dims)Xr   r  r   r  r+  r  rT   r  r  r   ro   r   default_int_dtyperp   rn   r  r  
token_typer  r  r5  r%   FunctionTyper@   func_dialectFuncOpr  r\  r[  r  insertflattenr_  r  rC  r  r  r  r   BoolAttrr  get_sharding_attr	unflattenr1  r]   r  r  r   r  	arg_attrsr  result_attrsr>  r?  rE  add_entry_blockr   new_name_stackr  	arguments
split_listjaxpr_const_argsr  rf   r   r  r  r  wrap_with_sharding_opr  r   r  r  replicate_trailing_dimsrH  r  r  r   jaxpr_subcompr3  r0   rO  r  _any_axis_autor   r   wrap_with_memory_kindreturn_)rr{   rE  r  r   r  r  r  r  r  r  rH  r  r  r  r  r  r  r  r  r  num_dim_varsdim_var_avalsdim_var_typesnr_argsinput_typesr  
num_tokensr   token_typestoken_avalsinput_avalsoutput_avalsprefix_input_output_aliasesrU   prefix_shardingstoken_shardingsprefix_replicated_argsprefix_memory_kindstoken_memory_kindsprefix_layoutstoken_layoutsflat_input_typesflat_output_typesftype	func_namefunc_opir_arg_shardingsrm  typesir_arg_memory_kindsmkir_arg_layoutslir_donated_args
is_donatedir_result_shardingsunconstrained_variantsir_result_memory_kinds"custom_call_ir_result_memory_kindsr`  custom_call_respomir_result_layoutsr^  
replicatedreplicated_ir_argsr  rO  r{  r  
output_idsaliasesitypesaliasvar_nametoken_arg_attrsconst_arg_attrsr_  token_result_attrsnamed_result_attrsname_uvmem_kindarg_locsr?  entry_blockrH  	flat_argsr  const_arg_valuesconst_args_and_avalsr  r   c_argr   entry_lowering_ctxa_avalors
token_argsunflattened_argsr  r}   r  unique_constsconsts_for_constvarsout_valsrz   outsr@  flat_outputso_avaltemp_flat_outputsunconstrained_dimssr                                                                                                                     r;   r  r    s   t //&-		.	.	4	4%  S))223,##B(@(@(BCD|S-o}5-S00'	CM	!6GX#66	!		 Gs?/C$C o	C		'S-?"? m	?		3{+; ; k	;		!W4D0E%E   	E		gY7M'99MM	7 OX.+_eoo6,7|*'./!/+/$$%
2++d8n<+<-<+<<+%%&3eooE,.;..,%#'&L:,E"FP8P;OP &:; ! %&ID!"Z0 ; ; v
!:;7&77M!fz)O<<+;< #Wz(ABA.AAO!&L:$=>@,@/?@$*,E.E1DEV|j89N1N1[1KFZ'M6}6~6N!+5'\J%>? +)+ &k2&|4
//

.0A
B%%f4)	5SVV<')+):):hY*0'%&'"||]KH	J 	JQ5 #31
-	.e1D	D 	JK !,,.<	>R U#	# 	>? .\\[+F	H 	HQ5 A
	,u"5	5 	HIN /!ll"%&6"D	FZ U+	+ 	FGO !,,.>M	O 	OQ5 #31
-	.e1D	D 	OP "\\.>M	O 	OQ5 &a
+	,|E/B	B 	OP  '+$$'!(3/B+C!CrC68K*, 	9R	RZ

C5<../

B4,u--. bTL$778	9 "\\#.)-o)F&lK	M 	MQ5 A
	,u"5	5 	MN !		%		(		#		)		$			a			!	 #./00202I 02 ""?K@BCT:uZL<+>> B B"9dll;M.NO M
%79{{t7L%3
4M  $4$@ ,<= A/%**00$5h$?E.!%6x%@E/"A & #I/B C E
%"&(mm&7&7&D%"
#E !y.9 @-%&(mm&7&7&?%"
#@ ""9o> <
%&(kkood&;%"
#< '>>U3()*+S|-LNj"$g{,@A ,-&%=
..$,v"66
7
..E*
+	,
 i1 8,%(0%$
%8 a !5!5!>!>!*=L!9; C/(E')}}'8'8'B#$C 
	 	 	<	< =%')}}'8'8'<#$= A~!,|j/HIo" 3%[[__T2k3 !,";".";n"LNo" 3%[[__T2k3 ((-67E	79G
 -./11b11, 11 !^%kz2# 1;;??40eK1 %jk2
#l"330,? <,%#%==#4#4U#; < "5"A"<1D#9 ; ?x		b&?&?((.."3H"=%
#4X#>%
 ? '|-CD @x		$&MM$5$5h$? !@ "\+<= >v		$&MM$5$5f$= !> ))+78%r{{u8:' ##%&,*CDH KooQbkk&&q)BKK4G4G4IJK))(3K				 .++-k. 
 %%d+**,  
% e'%%I-1__L*n=.?*NA'00= *>?%a &ad3 ? ? !$$8:J K IQu 
AuN  -zTTt//"%$>%' $(8(D ")-={KM MaF y!34F6STU
UMi M # !,<k!.0 
 aAr t***9y-b!4 ""4a
; =>>i  *.&y+>	z>2*4&Az1& Wj12I+D %,,(=(=>As 	1{1388,,M  ;@,,GQM"Q%0GG(U[[*i'#'4B%'Hj D '
kk*..%&'KK$T*L#(;(G ",0C\RT TaF y!34F6STU
UTl T &!,0C".0FH &
!QMb77$$**0021foo6J6JKA!% ,2??+?+?+N+NE&++&'  
"
"#8 !VQ1$3 4 
"
"1
%&  'l *5-  #> NO OaV z!4QFC
COl O
 &= !/BL!13 
 aAr t***9y-b!4 ""4a
; =>>l  &Ke'N 
.{
 0;R	J	>	H	F	O	O2	M 02B| 	811D 9. .(?M
 HT6Oye'N 
.s  <AD18AD6%AD;$AE
)$AE8AE
?%AE$AE$AE#8	AE*AE/&AE5	AE:&AE?=AFAAG*AFAGAF
/AAG1AF
AG)A)AF%A1AG+AF-.AG?AF3A=AGAF8
.C?AGA-AF?
B	AGB$A)AGDAGFAFF=AGGAGc                R   || j                   }n.t        |      }t        |t        j                        sJ |       |}t        d|g| gdd      }dt        j                  j                  |      i}t        j                  j                  |      |j                  d<   |j                  S )Nannotate_device_placementTrJ   )result_typesoperandshas_side_effectapi_version_xla_buffer_placementmhlo.frontend_attributes)rL   rn   r7   r%   r   custom_callr\  r@   r  r[  result)r:   r{  aval_outresult_typer   op	dict_attrs          r;   ri  ri    s    &&K
(
#Cc277#(S(#K.k]C1F"&(9(9+(FG).0kkooi.H"--*+	r=   c                
   t        |t        j                        sJ t        j                  j
                  rst        j                  |      j                  }t        d t        |      D cg c]&  }t        j                  g ||j                  k        ( c}      }t        | |||      S t        | ||t        j                  j                         j!                         t#        t        |j                                    S c c}w )N)axesis_open)
mesh_shapedim_shardingsrR  )r7   r   ro   r   r  r+  r   r   r.   r   r   SdyDimre  xcHloSharding	replicater  r   )r{   r   r   physical_ndimr\   rm  s         r;   rf  rf    s     
D$**	++	+""((&&t,11M =)
 !!r1tyy=A
	A !c433 	3bnn..099;5+,. .
s   )+D 
zset[core.Primitive]r&  c                  d| j                   vsJ 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        t!        |
j                              }t        d |D              s	J |
|f       t        d |
j                  D              }t#        t$        j&                  j)                  |
j*                              }|j-                  |      }||
j.                  j0                  z   }t3        | |
j4                  ||
j.                  j6                        }t        j8                  |
j.                  j6                        5  |5  |
j:                  j<                  5  |
j4                  t>        v}|rWt3        | d||
j.                  j6                        }|5  tA        | |
|t        |      g|i |
jB                  \  }}ddd       nd}tE        | |
j4                  ||
j.                  j6                  |t        d	 |
jF                  D              |d|
j:                  ||
      }tI        |
j4                  |
j:                  |
j*                  |g|i |
jB                  \  }}|jJ                  }t              t        |
jF                        k(  s	J ||
f       |r|jM                        }ddd       ddd       ddd       t        ||
jF                         t        jN                  |
|	        t        fd|jF                  D              |fS # 1 sw Y   xY w# 1 sw Y   vxY w# 1 sw Y   zxY w# 1 sw Y   ~xY w)a  Lowers a jaxpr into MLIR, inlined into an existing function.

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

  consts_for_constvars: the constants corresponding to jaxpr.constvars.
  dim_var_values: the list of dimension variables values in the current
    IR function, in the order of ctx.shape_poly_state.dim_vars.
  const_lowering: the lowering for constants, by constant id.
    See https://docs.jax.dev/en/latest/internals/constants.html
  gpuc                    t        |       t        j                  u r"t        | j                  | j
                        S t        | t        j                        sJ |    S )Nr   )rL   r   Literalrf   r   r   r7   Var)r   r   envs    r;   readzjaxpr_subcomp.<locals>.read  sG    Aw$,,~AFFKK488$$$Vmr=   c                    |J t        |t        j                        r|}n;t        |      dk(  r"t	        j
                  dt        d       |d   }nt        |      }|| <   y )NrJ   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   rT   r  r  DeprecationWarningr   )r   nodewr  s      r;   writezjaxpr_subcomp.<locals>.write  sb    $!
a	Ta% 1		.
 G$KCFr=   c              3  2   K   | ]  }t        |        y wr6   r   r   s     r;   r   z jaxpr_subcomp.<locals>.<genexpr>  s     ,!]1,   c              3  2   K   | ]  }t        |        y wr6   r  r   s     r;   r   z jaxpr_subcomp.<locals>.<genexpr>  s     <!]1<r  c              3  2   K   | ]  }t        |        y wr6   r  r   s     r;   r   z jaxpr_subcomp.<locals>.<genexpr>  s     2A}Q2r  c              3  4   K   | ]  }|j                     y wr6   r   r   s     r;   r   z jaxpr_subcomp.<locals>.<genexpr>   s     0QVV0   Nc              3  4   K   | ]  }|j                     y wr6   r   r   s     r;   r   z jaxpr_subcomp.<locals>.<genexpr>  s     8qAFF8r  )r  rG  rH  rI  r  r  r  rz   r  r  r
  r   c              3  .   K   | ]  } |        y wr6   rc   )r   r   r  s     r;   r   z jaxpr_subcomp.<locals>.<genexpr>+  s     .1tAw.rh  )r   z	core.Atomr   r   )r   zcore.Varr  r   )(r  r   r7   r   	NameStackrL   rT   invarsr  r  r  r3   r   	last_usedeqnsr   rp   r  r  r	  	filter_inr   rA  r  rH  rM  rG  rI  user_contextr{   managerr&  _cached_loweringr  r  outvars_uncached_loweringrz   rD  clean_up_dead_vars)r{   r  rH  r:  r  r  r   r}   r  r  r  in_nodesr  r	  r  eqn_name_stackrJ  can_cache_lowering	out_nodesrz   r
  rule_ctx_inliner  r  s         `                @@r;   rg  rg    s   $ 
cmm	##	#" #%#	,t,	,2d2	,	<';<	< 	<	J 0 : :	;MT*=MM	;	Tc%,,'	'6%6	'	!	"c%//&:	: "
 !"	:	S!!**	+s>/B	B 4""N34	B	%"67	%t$nnU#)ZZ .1cSszz*+H222CS(OC20SZZ00H;66@@MNOo.I#//"<"<<N
!#s}}n"%//";";=C

'
'(A(A
B  2C  2
''// 2 --6
6 	%c4&)oo&?&?A 	'"23	5#8.#'#'::#'
)Z	' 	' &#--%oo//8CKK88I377)')	+ 0MM377CKK<Djj	7 ((
^s3;;//A)S1AA/	%%j1A 2  2  2D E3;;	*Ci0].1^ 
..	.	669	' 	' 2  2  2  2  2  2sU   5Q)8Q9Q	*Q2C*Q	Q$Q)Q
Q	QQQ&"Q))Q2	c           
        t        d |j                  D              }t        t        j                  j                  |j                              }t        |j                  |j                  |t        |j                        t        j                  |j                        t        | j                              }		 | j                  j!                  |	d      }
|
zt'        d |j(                        }t+        t-        t.        |j                  |j                  |j                        | |j                  |j                  |||fi |}
|
| j                  |	<   t        fd|D              }t        fdt1        |
j2                  |
j4                        D              }t7        ||z   |z   |z         }|
j8                  rDt;        j<                  |
j>                  |t@        jB                  jD                  jF                        }notI        jJ                  tM        |
jN                        t@        jP                  j!                  |
j>                  jR                  jT                        |      jV                  }tY        ||
jN                        }t[        j\                  |t_        |      g      \  }}|ta        t1        ||            fS # t"        $ r t%        d|        w xY w)aw  Lowers a jaxpr equation, using a cache.

  The jaxpr equation's lowering is emitted as an out-of-line MLIR function, and
  that function's construction is cached in the event that we see a similar
  equation. For each such equation we either inline the function body or emit
  an out-of-line call to it, depending on whether any of the lowering rules
  opted out of inlining.c              3  4   K   | ]  }|j                     y wr6   r   r   s     r;   r   z#_cached_lowering.<locals>.<genexpr>:  s     .a166.r  )rG  r  r  r   r  r  NzUnable to hash key: c                    | j                   S r6   r   )r   s    r;   r   z"_cached_lowering.<locals>.<lambda>J  s
    aff r=   c              3  @   K   | ]  }j                  |        y wr6   )r@   )r   r@  r  s     r;   r   z#_cached_lowering.<locals>.<genexpr>Q  s     Gs+Gr   c              3  @   K   | ]  \  }}t        ||         yw)r   N)rf   )r   r  r   r   s      r;   r   z#_cached_lowering.<locals>.<genexpr>R  s)      
!T !N>>r   )1r   r  r  r  r	  r  r   r  rG  r{   	frozensetr   
FrozenDictr  r  r  r@   r   rv  rp   r  _emit_lowering_rule_as_funr   r  r_  r  r  r3  r  r"   inlined_func_callr'   r%   r  currentblockrW  CallOpr5  r  FlatSymbolRefAttrrQ  r+  resultsrH  r   rc  rT   r  )r{   r  r  r  r   r}   r  r  r	  	cache_keycache_entryr  tokens_in_argsr  r  r  
token_outss     ` `            r;   r  r  -  sX    .3::..(44>>s{{KL/gg$##CJJ/cmm$)
$$((D9K $ckk2I,"CMM377CKKH# K %0Cy!GGG. //1L1LM  
~%(884?
A$))$ 1 1 9 9 ? ?AD 112
  !1!1!:!:!@!@A g	 	
 -T;3K3KL)//)c/6J5KL*i	HS*=>	>>; 
 
	
 #&	
s   :J; ;Kc                f   t        |j                  j                        }t        t	        j
                  dt        j                                     g|z  }	t        j                  t	        j                  |            \  }
}t        t        ||z         }t        t        |      }|D cg c]  }t                }}g |	||}g ||}t        |      }t        |      }t        j                  j!                  ||      }t#        j$                  |j&                  ||j(                        }t        j*                  j!                  d      |j,                  d<   |j.                  j1                  |      j2                   |j5                         }t        j6                  |      5  t9        |j:                  |      }t        j<                  ||t        |      t        |
      g      \  }}}}t?        |
||      D ci c]  \  }}}tA        |      |f| }}}}tC        ||tE        jF                         d||tI        t?        ||            d|||      }t        jJ                  j'                  tM        |j&                              5   | |g|i |\  }}ddd       |jN                  r-g |D  cg c]  } |jN                  j!                  |        c} }tQ              }t#        jR                  |       ddd       tU        |||
|      S c c}w c c}}}w # 1 sw Y   xY wc c} w # 1 sw Y   5xY w)z<Emits the contents of a lowering rule as a private function.rc   rK  rM  rN  N)r  rG  rH  rI  r  r  r  rz   r  r  r   )+rT   r  r  rn   r   ro   r   rT  r   unzip2eqn_params_const_argsrp   rU  r5  r%   rV  r@   rW  rX  rE  r  r\  r[  r  rY  r+  r`  r  rH  rb  rc  r_  r   r  r   ra  r  r>  rF  rz   r3  rj  r  )!lowering_ruler{   r  rG  r	  r  r  r  rk  rm  r  r  ro  r  r   rq  r}  r~  r  r  r  r  r  r  r  r  r   r  r   sub_ctxr  r  r@  s!                                    r;   r  r  f  s    S))223, D$$R)A)A)CDE- !%D,F,Fv,N O*oO_x%?@+_i0,'67!7+7<-<+<<+.;..,%k2&|4
//

.0A
B%	#&66+')+):):9)E'%&'"(('')+	% 5[* oo&#S%93z?KM CNJ 02B
 "*o?OP AtU 
AuN  "i#224Y3
;<w~%'G 
		#inn-	. I"7H-=HHldFIOG""&&s+GO$OdT"D/0 
G\:"
$ $I 8$I I H+ sJ   'L
 AL'8L
A1L'LL'+"L"&L'L'L	L''L0c                ~    | j                   j                  y | j                   j                  D ]  \  }}||u s|c S  y r6   )r  r  )r{   rG  r  r!  s       r;   _get_override_lowering_ruler
    sI     	44<((@@ gaA~k 
r=   c                
   d}t        |j                  |       }i }d }	||}	t        |	t              rJ t	        |      xs |j                  j
                  D ]9  }
| t        |
   v st        |
   |    }|j                  ||
<   |xr |j                  }; | t        v r7t        |    }|j                  }	t        |	t              rJ |xr |j                  }t        |	t              rJ t        d |j                         D              rJ t        |t        |       ||	|g|i |}	 t        |      }|j"                  j%                         r|j&                  }|t!        d|  d| d      |j%                         |j"                  j%                         k7  rLt!        d|  dt        |j"                  j%                                d	t        |j%                                      ||fS # t        $ r}t!        d|  d|       |d }~ww xY w)
NTc              3  <   K   | ]  }t        |t                y wr6   )r7   r   r  s     r;   r   z%_uncached_lowering.<locals>.<genexpr>  s     SaA01Ss   -Output of translation rule must be iterable: , got output zLowering rule for `z4` needs to set `tokens_out` because it has effects: rh  z4` returns incorrect set of output tokens. Expected: z vs. Actual: )r
  r  r7   r   _platforms_for_eqn_ctxr  r#  r!  r  r"  r  valueslower_per_platformrF  r   r   r   r  r   rz   )rG  r  r   r{   r}   r  r  override_ruleplatform_rulesdefault_ruler  r  ansretsr  rz   s                   r;   r  r    s;    &-c.@.@)L-,..&*, L,(9::: $G,L0B0B0L0L %	215	5(+I6FFq$AHH	% J
Y
aVVlL*;<<<"!((f&788	8S>;P;P;RSSS	S3I"	5%)	5-3	5#?:D
 	]] J	{ +%%,IQ01 1 s}}4466	{ +  %cmm&;&;&= >? @J&&()*, 
 
v% 
 ?
D"=7 8=>??s   "G" "	H+G==Hc                J    | y| j                   dk(  ry| j                   dk(  ryy)zFReturns platforms to override based on compute type of jaxpr equation.rc   device_host)ru  tpu_sparsecore)rx  )compute_type)r  s    r;   r  r    s2     _]*--	r=   c                    t        t        | j                        xs$ | j                  xs | j                  j                        S )z*The lowering platforms for the current eqn)r   r  r  r  r  )r{   s    r;   _platforms_for_eqnr    s=    	%c&7&78 =}}= # 2 2 < <
> >r=   c           
     :
    t               }t        |      dk(  r+|j                  |d   |      }|t        d| d|d          g }	i }
|j	                         D ])  \  }}||vrt        |	      |
|<   |	j                  |       + |D cg c]  }||
vr|
 }}|r9|t        d| d|       |D ]  }t        |	      |
|<    |	j                  |       |	sJ t        |	      dk(  rO |	d    g|i |}t         fdt        t        t        |                   t         fdt        |             |S t        |      dkD  rt        |	      d	k\  s	J ||	f       t         j                        dk\  sJ d
        j                  d   }t        t        j                  dt        j                              }|j                   |k7  rt#        j$                  ||      }t#        j&                  |g|t        |            }t)        |      D ]  \  }}|j*                  |   j,                  j                         }t/        j0                  |      5  t#        j2                  t5        t        j                  |
|               g       ddd        t6        j8                  j;                  |      }t        j<                  gt        |      z   j>                  z   }tA        t        |      }t#        j&                  tC        |      |t        |	            }t)        |	      D ]8  \  }}|
j	                         D cg c]  \  }}||k(  r| }}} jE                  |      }|j*                  |   j,                  j                         }t/        j0                  |      5   ||g|i |}	 t        |      }t         fdt        t        |             t         fd|       |jJ                  Ot        |      t        |jJ                        k(  sJ |D cg c]  }|jJ                  j                  |       c}|z   }t#        j2                  |       ddd       ; |jL                  }|rltO        jP                  tS        ||      t        |      g      \  }} jT                  jW                  tY        t[        ||                  }  j]                  |        |S c c}w # 1 sw Y   xY wc c}}w # tF        $ r}tI        d| d|       |d}~ww xY wc c}w # 1 sw Y   xY w)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.
  rJ   r   Nz%MLIR translation rule for primitive 'z' not found for platform z' not found for platforms c                0    t        | j                        S r6   wrap_compute_type_in_placeownerr  r{   s    r;   r   z$lower_per_platform.<locals>.<lambda>B	      ,S!'': r=   c                0    t        | j                        S r6   wrap_xla_metadata_in_placer!  r"  s    r;   r   z$lower_per_platform.<locals>.<lambda>F	  r#  r=   r   z#Must have a platform_index variablerc   r   )rZ  num_branchesr  r  r  c                0    t        | j                        S r6   r  r"  s    r;   r   z$lower_per_platform.<locals>.<lambda>o	      .sAGG< r=   c                0    t        | j                        S r6   r%  r"  s    r;   r   z$lower_per_platform.<locals>.<lambda>s	  r*  r=   )/r  rT   r@   r  r  r  r3   filterr<   r3  r  rn   r   ro   rA   rg   rL   r(   rh   CaseOpr"  regionsblocksr%   r  rj  rf   r  r	  r  r  r  rp   r5  rx   r   r   rz   r  r   rc  rH  r  rD  r  r_  ry   )!r{   descriptionr  r  r   	rule_argsrule_kwargsr  r!  
kept_rulesplatform_to_kept_rules_idxr  pruleplatforms_without_specific_rulerx  current_platform_idxrk   rule_idx_opr\   branchr	  rule_out_avalsr  case_oprule_idxplatforms_for_this_rule	inner_ctxr  r  r@  r  r:  rz   s!   `                                r;   r  r    sS   R 04)^qilL9D|
/} ='l^	-. .
 $&*/1 &&( ha	$'
Oq!e	 1: %M1()1K(K &' %M! %M$
/} =>?	AB B - 6&)*o #6l#	_Z]3::k:F:%'8'@A :&! M	Y!	J1 4My*6MM	4	S	 A	%L'LL	% ++A.T--bAB((*;;x1EF

H:!5(+I8+ 	" Jda  #**113F			6	" J	kk;rxx(B1(EFGHIJ JJ  //99'B/''(3+??#--O._n5,JJ'5'#&z?4' :& ga2L2R2R2T1#.1h"*a-  ! 1 1 &=>I__Q&&--/F			6	" I9	9[9fF%f-	 
<
'
3 
<
 
			)?#s9+?+?'@@@@ /1 ))--c2 13<=		kk)' 6 OO'oo$Wl;
?OFG ,,Xc/:@7B .C DJz"	.e%MNJ J1  FH&-}VH> ?DE	FF1# sU   
S
6S2S"T"S(-AT"T.TS	(	T1TTTT	c           	         t        | |      D ci c]  \  }}t        |      t        ||       }}}| D cg c]  }|t        |          c}S c c}}w c c}w )Nr   )r_  r   rf   )r  avalsr  r   uniq_constss        r;   	ir_constsrB  	  s]    69&%6H+21dbe[&&+  '-	-+be
	-- 
.s
   !AA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`.c                    rnfd}t        j                  ||t        j                  d|i             }t	        j
                  || j                        \  }}}t        d |j                  D              rEddl	m
} t        j                  ||      }	|j                  |	      }	|	j                  |	j                  }}| j                   '| j"                  j%                  | j                         }
n| j"                  }
t'        |
|| j(                  | j*                  t-        ||j.                  D cg c]  }|j0                   c}      g|| j2                  | j4                  d\  }}| j7                  |       |S c c}w )	Nc                      | i |fS r6   rc   )r}   r  funs     r;   r   z.lower_fun.<locals>.f_lowered.<locals>.<lambda>	  s    #t:Jr:J9L r=   r]  )
debug_infoc              3  P   K   | ]  }t        |t        j                           y wr6   )r7   r   InternalMutableArrayEffect)r   r  s     r;   r   z/lower_fun.<locals>.f_lowered.<locals>.<genexpr>	  s     
Qa:a889
Qr   r   )pxlar(  rQ  )lu	wrap_initr   rG  petrace_to_jaxpr_dynamicr  r  r   jax._src.interpretersrJ  r   ClosedJaxpr_discharge_internal_refsr  r  r  r  rx   rg  rH  r  rB  r  r   r  r   ry   )r{   r}   r  fwrapped_funr  r   r  rJ  r  sub_contextr   r   r:  rF  rV  s                 r;   	f_loweredzlower_fun.<locals>.f_lowered	  sV   %LA,,q&&&{CrBDK &(%>%>S\\&#"E1" 
Q5==
QQ,%%e-ABl22<@l$0$6$68K8K!e }} &&...Gk&&kUCNNCMM&(IA(IJ+ 
+ ))))+KC vJ )Js   .E;r{   r  rc   )rF  rV  rU  s   `` r;   r]  r]  	  s    
: 
r=   c                4   |t        |j                        z   t        |      k(  sJ |j                  s8||cxu r1n n.||j                  t	        |      f}	 | j
                  |   \  }	}
}
|	S t        | |||||||      }	|	S # t        $ r t        | j                        }t        | |||||||      }	t        | j                        |kD  }|rd| j                  vr:|	|	j                  j                  |	j                  j                  f| j
                  |<   Y |	S w xY w)N)r  r  r  r  rx  )rT   r  r  r  r   r  r   r  r  r  rE  r+  rL   r  )r{   fn_name
call_jaxprr  r   r  r  r  r  r  r   num_callbackshas_callbackss                r;   _lower_jaxpr_to_fun_cachedr\  	  s-    
#j112	2c(m	CC	C			yL@J$$eGn
5C`44S9mgq!" 
.	 !Wj'%,8G 
.!  
`#,,-m"
w
GNy|Mg #,,-=me3==8.5w||7I7I7<<K_K_._&&s+ 
.!
`s   B   BDDc                    | y |^}}|rt        d      | |k7  r*|t        j                  |       vrt        d| d|  d      y y )Nz>Multi-platform lowering when a backend= parameter is specifiedz Outer-jit backend specification z5 must match explicit inner-jit backend specification rh  )r  r  r,  r   )inner_backendr  outer_backendmore_lowering_platformss       r;   check_backend_matchesra  	  sx     
,>)-)
HJ J}$233MBB

*=/ :++8/	<= = C %r=   c
           
     8   t        |t        j                        sJ t        |             t	        ||j
                         t        |j                               }
t        t        |      }t               gt        |
      z  |z   }t        || |||
|||	      }|||
fS )N)r  r  r  )r7   r   rP  rL   ra  r  r  r   rp   rn   rU  rT   r\  )rX  rY  r{   r  r  r  r  r  r  r  r   r  r  s                r;   lower_called_computationrc  	  s     
J 0 0	1C4
3CC	1/""$%'_i0,,#g,.=,&	7J(6' 
,	''r=   )r  r  r[  c                  t        |t        j                        sJ t        |             t        j                  |j
                        }t        j                  |      \  }}|D cg c]  \  }}t        |||       }}}t        |      |z   }|	dgt        |      z  |	z   }	g ||}t        | ||t        |      |||||	|

      \  }}}|j                  j                  }t        |      }|D cg c]  }|j                  |       }}g |||}t!        j"                  |t$        j&                  j                  |      t)        |            }|r6t$        j*                  j                  |      |j,                  j.                  d<   t1        |j2                  |      }t        j4                  |t        |      g      \  }}|j7                  t9        t;        ||                  }||fS c c}}w c c}w )Nr   r(  )r  r  r  r  )r7   r   rP  rL   rd  r  r   r  rf   r   rT   rc  rE  r+  r5  r@   rW  r  r%   r  r3  r  rZ  r[  rH  r  rc  rD  r  r_  )rX  rY  r  r{   r  r  r  r  r   r  r  r[  r}   r  r  const_avalsr  r   r  r  r  r   symbol_namer~  r@  r:  callr  rz   s                                r;   call_loweringrh  	  s    
J 0 0	1C4
3CC	1..z/?/?@ KK(<=*k%9;!!T "!NN ; ;		 4	'$s:&2I&{&X&(#;z3J9$G '< ""+&|4*123IMM#2&2	*>	*F	*T	*$			.1155kB.t4
6$ <>KKOOJ<WDNN89,T\\<H)ooi#g,@&)&&xGV0D'EF*	J	/; 3s   'G+2G1)r  c          	     0   t        |t        j                        rt        j                  |      }t        |||| j                  | j                  | j                  | j                  g|| j                  | j                  d\  }}| j                  |       |S )NrQ  )r7   r   JaxprrM  close_jaxprrh  r  r  r  r  r  r   ry   )r{   rE  r  rY  r}   r  r:  s          r;   core_call_loweringrl  
  s     
DJJ'
+J#
J!3!3	llCMM3==)37) ''''	))V
 V	r=   	core_callrK  closed_call)r.  c                D    | dk(  ry| dk(  ry| dk(  ryt        d|  d      )	Nr  hostdevicedenser  sparseoffloadzInvalid compute type zK. Current supported values are `device_host`, `device` and `tpu_sparsecore`r   )c_types    r;   map_compute_typeru  #
  sF    }!!*6( 3F F 	G Gr=   c                   | j                   P| j                   j                  8| j                   j                  j                  d      r| j                   j                  j                  d      d   }t        j
                  j                  |      t        j
                  j                  d      d}t        j                  j                  |      |j                  j                  d<   y dt        j
                  j                  t        | j                   j                              i}t        j                  j                  |      |j                  j                  d<   y y y )Nzgpu_stream:r<  rJ   ro  )_xla_stream_annotation
inlineabler  _xla_compute_type)r  r  
startswithrW  r%   r\  r@   r  rZ  r[  ru  )r{   r  streamr  s       r;   r   r   -
  s   "s'8'8'E'E'Q
%%00?  --33C8;f"$--"3"3F";mm''0i =?KKOOI<Vbll89&(9(9
3,,99
:)< =i<>KKOOI<Vbll89 (R"r=   c                <   | j                   | j                   j                  ri i }}| j                   j                  j                         D ]>  \  }}t        j                  j                  t        |      j                               ||<   @ t        |t        j                        rt|j                  D ]5  }|dk(  s	|j                  |   D ]  }|j                  ||j                  <    7 t        j                  j                  ||z        |j                  d<   y y y y )Nr  )r  xla_metadatar  r%   r\  r@   rF  ru  r7   	Operationr[  r   rE  r  )r{   r  ctx_attributesexisting_attributesr  r   r   rU   s           r;   r&  r&  <
  s    "s'8'8'E'E*,b'N!!..446 <1--++CFLLN;nQ<"bll#-- 1$--==& 1a*+&&'11 35++//
.
.3bmm./ $	 (F"r=   c                  t        j                  |j                  t         j                        rt	        j
                  |j                        j                  }t        t        |            D cg c]  }|j                  |z    }}g ||}t	        j                  |      }t        | |||      S t	        j                  |j                        sAt        | |j                        }t        j                  t!        |      ||t#        |            }	nLt%        d |j                  D              sJ |       t        j                  t!        |      |t#        |            }	t'        | |	j(                         |	S c c}w )Nbroadcast_dimensionsc              3  ^   K   | ]%  }|t         j                  j                         k7   ' y wr6   )r%   r   r   )r   rj   s     r;   r   z#broadcast_in_dim.<locals>.<genexpr>a
  s*      * bmm4466 *s   +-)r   r  r   r  r   physical_element_avalrM   r   rT   r   r   r   r   rr  r(   dynamic_broadcast_in_dimrn   r   r   r   r!  )
r{   r  r  r  	elt_shaper\   trailing_dimsphysical_aval_outrM   r   s
             r;   r   r   L
  sL   
 x~~v7**8>>:@@I05c)n0EF1X]]Q&FMFB1BMB**84R"9MO O !!(..1*3?e((
(
#R

.
/c  *"..* * 4+34 *  
(
#R
.
/1c sCII.J) Gs   .E4c           	        g }t        ||      D ]  \  }}|j                  }|j                  }	t        j                  ||j
                  |      }
t        j                  ||      r$|j                  |	|k(  r|nt        | ||
             zt        |      t        |      k  s	J ||f       t        t        t        |      t        |      z
  t        |                  }t        | ||
|      }t        | ||
      }|j                  |        |S )z)Broadcasts multiple ops to the out_shape.rN  r  )r_  rM   rO  r   ro   r   definitely_equal_shaper  lower_with_sharding_in_typesrT   r  r   r   )r{   ops	ops_avals	out_shapeout_shardingr   r  op_avalop_aval_shapeop_aval_shardingout_avalr  b_outs                r;   multi_broadcast_in_dimr  i
  s    	#i( kb'MMM''7==<9H""=)<	jj'<7-c2x@B 3y>1MM93MM1!%I]9K(KSQZ^"\]
r82FHe*3x@e	jj 
*r=   c                   t        j                  |      }t        j                  |j                        s6t	        | |j                        }t        j                  t        |      ||      S t        j                  t        |      |      S r6   )	r   r   r   rM   rr  r(   dynamic_reshapern   ri   )r{   r  r  rM   s       r;   ri   ri   
  sk    )(				/(hnn=E!2u  ;;x0"55r=   c               ~   t        j                  |j                  t         j                        rt	        j
                  |j                        j                  }dgt        |      z  }dgt        |      z  }g ||}g ||}g ||}t	        j                  |      }	t        | ||	|||      S t        d |||fD              rFt        | |      }t        | |      }t        | |      }t        j                  t        |      ||||      S t        j                  |t!        |      t!        |      t!        |            S )Nr   rJ   )start_indiceslimit_indicesr   c              3  H   K   | ]  }t        j                  |         y wr6   r   r   r  s     r;   r   zslice_op.<locals>.<genexpr>
  s     
ZQt%%a((
Z    ")r   r  r   r  r   r  rM   rT   r   slice_opr  rr  r(   real_dynamic_slicern   r   r   )
r{   r:   r  r  r  r   r  trailing_zerostrailing_onesr  s
             r;   r  r  
  s<   x~~v7**8>>:@@IS3y>)NS3y>)M5m5n5M0m0i0M((-(G**84C-]"/B B 
Z=-QX2Y
ZZ23Fm23Fm,S':g##!	=-2 2 YYq&}5&}5&w/1 1r=   c                  | j                   d   }t        j                  |j                  t        j                        rt        j                  |j                        j                  }| j                   dd  }|r|d   j                  nt        j                  }t        t        j                  d|            gt        |      z  }g ||}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.                  ||t1        |	            S )Nr   rJ   )r  r   r  r   r  r   r  rM   rA   rg   rf   rO   rT   r   r   rr  r(   clamprv   subtractr  rn   r'  dynamic_slicer   )r{   r  r:   r  x_avalr  index_avalsr   r  slice_sizesclamped_starts              r;   r  r  
  s   <<?&x~~v7**8>>:@@I,,qr"K$/KN  RXXE!"((1e"456YGN5m5n5M!!(+H'F+				, /sK@KIIA3]++,=!	ll$S&,,7M !!!1{+aS3}--.	  Q{/KLLr=   c                  t        j                  |j                  t         j                        rt	        j
                  |j                        j                  }| j                  dd  }|r|d   j                  nt        j                  }t        t        j                  d|            gt        |      z  }g ||}t	        j                  |      }	t        | |	|||      S t        j                  |||      S )Nr   r   r   )r  )r   r  r   r  r   r  rM   r  rA   rg   rf   rO   rT   r   dynamic_update_slicer(   )
r{   r  r:   rQ  r  r  r  r   r   r  s
             r;   r  r  
  s    x~~v7**8>>:@@I,,qr"K$/KN  RXXE!5123c)nDE,m,e,M**84%66.;= = ##Av}==r=   c           	     "   t        d |||fD              r4t        j                  ||t        |      t        |      t        |            S t	        | |      }t	        | |      }t	        | |      }t        j
                  t        |      |||||      S )Nc              3  F   K   | ]  }t        j                  |        y wr6   r  r  s     r;   r   zpad.<locals>.<genexpr>
  s!      	Nq			" 	Ns   !)r   r(   padr   rr  dynamic_padrn   )r{   r  r:   padding_valuepadding_lowpadding_highpadding_interiors          r;   r  r  
  s     	 	N[-9;K-M 	N N771m";/"<0"#346 6
 /sK@K/\BL3C9IJ??!	=+|5EG Gr=   c                  t        j                  |j                        s?t        | |j                        }t	        j
                  t        |      |t        |            S t	        j                  t        |      t        |            S r6   )	r   r   rM   rr  r(   dynamic_iotarn   ra   iota)r{   r  	dimensionrM   s       r;   r  r  
  sg    				/(hnn=E!  88OH-x	/BCCr=   c                r    t        t        j                  ||j                              }t	        | ||d      S )zAReturns an IR constant shaped full of `value` shaped like `aval`.rc   r  )rf   rA   rO   r   r   )r{   r+  r   zeros       r;   full_like_avalr  
  s,    	RXXeTZZ0	1$	#tT	CCr=   c                   t        | j                  d   xt        j                        rDt	        j
                  j                  t        j                        r t        fd      | ||      S t        j                  ||      gS )Nr   c                h    j                   j                  j                  j                   | |      gS r6   )r   _rulesr'  )r:   rG  rU   s     r;   r   z&add_jaxvals_lowering.<locals>.<lambda>
  s%    177>>#5#5aggq!#D"E r=   )r7   r  r   ro   r   r  r   r  r]  r(   r'  )r{   r:   rG  rU   s      @r;   add_jaxvals_loweringr  
  sa    cll1o%t'7'781F9EFsAqQQ
''!Q-r=   c                    |gS r6   rc   )r{   r:   s     r;   r   r   
  s    1# r=   c                X   |Vt        j                  | j                        j                  }t	        |t         j
                        r|j                  rdnd}nd}t        j                  | |t        j                  j                  |      t        j                  j                  |            S )zCreates CompareOp.UNSIGNEDSIGNEDFLOAT)compare_type)r%   rr   rL   r   r7   rR   is_unsignedr(   compareComparisonDirectionAttrr@   ComparisonTypeAttr)r:   rG  	directioncomparison_type	elem_types        r;   compare_hlor     s    ##AFF+88I)R^^,&/&;&;
oo		!!%%i0))--o>	
@ @r=   c                   t        j                  |j                        }t         j                  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%   rr   rL   r   r7   r   r(   realr  imagselect)r  cmpr:   rG  tensor_typerxryreal_eqreal_cmpimag_cmpwhichs              r;   _minmax_hlor    s    ##AFF++^^{778	!B	!B"b$0G2r30H388A;S'BHJJw(3E::eQ""a8Or=   LTGTc                   t        j                  |j                  t         j                        s|j                  t	        j                  t        j
                        k(  r}t        j                  |j                  t        j                        rd}n3t        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   r  r   r  rA   rP   inexactsignedintegerr  r  r(   rh   rn   )r{   r:   aval_inr  r  s        r;   convert_hlor  !  s    
 

HNNFOO
<nn**

3l			7=="*:*:	;llA~c1g6lKA	_X.	22r=   c           	     d   t         j                  j                  r:|r8t        j                  j                  ||j                               j                  S |r8ddj                  t        |      D cg c]  }t        |       c}      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 c c}w )Nzunspecified_dims=[rO  ]r(  rJ   )r  r  backend_configr  result_shapesr  )r   r  r+  r$   r  ShardingConstraintOpbuildr  r+  r)  rF  rn   r7   r%   r   r   r   rM   r   rr  r  set_sharding)rE  r{   r:   r  rO  rS  r  allow_shardy_loweringr\   r  r  r  r  r  s                 r;   _wrap_with_spmd_opr  2  s    ""((-B<<,,Q0@AHHH )CHH 012AQ2-4 469:N N)+	K	)6;6	)  *00)	I&M1#yABM4{mqc"0 !.#2	4"
 r8	# 	3s   0D-
r+   )r  SPMDFullToShardShapeSPMDShardToFullShapec                B   |j                   j                  j                  r|S |j                   j                  j                  r|S |j                   j                  j                  r|S t        j                  |j                  t
        j                        rt        j                  |      }t        j                  j                  rW|%|j                   j                  |j                        n|}t!        ||j                   j                        }t#        | |||      S |3|j                   j%                  |j                        j'                         n|}d }|j                   j                  j(                  rt+        t-        |j                              }t#        | ||||      S r6   )rO  r  emptyare_all_axes_manualare_all_axes_autor   r  r   r  r   r   r   r  r+  r  r   r0   re  r  r  rh  r   r   )r{   r  r   sharding_protoprotorS  s         r;   r  r  [  sB   	]]I	]]++I	]]))I tzz6??3d#D""((& ]]++DII6,: 
.udmm6H6HIE b$66 % ]]//		:CCE+9 
}}((U499-. b$7GHHr=   c                    t        |t        t        f      rt        |      | j                  d<   y t        |      | j                  d<   y )NrO  rP  )r7   r.   r/   r\  r[  )r  rO  s     r;   r  r  u  s7    8\23$5h$?BMM.!%6x%@BMM/"r=   c                b   t        | t        t        f      r| j                         S t	        | j
                        dkD  r-t        j                  j                  | j                               S t        j                  j                  t        t        j                  j                  |                   S )Nd   )r7   r.   r/   r  rT   tile_assignment_devicesr%   r\  r@   SerializeToStringreprr  r  
from_protorN  s    r;   r\  r\  |  s|     8\23>>
 8++,s2]]x99;<<]]tBNN$=$=h$GHIIr=   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LayoutConstraintrJ   r  )r  r  r  r  operand_layoutsr  )rn   r7   r%   r   r   r   rM   r   rr  r  r  r   r   major_to_minorr  )	r{   r:   r  r  r  r  r  r  r  s	            r;   wrap_with_layout_opr     s    
  )+	K	)6;6	)  *00)	I&M1#yABM%[MQC !.$(w||)<$=#> $*#8#82#>"?A" 
r=   c                   | j                   |j                   k(  sJ t        j                  |j                        }|xs t        j                  | j                        }t	               }i }|j
                  j                  D ]p  }|j                  j                  }||v xs |dk(  }	|	s'|dk(  r|n|}
|
}d}||v s||v s||v r|
 d| }|dz  }||v r||v r||v r|||<   |j                  |       r t        j                  j                  d      }|j
                  j                  D ]@  }|j                  j                  }||v r|j                  |||          ||j                  d<   B |j                         D ]3  \  }}|j
                  j                  D ]  }|j                  |||        5 |j
                  j                  D ].  }| j
                  j!                  |       |j#                  |       0 |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.
  rJ  r   r   rJ   rM  rN  )r  r%   r  rZ  r   r  
operationsrQ  r+  r'  r\  r@   set_symbol_namer[  r  replace_all_symbol_usesr  rY  )
dst_modulerQ  
src_module
dst_symtab
src_symtab
used_names	renamingsr  rE  should_rename	base_namenew_namer\   rM  old_names                  r;   merge_mlir_modulesr    s   2 
		z11	11	1~~j223*AR^^J,@,@A*u* )OO&& b;;DJ&8$&.M"fn($ih
a
 #x:'=#[!%	Q #x:'=# !ionnX!& MMi('OO&& .b;;Dy  Yt_5&-BMM"#	. &OO- Ahoo(( A((8R@AA OO&& bOO2b 
6	r=      c               4   t         j                  j                  | j                        }|rt	        d| j                         t        || d| j                  |g dgt        | j                  j                        z  ||t        d      
      }|j                  S )zGHelper to generate pmap-style XLA computations for custom partitioners.r  r   F)r  )r  r  r  r	  r
  r  r  r  )r  r  r  r   r   r  r  rT   r  r  r  r`  )r  rE  r  r  r  r  lowering_results          r;   build_mlir_module_helperr    s     $55CC
89M9M8NO
PP)$$$r7S!3!3!:!:;;9,UKM/ 
		r=   r(  rc   )	r  r  r  called_computationsr  operand_output_aliasesr  r  extra_attributesc               	   t        |      }| t        j                  j                  d      }nt	        |t
        t        f      r t        j                  j                  |      }nIt	        |t              r"t        j                  j                  d      }d}nt        dt        |      z         t        t        j                  j                  |       t        j                  j                  |      |t        |      t        j                  j                  |D cg c]!  }t        j                  j                  |      # c}            }|xt        j                  j                  |j                         xs dD cg c]8  \  }}t        j                  j                  t!        |      dkD  r|gng |g       : c}}      |d<   ||j#                  |       |t        j$                  j                  t'        j(                  t        t+        t!        |      t!        |      t!        |      z               t&        j,                              |d	<   |	;t!        |	      t!        |      k(  s	J |	|f       t        |	      d
gt!        |      z  z   }	t        |      t        |      z   }|	t        j                  j                  |	D cg c]u  }t        j$                  j                  t'        j.                  t'        j(                  |t&        j,                              t        j0                  j                               w c}      |d<   |
|
J t!        |
      t!        |      k(  s	J |
|f       t        j                  j                  |
D cg c]u  }t        j$                  j                  t'        j.                  t'        j(                  |t&        j,                              t        j0                  j                               w c}      |d<   t        j2                  j5                  |||      }t	        |t              r6t        j6                  j                  |      |j8                  j:                  d<   |S c c}w c c}}w c c}w c c}w )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.
  r(  rJ   z,custom_call backend_config unexpected type: )call_target_namer  r  r  r  rc   )output_tuple_indicesoperand_indexoperand_tuple_indicesoutput_operand_aliasesr   indices_of_shape_operands)r   )rL   r  r  )r  r  r[  zmhlo.backend_config)r  r%   r\  r@   r7   rF  bytesdictr   r[  r]   r   r  r  r(   OutputOperandAliasrT   rQ  r?   rA   rB   r   rC   
atleast_1d	IndexTypeCustomCallOpbuild_genericr  rZ  r[  )r  r  r  r  r  r  r  r  r  r  r  r  backend_config_attrrE  r[  r.  
output_idxr  r  s                      r;   r  r    s   B (^(--++B/.3,/--++N;.$' --++B/K
Cc.FYY
ZZ}}(()9:kkooo6(;',,**6I
Jd2##D)
J* '+-<<+;+; %;$@$@$B$Hb	=  )Z 
   03</@1/D
|"! " ! 	= 	,J'( !&' /1.E.E.I.I


4c(mS]S=O-OPQ	#/$J*+ "!S]2O_h4OO2_-]9K0KKoH~] 33H $&LL$4$4 />6 *+ 	##MM"**Qbhh78!!# 	$ 	%6 %J !
 %%%~#l"33 &6& &3#%<<#3#3 /=5 *+ 	##MM"**Qbhh78!!# 	$ 	%5 $J  %%lX1; & ="%57[[__6BLL12	)g K	=665s   &Q5:=Q:
 A:R A:Rc                   t        |D cg c]  }t        |       c}      }t        d ||||
g|	D              rt        t        j                  dt
        j                              d fd}t        j                  t        t        ||	            t        d            }t        j                  j                  ||z   |      }t        j                  j!                   j"                  j$                  j&                        5  t)        j*                  ||      }ddd        j"                  j,                  j/                         |j1                         }t        j                  |      5  t        j2                   ||             ddd       t5        dt        t        t        |            g ||t7         |      t7         |      t7         |
      t7         |      ||j8                  j:                  g      }nt        j<                  t        t        t        |            ||t?        |      t?        |      t?        |
      t?        |      t        j@                  j                  t        jB                  |	t
        jD                        tG        |	      dg	      
      } |jH                  d   jJ                  jL                  ||z    }t        j                  |      5  t        j2                   ||             ddd       tO        |jP                   jR                        D cg c]  \  }}tU         ||       c}}S c c}w # 1 sw Y   4xY w# 1 sw Y   xY w# 1 sw Y   fxY wc c}}w )z9Builds a ReduceWindowOp, with support for dynamic shapes.c              3  H   K   | ]  }t        j                  |         y wr6   r  r  s     r;   r   z reduce_window.<locals>.<genexpr>  s(      	a ##A&	& 	ar  )rJ   r   c                H    t        |       }t        j                  |      S r6   )rr  r(   ri   )	pad_lo_hipadsr{   int2ds     r;   prep_one_padz#reduce_window.<locals>.prep_one_pad  s     )#y9d[[%%r=   r   Nzstablehlo.dynamic_reduce_window)r  r  r  r   )rM   )window_stridesbase_dilationswindow_dilationspadding)r*  z!tuple[core.DimSize, core.DimSize])+r5  rn   r  r   ro   rA   rg   r(   rs   r  rp   ra   r%   rV  r@   r  at_block_beginr  r`  r  rW  rX  r  rY  r`  rj  r  rr  rE  r+  ReduceWindowOpr   r?   rB   rC   rT   r.  r/  r  r_  r  r  r  )r{   reducer_namereducer_bodyr  init_valuesinit_values_avalsr  window_dimensionsr.  r1  base_dilationwindow_dilationr   scalar_typesr-  	d_paddingreducer_typereducerr  rwr  r,  s   `                    @r;   reduce_windowr@  t  s    "EV"WT?4#8"WX, 	a$o~}_W^_	a a D,,VRXX>?E& Sw%? @(1+NI??&&\!<1L				)	)#*<*<*C*C*H*H	I @##L,?g@##**73))+K			;	' -	kk,{+,- 
'#C$CD	$S*;< 	%S.9 	%S-8	
 	%S/: 	 $LL../
B 
		S),-+)*&~6&}5(9''++BJJw,I36w<2C , E
FB *bjjm""))L<,GIG			7	# )	kk,w'() RZZ7
9a 'sAt
4 
9 9[ #X@ @- -4) )
9s/   MM-M&M'-M3MM$'M0c                |   	 t        t        j                  j                  t	        |       dd      } |t
        j                  j                        }t               }|5  t        j                  j                  |      cddd       S # t        $ r}t        dt        | d      z         |d}~ww xY w# 1 sw Y   yxY w)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)mlir_moduleenable_shape_assertionsvalidate_static_shapes)enable_shardyzError refining shapes.  before_refine_polymorphic_shapesN)r   r!   mlirrefine_polymorphic_shapesr}  r   r  r+  r  r   rj  r  r%   r  r  )r`  rH  refined_module_strr  r  s        r;   rH  rH    s    
P '		(K(K*62$(#'!) 33399; ' /99??-./ / 
 P
!F$FG	HINOPP/ /s$   AB B2	B/B**B/2B;)r:   r   r   r  )r   zir.DenseIntElementsAttr)rD   Sequence[bool]r   zir.DenseElementsAttr)rt   z#Sequence[int | ir.RankedTensorType]r   zir.RankedTensorType)r   z!core.bint | np.dtype | np.genericr   ir.Type)r   core.ShapedArrayr   rK  )r   core.AbstractValuer   IrTypes)r   rM  r   ztuple[ir.Type, ...])r   rL   r   r   )r   rL   r   r   )r   r	   r   z5dict[tuple[int, core.AbstractValue], IrValues] | Noner   r   r   r   )r:   np.ndarray | np.genericr   r   )r   rO  r   r   r   r   )r   r   )r   z
core.Tokenr   r   )r   	type[Any]r   AttributeHandler)r   rP  r   rQ  )r   r	   r   ir.Attribute)r:   rO  r   rR  )r   rO  r   rR  )r   znp.dtype | np.genericr   rR  )r   zdict[str, Any]r   rR  )r   zSequence[Any]r   rR  )r.  rF  r/  r&  r   rF  )r{   r3  r7  zxc.Tracebackr   ir.Location)
r{   r3  rG  r  rH  r  rI  r  r   rS  )r`  r  ra  rF  r   
str | None)r`  r  ra  rF  r   rF  )rm  rF  r   rF  r6   )r`  r  r   rF  )r`  r  r   r  )r   r  )r  ztypes.CodeTyper   rT  )NTT)r-  r  r!  r  r  rT  r  r  r.  r  r   r  )rD   zIterable[IrValues]r   zlist[ir.Value])rD   zIterable[IrTypes]r   zlist[ir.Type])rD   zIterable[ir.Type]r>  zSequence[int]r   zlist[IrTypes])r:   rN  r   re   )rD   zIterable[ir.Value]rF  r  r   list[IrValues])rE  rF  r   rF  )r   rM  rO  JSharding | AUTO | Noner   rM  )r{   r  rM   
core.Shaper   ztuple[int | Value, ...])r{   r  rM   rW  r   ztuple[Value, ...])r{   r  rM   rW  r   r   )r  zsharding_impls.SPMDAxisContext)r{   r3  r   rM  rO  rV  r   zxc.OpSharding | SdyArray | None)r  zLayout | None | AutoLayoutr   rM  r   rT  )rm  rV  r   rT  )r   r  )r  core.ClosedJaxpr).rd  rF  r  rX  r  re   r  r  r	  zlist[core.Effect]r  r  r  r  r  r  r
  rJ  r  Sequence[bool] | Noner  (Sequence[JSharding | AUTO | None] | Noner  rZ  r  +Sequence[Layout | None | AutoLayout] | Noner  r[  r  r  r  r  r  re   r  re   r  r  r  zNone | tuple[int | None, ...]r  tuple[None | str, ...] | Noner  r  r   rt  )*r{   r3  rE  rF  r  rX  r   rF  r  re   r  r  r  rY  r  r  r  rZ  r  rZ  rH  r  r  zSequence[int | None] | Noner  rY  r  Sequence[str | None] | Noner  r  r  r]  r  r]  r  r[  r  r[  r  r\  r   r  )r:   ir.Valuer{  rF  r  rM  r   r^  )r   r^  r   r^  )r{   r3  r  z
core.JaxprrH  r  r:  r  r  zSequence[IrValues]r}   r   r  r  r   r	  r   #tuple[Sequence[IrValues], TokenSet])r{   r3  r  zcore.JaxprEqnr  r  r  ztuple[ir.Value, ...]r   r	  r   r_  )r  r  r{   r3  r  r  rG  r  r	  rF  r  r  r  r  r   r  )r{   r3  rG  r  r   LoweringRule | None)rG  r  r  r  r   r  r{   r  )r  r  r   r  )r{   r  r   r  )r{   r  r0  rF  r  zdict[str, LoweringRule]r  r`  r   r  r1  zir.Value | tuple[ir.Value, ...]r   r  )r@  r  r   rU  )T)rF  r   rV  r  r   r   )NN)r{   r3  rY  rX  r  re   )r^  rT  r  r  )NNN)rY  rX  r{   r3  r  re   )
rY  rX  r{   r3  r  r  r   r	  r[  zNone | dict[str, Any])r{   r  rY  zcore.ClosedJaxpr | core.Jaxpr)rt  rF  r   rF  )r{   r  r  ir.Operationr   r  )r{   r  r  rM  r   r^  )
r{   r  r  r  r  r  r  rW  r   r  )r{   r  r   r^  )r{   r  r  re   )r{   r  r   rL  r   r^  )r  rF  r  rT  rV  )NFF)rE  rF  r{   r  r:   r^  r  rM  rO  zxc.OpSharding | SdyArrayrS  zset[int] | Noner  r  r  r  )rO  'xc.OpSharding | SdyArray | SdyArrayList)rO  rb  r   rR  )
r{   r  r:   r^  r  rM  r  r    r  rM  )
r  r  rQ  rF  r  r  r  r  r   rF  )r  rX  rE  rF  r  r  r  r  r  r  r   r  )r  rF  r  zSequence[ir.Type]r  r  r  z%str | bytes | dict[str, ir.Attribute]r  r  r  zSequence[ir.Value] | Noner  r  r  re   r  zdict[int, int] | Noner  Sequence[Sequence[int]] | Noner  rc  r  zdict[str, ir.Attribute] | Noner   ra  )r{   r  r4  rF  r5  z(Callable[[ir.Block], Sequence[ir.Value]]r  r  r6  r  r7  r  r  r  )r`  r  r   r  (S  
__future__r   r  collections.abcr   r   r   r   r  r  r   r  rr  r  rY  r,  r  typingr	   r
   r   r   r   rq   r  jax._srcr   r   r   r   r   r   r  r   r   r   r   r   rK  r   r   r   r   r   r  rO  r   rM  jax._src.layoutr   r    jax._src.libr!   r"   r#   r  jax._src.lib.mlirr$   r%   r&   jax._src.lib.mlir.dialectsr'   rW  r(   jax._src.meshr)   jax._src.partition_specr*   jax._src.shardingr+   r  jax._src.sharding_implsr,   r-   r.   r/   r0   jax._src.state.typesr1   jax._src.typingr2   jax._src.utilr3   numpyrA   safe_maprp   
unsafe_mapsafe_zipr_  
unsafe_zipTypeVarr4   r   MYPYr   r   r<   rE   DenseI64ArrayAttrr@   r   rV   r]   ra   rv   r   r   rN  r   r   r   rR   rS   rP   int4int8int16rg   rC   uint4get_unsigneduint8uint16uint32uint64float8_e4m3b11fnuzFloat8E4M3B11FNUZTypefloat8_e4m3fnFloat8E4M3FNTypefloat8_e4m3fnuzFloat8E4M3FNUZTypefloat8_e5m2Float8E5M2Typefloat8_e5m2fnuzFloat8E5M2FNUZTypebfloat16BF16Typefloat16F16Typefloat32r   float64r   	complex64
complex128int2uint2float8_e3m4Float8E3M4Typefloat8_e4m3Float8E4M3Typefloat8_e8m0fnuFloat8E8M0FNUTypefloat4_e2m1fnFloat4E2M1FNTyper   r  r   r   r   r   rn   ro   rP  r   r   r   r   r   rf   r   r   maMaskedArrayr   ShapeDtypeStructr   ndarrayTypedNdArraylonglong_scalar_typer   python_scalar_typesptyper  rE  r%  rQ  r  r  r  r  r  r  HashableArrayr  r   r  python_scalar_types_to_dtypesr  rF  r\  r  r  r  r!  r  r  r2  r8  rM  DialectRegistryr  register_dialectsr  rY  rf  rj  r]  r_  r}  
ThreadPoolr  r  r  r  r  ReplicaAxisContextShardingContextr  r  	dataclassr  r  r&  r  r  r3  r  r  r   r"  r   r#  r/  r3  r5  objectr=  rA  rC  rH  compilerJ  rL  rS  rU  rl  rp  rr  rt  r  r  r  r  r  r  r  r  r  r  r  r  r  r   rU  r   r  r  ri  rf  r   r&  rg  r  r  r
  r  r  r  r  rB  r]  r\  ra  rc  rh  rl  call_pclosed_call_pru  r   r&  r   r  ri   r  r  r  r  r  r  r  add_jaxvals_pstop_gradient_pr  r  minimummin_hlomaximummax_hlor  r  re  wrap_with_full_to_shard_opwrap_with_shard_to_full_opr  r  r\  r   r  DEVICE_TO_DEVICE_TYPESEND_TO_HOST_TYPERECV_FROM_HOST_TYPEr  r  r@  rH  rc   r=   r;   <module>r     sw    #  B B     	   	   F F       +   #   &  # %  % 4 .  % ) 7 7 @ " 1 3I I - % !  --Z--ZFNN3  53//0-? &&**?
 O N,),( rww|,,
-7=
"((6==72>>#>#>B=
"((288gbnn991== "((6;;!<!<a@= "((277WR^^88!<	=
 "((288gbnn992>= "((288gbnn992>= "((288gbnn992>= "((6<<'".."="=qA= "((288gbnn991== "((299wr~~::B?= "((299wr~~::B?= "((299wr~~::B?= "((6$$%r'?'?'C'C= "((6 ""5"5"9"9= "((6!!"B$9$9$=$==  "((6 1 1 5 5!=" "((6!!"B$9$9$=$=#=$ "((6??R[[__
"((2::


"((2::


"((2::


"((2<<F
"((2==G
"((6;;!<!<a@
"((6<<'".."="=qA
"((6 1 1 5 5
"((6 1 1 5 5
"((6  !2#7#7#;#;
"((6 ""5"5"9"9;= 9 @KF
 NP J OP &5 !! "'D ## $5=h = 46 0 5*#
 KO$(@
@G@ 	"@ 	@@W "%%++-K L; $//>@ &$= &BJ &D "**&? @ (//1J KWWbhh"((XXryy"))RYYZZRZZ\\2==XXr{{FOO	= EL
 L*CDE:
 '' ;EE#9:; $**&= > SE2<</0 9; 6 ;+$GE	' 2::'E F >77JL WWbhh"((XXryy"))RYYZZRZZ\\2==XXr{{FOO	= KL
 \+IJK2 288%= > 2::'? @7 88>>@ ?LE5W5u=?? 3 1 1 2 5"--"3"3 4G 4!8 9: 4!< = 5"= > 2<< 5 277K 0
I	B	#8* # (38 'B&&(    0 1 #9??$ :7& #R]]_ 62:: 6, ""%%""$ > d#N N $N4P 	# 	# 	# d#  $ d#  $ sB sB sBl 1F 1F 1Fh ;X ; ,d#  $ 79
3 8P P6{66t<  CG(,#N *#N;?#N!%#N15#NL (
1'8=K  RZZ
+ +
F2F7IFB(B-DBG&0G5FGG	G%/G	 G6(267<6
+Z + D 4=	=
=(?= %=80+00:05
J O,Qf   .2>BAE>B?C&*)-!%:>>B/f.f.f. 	f.
 +f. 'f. f. f. f. !f. +f. <f. ?f. <f.  =!f." $#f.$ '%f.& 'f.( )f.* +f., 8-f.. </f.0 ,1f.2 3f.PcBJ 	]]
*  * f  -1>BAE%)8<.2-1)-487;?CBF>B+r	r
r r #	r r r +r +r <r ?r #r 6r ,r +r  '!r" 2#r$ 5%r& ='r( @)r* <+r, -rj!-?DL.0 03u , 4i78i7"i7 );i7 "	i7
 #5i7 #Qi7 ;i7V6?	6?*6?6? )6? C	6?
 <6?r:$:$	:$ ":$ 	:$
 +:$ +:$ ,:$ :$z	#166!6 !6 
	6r	 /	>M$'M'>M &9M !4	M
 $CM *<M^."N JN	-=8=.;=& BF!%()(0=((( !t6:!$! #5! #Q	! 4!H -1#@ $++w'9L M $$$,=A!#GW .6: 2&B '1 -?	0617?14M$,MB>+3>G8@G"	DD

 ''')= > '))+= >@ +s{{D
1
+s{{D
13, <@/45: / "  "4  ":	 
 *9  )-  /3 F   2J6:< $%79OP $%79OP I4AJ5JJ#"4 !' "4	: <@E!$E#,E $9E EHEP     " -0     	  #,	 0 =?!/3)+486:597;mm $m !	m
 :m m -m 'm m 2m 4m 3m 5m m`<9	<9 	<9 ;<9 !<9 $<9 4<9 ,<9~/r=   