
    ukiv                       d Z ddlmZ ddlZddlZddlmZmZ ddlZddl	Z	ddl
Z
ddlZddlmZmZ ddl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# e#jH                  Z$ ejL                  ddd      Z'd=dZ( ejR                  d      Z*de*_+         ejX                  e*       d Z-e-ej\                  e*<    G d de	j^                        Z0 G d de      Z1 G d d e	j^                        Z2 G d! d"e	j^                        Z3 ejh                  d#       G d$ d%             Z5d>d&Z6e*jn                  d'        Z8d?d(Z9	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d@d)Z: ejv                  e*e:d*+       dd,	 	 	 	 	 dAd-Z<dBd.Z=dCd/Z>dddddddddd0		 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dDd1Z?ddddddddd2	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dEd3Z@dddddddd4	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dFd5ZAdddddd6dde2j                  d7dddddddd8	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dGd9ZCdddddddddddd6dddddd:	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dHd;ZD	 	 	 	 	 	 	 	 	 	 	 	 	 	 dId<ZEy# e%$ r i Z$Y _w xY w)JzJAX bindings for Mosaic.    )annotationsN)CallableSequence)Any	TypedDict)api)config)core)dispatch)sharding_impls)is_cloud_tpu_older_than)
FrozenDict)batching)mlir)tpu)ir)PassManager)flagsjax_mosaic_allow_hloFzAllow hlo dialects in Mosaic)namedefaulthelpc                ~    | j                   j                  d      }| j                         s|t        ddd|      ryy )NT)optionali           )module_contextget_backendis_forward_compatr   )ctxbackends     S/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/_src/tpu_custom_call.pyget_ir_versionr$   B   sB    **D*9' 
		 r1g	6	    tpu_custom_callTc                    | j                   dk7  rt        d      t        d t        ||d      D              }t	        j
                  |i |}t        d |D              }|dt        |      z  fS )N   z6tpu_custom_call does not support non-trivial batching.c              3  X   K   | ]"  \  }}|t         j                  u s||n||    $ y wN)r   
not_mapped).0ads      r#   	<genexpr>z*tpu_custom_call_batcher.<locals>.<genexpr>X   s6      
!Q $$$	a!<s   (*T)strictc              3  &   K   | ]	  }|d      y wr*    )r,   os     r#   r/   z*tpu_custom_call_batcher.<locals>.<genexpr>]   s     -!ag-s   )r   )sizeNotImplementedErrortupleziptpu_custom_call_pbindlen)	axis_dataargsdimskwargsunbatched_argsout_unbatchedouts          r#   tpu_custom_call_batcherrB   S   sz    ^^q
@   dD. . $((.CFC--}--#	dSXo	r%   c                      e Zd Z ej                         Z ej                         Z ej                         Z ej                         Z ej                         Z	 ej                         Z
edd       Zy)MemorySpacec                    | t         j                  k(  ry| t         j                  k(  ry| t         j                  k(  ry| t         j                  k(  ry| t         j
                  k(  ry| t         j                  k(  ryt        dt        |       z         )Nr   r(      r      r   zinvalid memory space: )	rD   HBMVMEMSEMAPHORE_MEMSC_SCALAR_SEMAPHORE_MEMSMEMHOST
ValueErrorstrselfs    r#   colorzMemorySpace.colorj   s{    {	!!	!	**	*	44	4	!!	!	!!	!/#d);<<r%   N)returnint)__name__
__module____qualname__enumautorH   rI   rJ   rL   rM   rK   propertyrR   r2   r%   r#   rD   rD   b   sc    		#	$$))+-	$	$%DIIK= =r%   rD   c                  B    e Zd ZU ded<   ded<   ded<   dZded<   d	dZy)
CostEstimaterT   flopstranscendentalsbytes_accessedr   remote_bytes_transferredc           	     V    d| d    d| d    d| d    d| d    d		j                  d
      S )Nz
{"flops": r]   z, "transcendentals": r^   z, "bytes_accessed": r_   z, "remote_bytes_transferred": r`   }ascii)encoderP   s    r#   to_jsonzCostEstimate.to_json   sW    
d7m_ %"#$ %!"# $+,-R	1 fWor%   NrS   bytes)rU   rV   rW   __annotations__r`   re   r2   r%   r#   r\   r\   |   s#    	*"#C#r%   r\   c                      e Zd ZdZdZdZy)TpuSideEffectTypepuredataflow_side_effectingside_effectingN)rU   rV   rW   PUREDATAFLOW_SIDE_EFFECTINGSIDE_EFFECTINGr2   r%   r#   rj   rj      s    	$5#.r%   rj   c                      e Zd ZdZdZy)TilingTILING_COMPACTTILING_SPARSE_COREN)rU   rV   rW   COMPACTSPARSE_COREr2   r%   r#   rr   rr      s    '$+r%   rr   )frozenc                      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<   ded<   ded<   dZded<   d Zd  Zd"d!Zy)#CustomCallBackendConfigz;Represents an unserialized backend config for custom calls.rg   lowered_module_asmboolhas_communication
int | Nonecollective_id
str | Nonedevice_typeCostEstimate | Nonecost_estimateneeds_hlo_passesneeds_layout_passesvmem_limit_bytes$dict[str, bool | int | float] | Noner   Sequence[bool] | Noneallow_input_fusionserialization_formatinternal_scratch_in_bytes%tuple[MemorySpace | None, ...] | Noneoutput_memory_spacesdisable_bounds_checksactive_core_countinput_memory_spacesskip_device_barriershape_invariant_numericsNTiling | Nonetilingc                    | j                   *t        j                  | dt        | j                                | j                  +t        j                  | dt        | j                               y y )Nr   r   )r   object__setattr__r6   r   r   rP   s    r#   __post_init__z%CustomCallBackendConfig.__post_init__   s`    *3t6679%#D$6$679 &r%   c                     y)Nz"CustomCallBackendConfig(<omitted>)r2   rP   s    r#   __repr__z CustomCallBackendConfig.__repr__   s    /r%   c                   t        j                         } |j                  d        |j                  t        j                  | j
                                |j                  d       | j                  rT |j                  d        |j                  t        | j                        j                         j                  d             | j                  F |j                  d        |j                  t        | j                        j                  d             | j                  7 |j                  d        |j                  t        d/i | j                         | j                  rT |j                  d        |j                  t        | j                        j                         j                  d             | j                  T |j                  d        |j                  t        | j                        j                         j                  d             | j                  rT |j                  d	        |j                  t        | j                        j                         j                  d             | j                   sT |j                  d
        |j                  t        | j                         j                         j                  d             | j"                   |j                  d       t%        | j"                        D ]I  \  }} |j                  |rdnd       |dz   t'        | j"                        k7  s8 |j                  d       K  |j                  d       | j(                  F |j                  d        |j                  t        | j(                        j                  d             | j*                   |j                  d       t%        | j*                        D ]S  \  }}|r |j                  d       ||j,                  nd} |j                  t        |      j                  d             U  |j                  d       | j.                  d}t%        | j.                        D ]  \  }}|	|t0        j2                  u r|t0        j4                  t0        j6                  t0        j2                  fvrt9        d      |r |j                  d       n |j                  d        |j                  d| d|j,                   dj                  d             d} |r |j                  d       | j:                  rT |j                  d        |j                  t        | j:                        j                         j                  d             | j<                  rT |j                  d        |j                  t        | j<                        j                         j                  d              |j                  d       | j>                  B |j                  d        |j                  t        | j>                  j@                               | jB                  Q |j                  d         |j                  d!| jB                  jE                         z   d"z   j                  d             | jF                  X |j                  d#        |j                  t        | jF                        j                  d              |j                  d$       | jH                   |j                  d%       t%        | jH                  jK                               D ]y  \  }\  }} |j                  d&        |j                  |j                  d              |j                  d'       tM        |tN              r) |j                  d(        |j                  |rdnd       ntM        |tP              r= |j                  d)        |j                  t        |      j                  d             ndtM        |tR              r= |j                  d*        |j                  t        |      j                  d             ntU        d+t        |      z          |j                  d,       |dz   t'        | jH                        k7  sh |j                  d       |  |j                  d       | jB                  d-k(  r!| jV                  dk(  r |j                  d.        |j                  d        |jX                         S )0z(Serializes the backend config into JSON.s!   {"custom_call_config": {"body": "   "s   , "has_communication": rc   s   , "collective_id": s   , "cost_estimate": s   , "needs_hlo_passes": s   , "serialization_format": s   , "needs_layout_passes": s   , "shape_invariant_numerics": s   , "allow_input_fusion": [s   trues   falser(      ,   ]s   , "internal_scratch_in_bytes": s   , "output_memory_colors": [Fz:input_memory_space_colors only supports HBM, VMEM and SMEMs    , "input_memory_space_colors": [z{"operand_index":z	,"color":rb   Ts   , "disable_bounds_checks": s   , "skip_device_barrier":    }s   , "sparse_core_config": )r   s   , "device_type": z"DEVICE_TYPE_"sD   , "scoped_memory_configs": [{"memory_space":1, "offset": 0, "size": s   }]s   , "flag_configs": [s   {"flag_type": "s   ", "value": {s   "boolean_value": s   "integer_value": s   "double_value": zinvalid flag value: s   }}
sparsecores1   , "megachip_parallelism_config": {"cores": ["0"]}r2   )-ioBytesIOwritebase64	b64encoderz   r|   rO   lowerrd   r~   r   _compact_json_objectr   r   r   r   r   	enumerater:   r   r   rR   r   rD   rL   rH   rI   r5   r   r   r   valuer   upperr   r   items
isinstancer{   rT   floatrN   r   getvalue)	rQ   r	   ir   memory_spacerR   commainput_memory_spaceflags	            r#   re   zCustomCallBackendConfig.to_json   sw    ZZ\FFLL56FLL!!$"9"9:;FLLfll-.fll3t--.446==gFG%fll)*fll3t))*11':;%fll)*fll'=$*<*<=>fll,-fll3t,,-335<<WEF  ,fll01fll3t001779@@IJfll/0fll3t//0668??HI((fll45fll3t445;;=DDWMN*fll/0 7 78 (!UW84q5C//00
&,,t
 fll4%%1fll56fll3t556==gFG  ,fll12&t'@'@A 1/!\
&,,t
&2&>""BSZ&&w/0	1
 fll4+e#,T-E-E#F 
!%
!1!11 OO&
 

 $J  
&,,t

&,,:
; 9-?-E-E,FbIVG_	
 /0 
T!!fll12fll3t11288:AA'JKfll/0fll3t//0668??HIFLL{{fll./fll't{{/@/@AB#fll'(fllT--3355;
C
CG
L (fll fll3t,,-44W=>fll5zzfll)*'

(8(8(:; 
!]dE'(T[[)*%&eT"
&,,+
,
&,,%wX
6s#
&,,+
,
&,,s5z((1
2u%
&,,*
+
&,,s5z((1
21CJ>?
?Uq5C

O#
&,,t
#$ fll4<'D,B,Ba,GfllGHFLL6??r%   rf   )	rU   rV   rW   __doc__rh   r   r   r   re   r2   r%   r#   ry   ry      s    C$$	--++""''==<<  &-90xr%   ry   c                 R    t        j                  | ddd      j                  d      S )NTr   ),:)	sort_keysindent
separatorsrc   )jsondumpsrd   )r>   s    r#   r   r   8  s%    	Q:

F7Or%   c                    | S r*   r2   )	out_avals___s      r#   _tpu_custom_call_abstract_evalr   >  s    	r%   c           
     n    | D cg c]%  }t        t        |j                  dz
  dd            ' c}S c c}w )Nr(   r   )r6   rangendim)avalsr-   s     r#   _avals_to_layoutsr   C  s,    49	:q%affqj"b)
*	::	:s   *2c               @   |D cg c]  }t        j                  |       }	}| j                  j                  }
t	        |
t
        j                        rC|
j                  r|
j                  t        |
j                  j                        k7  rVt        d      t	        |
t
        j                        r|
j                  dk7  r"t        d      |j                  rt        d      t        d | j                   D              rd }nJ| j                   D cg c]5  }t        j"                  t        j$                  | |j&                              7 }}d }|)t)        t*        j,                  j/                  |            }t        j0                  d|	| |j2                         d|t4        j6                  k7  t)        |      t9        | j:                        t9        | j                         ||      }i }|+t*        j,                  j/                  t=        di |      |d<   t	        |t4              sJ |t4        j>                  k(  r"t*        j,                  j/                  d	      |d
<   |r,t*        j@                  j/                  |      |jB                  d<   |jD                  S c c}w c c}w )NzXMosaic kernels cannot be automatically partitioned. Please wrap the call in a shard_map.r(   z4Replica lowering for Mosaic kernels not implemented.c              3  Z   K   | ]#  }t        j                  |j                         % y wr*   )r
   is_constant_shapeshape)r,   aval_outs     r#   r/   z,_tpu_custom_call_lowering.<locals>.<genexpr>d  s     NH			/Ns   )+)kernel_namer&   )
result_typesoperandsbackend_configapi_versionhas_side_effectoperand_output_aliasesoperand_layoutsresult_layoutsresult_shapesextra_attributeskernel_metadatatruexla_allow_dce_side_effecting_opzmhlo.frontend_attributesr2   )#r   aval_to_ir_typer   axis_contextr   r   SPMDAxisContextmanual_axes	frozensetmesh
axis_namesr5   ShardingContextnum_devicesr|   all	avals_outshape_tensoreval_dynamic_shaper   dictr   
StringAttrgetcustom_callre   rj   rn   r   avals_inr   ro   DictAttr
attributesresults)r!   r	   has_side_effectsr   r   input_output_aliasesmetadatain_nodesavalr   r   r   r   r   callmetadata_dicts                   r#   _tpu_custom_call_loweringr   G  sO    :CC$&&t,C,C##00,n<<=    Il.?.?.J.J$KK"  , > >?1$"  
>  	NNNM ' 	$11#x~~FG'M '  (9(9+(FG			#V^^%&*;*@*@@!"67'5&s}}5!'
$ -')}}'8'8(x((M#$ 
$&7	88	8*BBB79}}7H7H7PM3424++//-2PDOO./	m D,'s   J:Jr   )platform
ir_versionc               Z   t        j                  | j                        \  }}| j                  5 }| j                  j                  5 }| j                  j                         }|j                  }d|_        |d| nd}	 t        j                  d|z   dz         }	|	j                  |       ||_        t        j                         }
|j                  |
d       |
j                         }|||ffcd d d        cd d d        S # ||_        w xY w# 1 sw Y   nxY wd d d        y # 1 sw Y   y xY w)NTztarget-version= z+builtin.module(mosaic-serde{serialize=true z})r   )desired_version)r   private_has_communication	operationcontextlocationcloneallow_unregistered_dialectsr   parserunr   r   write_bytecoder   )moduler   r|   has_custom_barrierr!   r   	module_op prev_allow_unregistered_dialectstarget_versionpipelinebytecode_bufferasms               r#   _lower_mosaic_module_to_asmr    s6   
 +.*G*G+'' ~~ f..77 1  &&(I'*'F'F$&*C#*4*@/*&b I""
7.
H4
Oh ll9(Hc%jjlO_a@

"
"
$C #   )Ic%    s<   D!7D>,D *AD-	D! 	D		DD	D!!D*c                    dddfd}| j                   j                  |t        j                  j                         rrt        d      ryy)z>Determines the device type based on the core_type annotations.Fc                   | j                   dk(  rd| j                  v r| j                  d   }t        |      dD cg c]  }d| d
 c}v r8drt        j                  j
                  S t        j                  j                  S t        |      dk(  rdt        j                  j                  S t        d|       t        j                  j                  S c c}w )	N	func.funcztpu.core_type)sc_scalar_subcoresc_vector_subcorez#tpu.core_type<>Tz#tpu.core_type<tc>zUnknown core type: )	r   r   rO   r   
WalkResult	INTERRUPTSKIPrN   ADVANCE)op	core_typecsparsecore_func_foundtensorcore_func_founds      r#   %assign_device_type_based_on_core_typez?_get_device_type.<locals>.assign_device_type_based_on_core_type  s     
ww+	BMM	)MM/2	y>?
 aS"
 
 #'
"==***##
#y>11"&
##
#.yk:;;==   
s   C)
walk_orderzOA single Mosaic kernel cannot contain both TensorCore and SparseCore functions.r   N)r  zir.OperationrS   zir.WalkResult)r   walkr   	WalkOrder	PRE_ORDERrN   )r  r  r  r  s     @@r#   _get_device_typer$    sc    !( 	+8N8N   4
	  	r%   c                   	 d	 	 	 	 	 	 	 d	d}d }| j                   j                  D ]  }|j                  j                  dk7  rd|j                  vsd|j                  vr:	 t        j                  |j                  d         }	 t        j                  |j                  d         } ||||      } |S # t        $ r}|j                  d        d }~ww xY w# t        $ r}|j                  d        d }~ww xY w)
Nc                J   t        |      t        |       k7  rt        d      d }t        t        ||             D ]f  \  }\  }}t	        |      dk7  rt
        j                  j                  |      rt        d| d      |t        d      |||k7  rt        d      |}h |S )Nz^The iteration bounds and dimension semantics attributes must have the same number of elements.z'#tpu.dimension_semantics<core_parallel>zAThe iteration bound corresponding to the core-parallel dimension z must be statically known.zKA single Mosaic subkernel cannot contain multiple core sharding dimensions.zcThe iteration bound corresponding to the core-parallel dimension be the same across all subkernels.)r:   rN   r   r7   rO   r   
ShapedTypeis_dynamic_size)dim_semanticsiter_boundsother_subkernel_core_dim_sizesubkernel_core_dim_sizedim_idxdim_sizedim_sems          r#   get_core_parallel_dim_sizez:_get_active_core_count.<locals>.get_core_parallel_dim_size  s    
 ;3}--* 
 #(1K') )$$(G 
WB	B		&	&x	0Oi13
 	
 
!	,
 	

 (
3+x71
 	
 !)1)4 #"r%   r  iteration_boundsdimension_semanticsz0The iteration bounds attribute must be an array.z3The dimension semantics attribute must be an array.)r)  r*  r+  r*   )r)  zir.ArrayAttrr*  zir.DenseI64ArrayAttrr+  r}   rS   r}   )
body
operationsr   r   r   r   DenseI64ArrayAttrrN   add_note	ArrayAttr)r  r0  core_parallel_dim_sizer  r*  er)  s          r#   _get_active_core_countr:    s"   
 37'#!'#''# &0'# <F'#R  KK"" b	||K' 	"--/ 5((7I)JKkll2==1F#GHm
 8#&<+6 
   jjCD
  jjFGs0   !"B5"C5	C>CC	C8!C33C8)	r   r   r   r   r   *allow_collective_id_without_custom_barrierr   r   r   c               
   t        |       }t        j                  }|| }t        | |	      \  }\  }}t	        |       }t        |fi d|d|d|d|d|d|d|d	|d
|d|d|d|d|d|
d|d|d|d|d|d|S )Nr   r   r   r   r   r   r~   r   r   r  r|   r   r   r   r   r   r   r   r;  r   r   )r$  _MOSAIC_ALLOW_HLOr   r  r:  _lowered_to_custom_call_config)r  r   r   r   r   r   r~   r   r   r   r   r   r   r;  r   r   r   r   r   rz   r|   r  r   s                          r#   _lower_to_custom_call_configr?    s/   ( !(+&,, )/ " 
 -V4	'
'
 "
 	

 ,
 !:
 "
 
 0
 ,
 *
 (
 .
 0
 2
  *!
" .#
$ .%
& 2\'
(  8)
* +
 r%   )r   r   r   r   r   r;  r   r   c                   |r|t        d      ||st        d      |(t        |t              st        dt        |       d      t	        | |	||||
|||||||||||||      S )Nz=collective_id has to be specified when using a custom barrierzKcollective_id has to be unspecified or None when not using a custom barrierz1vmem_limit_bytes must be an int: provided with a .)r   r   r   r   r   )rN   r   rT   typery   )rz   r   r   r   r   r   r~   r   r  r|   r   r   r   r   r   r   r   r   r;  r   r   s                        r#   r>  r>  X  s    0 
I   )S
	  !*5Es*K
	!"#1	&  
!)--7'
 r%   )r   r   r   r;  r   r   r   c          
        t        |t              r"|st        j                  nt        j                  }t        |fi d|d|d|d|d|	d|
d|d|d	t        |       d
|d|d|d|d|d|d|}t        | g|||||||dS )Nr   r   r   r   r   r~   r   r   r   r   r   r   r;  r   r   r   r	   r   r   r   r   r   )r   r{   rj   rn   rp   r?  r$   r   )r!   r  out_typer   r   r   r   r   r   r   r~   r   r   r   r   r   r   r   r;  r   r   r   r   r	   s                           r#   lower_module_to_custom_callrF    s   2  $'   	-- 
 (' " 	
 , !: " 0 0  $ 2 . . 2\  8  .!" #&& 
#		
	
 '/	
 	r%   r2   r(   )r   r   r   r   r   r   r   r~   r   r   r   r   r   r   r   r   _ir_versionc               V    t        | ||||||	|||||||      }t        ||
||||      S )z;Turns an MLIR Mosaic kernel into a JAX-compatible function.)r   r   r   r   r   r~   r   r   r   r   r   r   r   r   r   r   )r?  _as_jax_callable)r  rE  r   r   r   r   r   r   r   r~   r   r   r   r   r   r   r   r   rG  r	   s                       r#   as_tpu_kernelrK    sY    . ('!+ 9!//1-7-&  
/
 r%   )r~   r   r   r   r|   r   r  r   r   r   r   r   r   r   r   r   r;  c                  t        |       }| j                  j                  dd      }t        |t              r"|st
        j                  nt
        j                  }t        ||
|||||||||||||      }t        ||||	||      S )NT)binaryenable_debug_info)r   r   r   r   r   r~   r   r   r  r|   r   r   r   r;  rI  )
r$  r   get_asmr   r{   rj   rn   ro   r>  rJ  )lowered_modulerE  r~   r   r   r   r|   r   r  r   r   r   r   r   r   r   r   r   r;  r   rz   r	   s                         r#   lowered_as_tpu_kernelrQ    s    , !0+%//77T 8   $'   	66 
 *'!+ 9!/+)'-11[&" 
/
 r%   c                    dt        |t        j                  j                        s|f}dt	        d |D               fd}t        j                  |      S )NFTc              3  p   K   | ].  }t        j                  |j                  |j                         0 y wr*   )r
   ShapedArrayr   dtype)r,   tys     r#   r/   z#_as_jax_callable.<locals>.<genexpr>M  s%     LRD$$RXXrxx8Ls   46c            
     J    t        j                  | d}r|d   S |S )NrD  r   )r8   r9   )	r<   resultr	   r   r   r   r   r   unpacks	     r#   apply_kernelz&_as_jax_callable.<locals>.apply_kernelP  s;    ##	)1F 6!9*F*r%   )r   collectionsabcIterabler6   r   jit)	r	   r   rE  r   r   r   rZ  r   rY  s	   `` ``` @@r#   rJ  rJ  @  sU     &	Hkoo66	7{HFL8LL)
+ 
+ 
	r%   )r!   mlir.LoweringRuleContextrS   r}   )r>   r   rS   rg   )rS   zSequence[Sequence[int]])r!   r_  r	   ry   r   rj   r   r   r   r   r   tuple[tuple[int, int], ...]r   
Any | NonerS   zir.OpResultList)r  	ir.Moduler   r}   rS   z#tuple[ir.Module, tuple[bool, bool]])r  rb  rS   r   )r  rb  rS   r}   )$r  rb  r   r}   r   r   r   r   r   r   r   r}   r~   r}   r   r}   r   r   r   r}   r   r{   r   r   r   r{   r;  r{   r   r{   r   bool | Noner   r   rS   ry   )*rz   rg   r   r}   r   r   r   r   r   r   r   r}   r~   r}   r   r}   r  r{   r|   r{   r   r{   r   r{   r   r   r   r   r   r{   r   r}   r   r   r   r{   r;  r{   r   r{   r   r   )0r!   r_  r   zir.Valuer  rb  rE  r   r   rO   r   r   r   r}   r   r   r   r   r   r`  r   r}   r~   r}   r   bool | TpuSideEffectTyper   r}   r   r   r   r{   r   r   r   ra  r   r{   r;  r{   r   r{   r   rc  r   r   rS   zSequence[ir.Value])(r  rb  rE  r   r   r   r   r   r   r}   r   r   r   r   r   r`  r   r}   r~   r}   r   rj   r   r}   r   r   r   r{   r   r   r   r{   r   rc  r   ra  rG  r}   rS   Callable[..., Any])(rP  rb  rE  r   r~   r}   r   r   r   r{   r   r{   r|   r{   r   rd  r  r{   r   r   r   r}   r   r   r   r   r   r`  r   r}   r   r}   r   r{   r   ra  r;  r{   rS   re  )r	   ry   r   rj   rE  r   r   r   r   r`  r   ra  rS   re  )Fr   
__future__r   r   collections.abcr[  r   r   dataclassesrX   r   r   typingr   r   jax._srcr   r	   r
   r   r   jax._src.cloud_tpu_initr   jax._src.frozen_dictr   jax._src.interpretersr   r   jax._src.libr   jaxlib.mlirr   jaxlib.mlir.passmanagerr   abslr   FLAGSImportError
bool_stater=  r$   	Primitiver8   multiple_resultssimple_implrB   fancy_primitive_batchersEnumrD   r\   rj   rr   	dataclassry   r   def_abstract_evalr   r   r   register_loweringr  r$  r:  r?  r>  rF  rn   rK  rQ  rJ  r2   r%   r#   <module>r}     s4    #   .   	  !     # ; + * &   /
++% &F%%		' $	 #DNN#45 %)  "   & ' 8O ! !"3 4=$)) =49 $		 $%TYY %
 d#[ [ $[| $$ %;@	!@ $@ (	@
 @ @ 6@ @ @F   (*C %' "  )	@#LH j CG!"'AE %7<%*'+ %66 !6 '	6
 06 .6  *6 6 %6 @6 6  6 ?6 6 156  #!6" %#6$ %6& '6P CG"'$(AE %7<%* -;; !; '	;
 0; .;  *; ; %; ; ; ; ; ; @;   !;" "#;$ ?%;& ';( 15);* #+;, -;\ #( %7<%*'+ /;	!;; ; 	;
 ; '; !; 0; .; 6;  *; ; /; %; @;   !;" ?#;$ %;& ';( 15);* #+;, %-;. /;0 1;D *."#'26048:,0 $*;*@*@'(BF"'AE%*'+")... '	.
 . !. 0. .. 6.  *. . (. %. @.  .  ?!." ##.$ %%.& '.( ).* +.j !%)-" %#16$"#'26048:'+,0"'7<)888 	8
 '8 8 8 8 /8 8 8 !8 08 .8 68  %!8"  *#8$  %8& '8( 15)8* +8v#' 
  6  g  
%s   :K! !K,+K,