
    VpfsT             "      )   d Z ddlmZ ddlmZmZ ddlZddlZddlZddl	Z	ddl
mZmZmZ ddlZddlmZ ddlmZ ddlmZ dd	lmZ dd
lmZ ddlmZ ddlmZ ddlmZ ddlmZ ddlmZ ddlmZ ddlmZ ddlmZ ddlm Z  ddlm!Z" ddl#m$Z$ ddl%m&Z& ddl'm(Z) ddl'mZ* ddl'm+Z, ddl-m.Z/ ddl0mZ1 ddl0m2Z2 ddl0m3Z4 ddl5m6Z6 ddl5m7Z7 ddl5m2Z8 ddl9m:Z: dd l9m;Z; dd!l9m<Z< ddl=m>Z? ddl>Z@ ed"          ZAejB        eCcZCZDejE        eFcZFZGe7jH        ZHe1jI        ZIe1jJ        ZJe1jK        ZKejL         G d# d$                      ZMejL         G d% d&                      ZNejL         G d' d(                      ZOejL         G d) d*                      ZP G d+ d,eQ          ZRdd0ZSdd6ZTdd=ZUi ZVddAZWddDZXddFZYddLZZddPZ[ddSZ\ eWe2j]                  ddT            Z^ eWe2j_                  ddU            Z`de/ja        jb        e/jc        jd        fdd`Ze eWe2jf                  ddc            Zg eWe2jh                  ddd            ZiddeZj eWejk                  ddh            Zl eWejm                  ddi            Zn ejL        djk           G dl dm                      Zo ejL        djk           G dn do                      ZpddsZq eqdt eodugdvdu           eodwgdxdw           eodygdzdy           eod{gd|d{          g epdugd}            epdwgd~            epdygd            epd{gd           g          Zr eWejs                  dd            ZteVu                    i ejv        d ejw         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ejx         eqd eodygddy           eod{gdd{           epdgd            epdgd           g eodygddy           eod{gdd{           epdgd            epdgd           g          ejy         eqd eodygddy           eod{gdd{           epdgd            epdgd           g epdygd            epd{gd            epdgd            epdgd           g          ejz         eqd eodygddy           eod{gdd{           epdgd            epdgd           g eodygddy           eod{gdd{           epdgd            epdgd           g          ej{         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej|         eqd eodygddy           eod{gdd{           epdgd            epdgd           g eodygddy           eod{gdd{           epdgd            epdgd           g          ej}         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej~         eqd eodygddy           eod{gdd{           epdgd            epdgd           g eodygddy           eod{gdd{           epdgd            epdgd           g          ej         eqd eodydugddy           eod{dugdd{           eodydygddy           eod{d{gdd{          g eodydugddy           eod{dugdd{           eodydygddy           eod{d{gdd{          g          ej         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej         eqd eodygddy           eod{gdd{           epdgd؄            epdgdل           g eodygddy           eod{gdd{           epdgd܄            epdgd݄           g          ej         eqd eodygddy           eod{gdd{           epdgd            epdgd           g eodygddy           eod{gdd{           epdgd            epdgd           g          ej         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej         eqd eodydygddy           eod{d{gdd{          g eodydygddy           eod{d{gdd{          g          ej         eqd  eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej         eqd eodygddy           eod{gdd{          g eodygddy           eod{gd	d{          g          ej         eqd
 eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej         eqd eodygddy           eod{gdd{          g eodygddy           eod{gdd{          g          ej         eqd eodugddu           eodwgd du          g epdugd!            epdwgd"           g          ej         eqd# eodugd$du           eodwgd%du          g epdugd&            epdwgd'           g          ej         eqd( eodydygd)dy           eod{d{gd*d{          g eodydygd+dy           eod{d{gd,d{          g          i           d̐d-Zd͐d.Zdΐd/Zdΐd0Zdϐd2Zdϐd3Zdϐd4ZdАd:Z ej        ee)j        j        e)j        j        e)j        j        ;          Z ej        ee)j        j        e)j        j        e)j        j        ;          Z ej        ee)j        j        e)j        j        e)j        j        ;          Z ej        ee)j        j        e)j        j        e)j        j        ;          Z ej        ee)j        j        e)j        j        e)j        j        ;          Z ej        ee)j        j        e)j        j        e)j        j        ;          Zej        eej        eej        eej        e)j        ej        e)j        ej        e)j        ej        e)j        ej        e)j        ej        e)j        ej        ei
Zeà                                D ]\  ZZefdǐd<ZeeVe<   ej        eej        eej        eej        eej        eej        eej        eiZeϠ                                D ]\  ZZefdǐd=ZeeVe<    eWe2j                  dѐdA            ZҐdҐdFZ eWe2j                  dӐdI            Z eWe2j                  dӐdJ            Z eWe8j                  dԐdK            Z eWej                  dՐdL            Zېd֐dPZej        dQ ej        dR iZeߠ                                D ]\  ZZ eeƐdST          eVe<    eWej                  dǐdU            Z eWej                  dǐdV            Z eWej                  dǐdW            Z eWej                  dǐdX            Z eWej                  dǐdY            Zdאd\Zdؐd_ZdِdaZdڐdcZdېddZdܐdgZdݐdhZdݐdiZdݐdjZdސdmZdݐdnZ eWej                  dǐdo            Z eWej                  dǐdp            Z eWej                  dǐdq            Z eWej                  dǐdr            ZdߐdsZ eWej                  dǐdt            Zdd{Z  eWe8j                  dǐd|            Zd} e/j        D             Zd~ e/j        D             Z	 	 ddddSdddZ eWe2j                  dǐd            Z	 eWe8j
                  dǐd            Z	 ddddddZ eWe2j
                  dǐd            Z eWe8j                  dǐd            Z eWej                  dǐd            ZddZ	 ddjdddddZej        j        ej        j        fZ eWej                  dǐd            ZdǐdZdǐdZ ej        ee?j                  eVej        <    ej        ee?j                  eVej        <    ej        ee?j                   eVej!        <   dǐdZ"d Z# ej        e"e#          eVej$        <   d Z% ej        e"e%          eVej&        <    eWej'                  dǐd            Z( eWej)                   eWej*                  dǐd                        Z+ eWej,                  dǐd            Z-d eVej.        <    eWej/                  dd            Z0ddZ1 eWe$j2                  dǐd            Z3dddddZ4 eWej5                  dǐd            Z6dǐdZ7 eWej8                  dǐd            Z9 eWej:                  dǐd            Z;ddZ<ddZ=ddZ>ddZ?ddZ@dS (  z0Module for lowering JAX primitives to Triton IR.    )annotations)CallableSequenceN)AnyHashableTypeVar)lax)	tree_util)ad_checkpoint)ad_util)api_util)config)core)custom_derivatives)linear_util)pjit)source_info_util)state)util)mlir)partial_eval)for_loop)ir)arith)math)scf)dialect)
primitives)utils)	discharge)indexing)merge_lists)partition_list)
split_list_Tc                  d    e Zd ZU ded<   ded<   ded<    ej        d          Zd	ed
<   ded<   dS )ModuleContextstrnameGridMappinggrid_mappingSequence[ir.Value]program_idsF)reprzmlir.TracebackCachestraceback_cachesplatformN)__name__
__module____qualname____annotations__dataclassesfieldr/        _/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/_src/pallas/triton/lowering.pyr'   r'   L   s\         )))!!!!+<;+<%+H+H+HHHHH-----r8   r'   c                  .    e Zd ZU ded<   ded<   ded<   dS )	BlockInfozjax.ShapeDtypeStructfull_shape_dtypeSequence[Any]start_indicestuple[int, ...]block_shapeN)r1   r2   r3   r4   r7   r8   r9   r;   r;   U   s6         ((((r8   r;   c                  F    e Zd ZU ded<   ded<   ded<   ded<   ej        ZdS )	LoweringRuleContextr'   contextSequence[jax_core.ShapedArray]avals_in	avals_outzSequence[BlockInfo | None]block_infosN)r1   r2   r3   r4   r5   replacer7   r8   r9   rB   rB   \   sH         ****++++))))'''r8   rB   c                  (    e Zd ZU dZded<   ded<   dS )LoweringResultzKeeps pybind11 objects alive.z	ir.Modulemoduler?   gridN)r1   r2   r3   __doc__r4   r7   r8   r9   rJ   rJ   f   s0         %%r8   rJ   c                      e Zd ZdS )LoweringErrorN)r1   r2   r3   r7   r8   r9   rO   rO   n   s        $r8   rO   ctxblock_mappingBlockMappingc                    t          | |j        j        d g|R  }d |D             }t          d t	          ||j                  D                       S )Nc              3  p   K   | ]1}t          |t          j        d t          j                            V  2dS )r7   N)_ensure_ir_valuejax_coreShapedArrayjnpint32).0is     r9   	<genexpr>z"_eval_index_map.<locals>.<genexpr>x   sO        
 q(.r39==>>     r8   c              3     K   | ]:\  }}|t           j        u r|n"t          |t          ||j                            V  ;d S N)pallas_coremapped_mul_ir_constanttype)rZ   r[   bs      r9   r\   z"_eval_index_map.<locals>.<genexpr>|   sa        
!Q """aaQQ0G0G(H(H     r8   )lower_jaxpr_to_triton_irindex_map_jaxprjaxprtuplezipr@   )rP   idxrQ   block_indicess       r9   _eval_index_maprl   r   s     +	=(.7:  -   - 
  m]%>??   
 
 r8   air.Valueshaper?   returnc                H   t           j                            | j                  s<|s| S t	          j        t           j                            || j                  |           S t          j        | j                  j        g |k    r| S j        t          |          k    s(t          fdt          |          D                       st          dj         dg |           t	          j        t           j                            |j        j                  |           S )Nc              3  @   K   | ]\  }}j         |         |d fv V  dS    Nro   )rZ   r[   dima_types      r9   r\   z_bcast_to.<locals>.<genexpr>   sG       , ,(.3QC8#, , , , , ,r8   zCannot broadcast from  to )r   RankedTensorType
isinstancerc   
tt_dialectsplatgetro   ranklenall	enumerate
ValueError	broadcastelement_typeencoding)rm   ro   rw   s     @r9   	_bcast_tor      s&   			'	'	/	/  hB/33E16BBAFFF ((F|xxh{c%jj   , , , ,2;E2B2B, , , ) )  LLL(E(LLMMM
v':FOLLa  r8   xyx_avaljax_core.ShapedArrayy_avalout_avalc                .   t          | t          j        t          j        t          t
          f          r2|j        }|j        r|j        }t          | t          |                    } t          |t          j        t          j        t          t
          f          r2|j        }|j        r|j        }t          |t          |                    }|j
        |j
        k    rt          | |j
                  } |j
        |j
        k    rt          ||j
                  }| |fS r^   )rz   npndarraynumberintfloatdtype	weak_typerb   _dtype_to_ir_typero   r   )r   r   r   r   r   x_dtypey_dtypes          r9   _bcastr      s     BJ	3677 4lG gQ)'2233ABJ	3677 4lG gQ)'2233A\X^##!X^$$A\X^##!X^$$A	
A+r8   	primitivejax_core.PrimitiveCallable[[_T], _T]c                      fd}|S )Nc                    | t           <   | S r^   )triton_lowering_rules)fnr   s    r9   wrapperz"register_lowering.<locals>.wrapper   s    '))$Ir8   r7   )r   r   s   ` r9   register_loweringr      s#         
.r8   r+   r*   c                    g }g }t          | j                  D ]8\  }}|| j        vr*|                    |           |                    |           9t	          | j                  D ]9}| j        |         }|                    |           |                    |           :t          |d d                   }d}|dz   t          |          k     r||dz            |k    r|dz  }n$|t          |          k     r||         |k    r|dz  }|d |         }||d          }	t          |          dk    rPd gt          |	          z  }
t          t          |	                    D ]}||         }t          |          |
|<   |	|
fS t          j	        |          g|	}|d         dk     s
J d            d gt          | j                  z  }t          d          }t          |          D ]C\  }}||         }t          |          }t          ||d	          ||<   t          ||d	          }Dt          t          |	                    D ]"}|||z            }t          |dz             ||<   #t          |          t          | j                  k    sJ ||fS )
Ni  rt      r   iz7Cannot fix pallas kernel launch grid within CUDA limitsFsigned)r   rL   vmapped_dimsappendreversedr   range_program_idr   prod_i32_constant_mod	_floordiv)r+   launch_gridlaunch_grid_to_pallas_gridr[   srv   num_collapsecuda_yz_limitcollapse_dimsprog_id_dimsprog_idsout_idxnew_gridout_indicesgrid0s                  r9   _process_grid_to_3d_gridr      s   +! )** + +da))) ''*** l/00 + +c#Aq%%c****["%&&,- Q[)))),"#m33ALLs;''''L!M11ALm|m,-\]]+,1vL)))H3|$$%% ) )*1-g%a..hw!!	-((8<8H	!y	 	 	 
C 
!	 	  \.///+
a..%&& . .da(+GaAq777KeQu---EE\""## . .a()9:G&q1u--K	[		S!233	3	3	3	3	;	r8   
ir.Contextc                 |    t          j                    } t          j        |            |                                  | S r^   )r   Contextr{   register_dialectload_all_available_dialects)rP   s    r9   _new_ir_contextr      s5    

#c"""!!###	*r8   rg   jax_core.Jaxprr)   r(   r0   c                   j         rt          d          j        rt          d                                          5  t	          j        | dgt          | j                  z  d          \  } }d d d            n# 1 swxY w Y   t                      5  t          j
                                        5  t          j                                        }d | j        D             }t          | j                  dk    sJ t          j                            |g           }t#          j        |t          j                            |          dt          j                            t+          t          j                            d          	                    t          j                            |j                  
          }t          j                            t          j                            dt7          j        d          i          gt          |          z            |_         |j        j        j        |j           |j        j        \  }	t          j        |	          5  tC                    \  }
}fdtE          |          D             }tG          ||t7          j$                    |          }j        rt          d          tK          d j&        D                       st          d          tO          tQ          j)        tT          ||          j&                  }d tW          j&        |          D             }tY          || |g|	j-        R  \   t#          j.        g            d d d            n# 1 swxY w Y   t_          ||
          cd d d            cd d d            S # 1 swxY w Y   d d d            d S # 1 swxY w Y   d S )Nz7dynamic grid bounds not supported in the Triton backendz5scalar prefetch not implemented in the Triton backendT)instantiatec                |    g | ]9}t           j                            t          |j        j                  d           :S rt   )r{   PointerTyper}   r   avalr   )rZ   vars     r9   
<listcomp>z0lower_jaxpr_to_triton_module.<locals>.<listcomp>  sH        	""#4SX^#D#DaHH  r8   r   publicF)noinline)sym_visibility	res_attrsiptt.divisibility    c                ,    g | ]\  }}|j         v|S r7   )r   )rZ   r[   pidr+   s      r9   r   z0lower_jaxpr_to_triton_module.<locals>.<listcomp>/  s5       al/// ///r8   z1Scalar prefetch not supported in Triton lowering.c              3  J   K   | ]}t          |j        t                    V  d S r^   )rz   indexing_modeBlocked)rZ   bms     r9   r\   z/lower_jaxpr_to_triton_module.<locals>.<genexpr>;  sC       8 8 B,g66 8 8 8 8 8 8r8   z;Only Blocked indexing mode is supported in Triton lowering.c                J    g | ] \  }}t          |j        ||j                  !S r7   )r;   array_shape_dtyper@   )rZ   rQ   	start_idxs      r9   r   z0lower_jaxpr_to_triton_module.<locals>.<listcomp>D  sG     
 
 
 'mY -' 
 
 
r8   )0num_dynamic_grid_boundsNotImplementedErrornum_index_operands	trace_envpe	dce_jaxprr   outvarsr   r   LocationunknownModulecreateinvarsFunctionTyper}   r{   FuncOpTypeAttrDictAttrdictBoolAttrInsertionPointat_block_beginbody	ArrayAttrr   i32_attr	arg_attrsblocksr   inputsr   r   r'   TracebackCachesr   block_mappingsmap	functoolspartialrl   ri   re   	argumentsreturn_rJ   )rg   r+   r)   r0   _rK   param_typesfn_typer   entryr   r-   local_program_idsrP   r>   rG   s    `              r9   lower_jaxpr_to_triton_moduler     s     ) 
A   $ 
?     |vEM***  HE1                9, 9,"+--// 9, 9,YF <  K u}""""o!!+r22G		
  +//$0F0F"G"G"GHH++FK88
 
 
B <##	+T]2->->?	@	@A
k

	 BL BGN7>**gnGU		5	!	! # #6|DDh   !+..  
 
/1E1G1G c 
	( 
!?
 
 	
  8 8'68 8 8 8 8 
!I
 
 	
 

OS+
>
>

% m
 
 +.)+ +
 
 
k $CNeoNNNbG# # # # # # # # # # # # # # #H &(++s9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9, 9,sn   1A??BBO7F+N,"DM9-N,9M==N, M=N,O,N0	0O3N0	4OOOrG   !Sequence[BlockInfo | None] | Noner=   c                4   i i dfd}dfd}dfd}| t          |j        |          D ]
\  }}||<   t          ||j        |           |j        D ]}	t          ||	j                  }
|	j        t
          vrt          d|	j        j         d	          t
          |	j                 }d
 |	j        D             }d |	j        D             }t          ||	j                  }t          j
        | |	j        |	j        |	j                  }t          | |||          }	 t          j        |	j        j                  5  |5   ||g|
R i |	j        }d d d            n# 1 swxY w Y   d d d            n# 1 swxY w Y   nI# t"          $ r  t$          $ r2}t          d |
          }t#          d|	 d| d| d|           |d }~ww xY w|	j        j        rt          ||	j        |            ||	j        d         |           t          ||j                  S )Natomjax_core.Atomc                V    t          | t          j                  r| j        n|          S r^   )rz   rV   Literalval)r  envs    r9   read_envz*lower_jaxpr_to_triton_ir.<locals>.read_env]  s%    !$(899H488s4yHr8   c                h    t          | t          j                  rd S                     | d           S r^   )rz   rV   r  r}   )r  block_info_envs    r9   read_block_info_envz5lower_jaxpr_to_triton_ir.<locals>.read_block_info_env`  s3    $()) TdD)))r8   r   jax_core.Varc                    || <   d S r^   r7   )r   r	  r
  s     r9   	write_envz+lower_jaxpr_to_triton_ir.<locals>.write_enve  s    CHHHr8   z0Unimplemented primitive in Pallas GPU lowering: z?. Please file an issue on https://github.com/google/jax/issues.c                    g | ]	}|j         
S r7   r   rZ   vs     r9   r   z,lower_jaxpr_to_triton_ir.<locals>.<listcomp>v  s    +++1+++r8   c                    g | ]	}|j         
S r7   r  r  s     r9   r   z,lower_jaxpr_to_triton_ir.<locals>.<listcomp>w  s    ---A---r8   c                $    t          | dd           S )Nrc   )getattrts    r9   <lambda>z*lower_jaxpr_to_triton_ir.<locals>.<lambda>  s    '!VT":": r8   z Exception while lowering eqn:
  z
With context:
  z
With inval types=z
In jaxpr:
r   )r  r  )r   r  )ri   r   r   eqnsr   r   r   r)   r   r   _source_info_to_locationparamssource_inforB   r   user_context	tracebackrO   	Exceptionmultiple_results)rP   rg   rG   argsr  r  r  invar
block_infoeqninvalsrulerE   rF   eqn_block_infoslocrule_ctxoutvalseinval_typesr  r
  s                       @@r9   re   re   T  s|    	#.I I I I I I* * * * * *
       {;; ) )z(nUit$$$Z ) )c3:&&F
}111J]J J JK K K !/D++
+++H-----I-sz::O

'S]CJ C #3)_MMH
()BCC 8 8S 8 8$x7&777CJ778 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8      ::FCCkMc M MM M+6M MEJM M   }% )	)S['****iA((((	Xu}	%	%%s`   E?;E3>EE3E  E3#E $E3'E?3E7	7E?:E7	;E??G-G  Gaxisr   c                p    | t          d          vrt          d|            t          j        |           S N   z!axis must be in [0, 3), but got: )r   r   r{   get_program_idr0  s    r9   r   r     s;    	q
???
@
@@		"4	(	((r8   c               &    | j         j        |         S r^   )rC   r-   rP   r0  s     r9   _program_id_lowering_ruler8    s    		 	&&r8   c               p    |t          d          vrt          d|           t          j        |          S r2  )r   r   r{   get_num_programsr7  s     r9   _num_programs_lowering_ruler;    s;    	q
???
@
@@		$T	*	**r8   optt_dialect.RMWOpptrr	  maskir.Value | Nonesemantictt_dialect.MemSemantic
sync_scopett_dialect.MemSyncScopec           	        t           j                            |j                  rct          j        |j                  }t	          j        |j                  }t           j                            |j        |j	        |j
                  }nt	          j        |j                  j	        }t	          j        || |||||          S )N)r?  semscope)r   ry   rz   rc   r{   r   r   r}   ro   pointee_typer   
atomic_rmw)	r<  r>  r	  r?  rA  rC  ptr_typer   result_types	            r9   _atomic_rmwrL    s     ##CH-- @"38,,H)(*?@@L%))183D KK (22?K		2sCd

 
 
 r8   atomic_typeprimitives.AtomicOpTypec               4   |                     |          \  }}}}|                     | j                  ^ }}	}
t          |          dk    rt          d          |d         }t	          || j        d         || j        d         j                  }t          ||	          }|t          ||
          }|t          j	        j
        k    rt          j        j
        }n-|t          j	        j        k    rCt          |j        t           j                  rt          j        j        }nt          j        j        }n|t          j	        j        k    rt          j        j        }n|t          j	        j        k    rt          j        j        }n|t          j	        j        k    rt          j        j        }n`|t          j	        j        k    rt          j        j        }n9|t          j	        j        k    rt          j        j        }nt          d|           t1          ||||          S )Nrt   z!Only single indexer is supported.r   zunsupported atomic operation: )r?  )	unflattenrE   r   r   _compute_pointers_from_indicesrG   ro   rU   r   AtomicOpTypeXCHGr{   RMWOpADDrz   rc   r   IntegerTypeFADDMINMAXANDORXORrL  )rP   	args_treerM  	args_flatr>  indexersr	  r?  r   
value_aval	mask_avalrj   r<  s                r9   _atomic_lowering_rulerb    s    '00;;#xd'11#,??1j)]]a
A
B
BB#&	3?1sCLO$9	 	# 	j))#	D),,DJ+000			BBj-111#(BN++ !bb bbj-111			BBj-111			BBj-111			BBj-000			BBj-111			BB
L{LL
M
MM	Rc	-	-	--r8   c           	     
   | j         \  }}}t          j                            |j                  rct          j        |j                  }t          j        |j                  }t          j                            |j	        |j
        |j                  }	nt          j        |j                  j
        }	t          j        |	|t          ||          t          ||          t
          j        j        t
          j        j                  S )N)rF  rG  )rE   r   ry   rz   rc   r{   r   r   r}   ro   rH  r   
atomic_casrU   MemSemanticACQUIRE_RELEASEMemSyncScopeGPU)
rP   r>  cmpr	  r   cmp_avalval_avalrJ  r   rK  s
             r9   _atomic_cas_lowering_rulerl    s    ,!Xx##CH-- @"38,,H)(*?@@L%))183D KK (22?K			sH%%sH%%

 
0#'
 
 
 r8   c                4   t          j        |          }|\  }|j        d         j        }t	          j        d|          t	          j        d|          g}t          j        ||f          }t          j        t          j
        |           |          \  }	}
t          j        |	|          \  }}}\    |
            }~|rt          d          d |D             }t          j        ||          }|dz  } |j        d         j        j        | }t&          j                            |          5  t-          |j        |d g|j        R  }t          j        |           d d d            n# 1 swxY w Y   |                                 t7          |j                  S )Nr   r7   r   z.Associative scan with constants not supported.c                6    g | ]}t          |j                  S r7   _element_typerc   rZ   args     r9   r   z._associative_scan_lowering.<locals>.<listcomp>
  "    @@@s=**@@@r8   r   )r
   tree_leavesrE   r   rV   rW   tree_structurer   flatten_fun_nokwargslu	wrap_initr   trace_to_jaxpr_dynamicr   r{   ScanOpregionsr   r   r   r   r   re   rC   r   scan_returnverifylistresult)r   rP   r$  axes	flat_argsr0  r   in_avalsin_treeflat_funout_tree_thunkcombine_jaxprr   constsout_treeelement_typesscan_opr   r   resultss                       r9   _associative_scan_loweringr    s   #D)))'4
,q/
%2U+++2U+++( $dD\22'%:l4' (N "$!:" "-FB ^( P
N
O
OO@@i@@@-i..'!+
*'/!

#
*K
8%	''.. $ $&]D+0?  G 7###	$ $ $ $ $ $ $ $ $ $ $ $ $ $ $
 
..	gn		s   -.E''E+.E+reverseboolc               j    |rt          d          t          t          j        | ||f          d         S )Nz Reverse cumsum is not supported.r   )r   r  rX   add)rP   r   r0  r  s       r9   _cumsum_lowering_ruler    s9      B
@
A
AA	#CGS!dW	=	=a	@@r8   c           	         | j         \  }t          j        |t          |j        |j                            d                               S Nr   )rE   arith_dialectxori_fullrc   r   )rP   r   r   s      r9   _not_lowering_ruler     s=    \(6		AuQVfl.?.?.B.B-BCC	D	DDr8   T)frozenc                  >    e Zd ZU ded<   ded<   ded<   dd
ZddZdS )_ExternSequence[str]	arg_typesr(   symbolrK  avalsrD   rp   r  c                    t          |          t          | j                  k    rdS t          d t          || j                  D                       S )NFc              3  J   K   | ]\  }}|j         p|j        j        |k    V  d S r^   )r   r   r)   )rZ   r   arg_types      r9   r\   z"_Extern.matches.<locals>.<genexpr>/  sJ        D( 	5$*/X5     r8   )r   r  r   ri   )selfr  s     r9   matchesz_Extern.matches,  s\    
5zzS((((U  !%88     r8   rP   rB   r$  r,   c                    |j         \  }t          t          j        | j                            }|j        r%t          j                            |j        |          }t          j
        ||dd| j        d          S )N T)libnamelibpathr  pure)rF   r   rX   r   rK  ro   r   ry   r}   r{   extern_elementwiser  )r  rP   r$  r   rK  s        r9   lowerz_Extern.lower4  sx    JX#CId.>$?$?@@K~ I'++HNKHHk({   r8   N)r  rD   rp   r  )rP   rB   r$  r,   )r1   r2   r3   r4   r  r  r7   r8   r9   r  r  &  s_         +++        r8   r  c                  2    e Zd ZU ded<   ded<   ej        ZdS )	_Fallbackr  r  Callable[..., ir.Value]r  N)r1   r2   r3   r4   r  r  r7   r8   r9   r  r  C  s/             O'''r8   r  tablesSequence[_Extern | _Fallback]r  c                     d fd}|S )NrP   rB   r$  rn   rp   c                0     j         j                 }t           fd|D             d           }|3t          d  j        D                       }t          d d|            j        \  }g }t           j        ||j                  D ]}\  }}}	t          t          ||          |j                  }
|j        r8|j        j        |	k    r(t          |
|j        t!          j        |	                    }
|                    |
           ~ |j         g|R  S )Nc              3  P   K   | ] }|                     j                  |V  !d S r^   )r  rE   )rZ   r.  rP   s     r9   r\   z6_make_dispatch_table.<locals>.inner.<locals>.<genexpr>Q  s6      ::A!))CL"9"9:a::::::r8   c              3  .   K   | ]}|j         j        V  d S r^   )r   r)   rZ   r   s     r9   r\   z6_make_dispatch_table.<locals>.inner.<locals>.<genexpr>S  s&      GG$djoGGGGGGr8   zunsupported types for z: )rC   r0   nextrh   rE   r   rF   ri   r  r   rU   ro   r   r   r)   _castrX   r   r  )rP   r$  tableharg_aval_dtypesr   
bcast_argsr   rs  r  	bcast_argr)   r  s   `          r9   innerz#_make_dispatch_table.<locals>.innerO  s8   3;'(E:::::::DAAAyGG#,GGGGGo
<4
<
<?
<
<   JXJ"3<q{CC # #c8,S$77HHi	 FDJOx77)TZ81D1DEE		""""173$$$$$r8   )rP   rB   r$  rn   rp   rn   r7   )r)   r  r  s   `` r9   _make_dispatch_tabler  K  s.    % % % % % % %$ 
,r8   absrY   __nv_absint64
__nv_llabsfloat32
__nv_fabsffloat64	__nv_fabsc                *    t          j        |          S r^   math_dialectabsirP   r   s     r9   r  r  m      L,=a,@,@ r8   c                *    t          j        |          S r^   r  r  s     r9   r  r  n  r  r8   c                *    t          j        |          S r^   r  absfr  s     r9   r  r  o      l.?.B.B r8   c                *    t          j        |          S r^   r  r  s     r9   r  r  p  r  r8   )cudarocmc                @   	 t          | |          S # t          $ r}| j        \  }t          j        |t          j                  rt          j        |          cY d }~S t          j        |t          j                  rt          j	        |          cY d }~S |d d }~ww xY wr^   )
_abs_dispatch_tabler   rE   rX   
issubdtypeintegerr  r  floatingr  )rP   r   r.  r   s       r9   _abs_lowering_ruler  u  s    	sA&&&	   |HV
~fck** q!!!!!!!!		-	- q!!!!!!!!4s,    
B;BB2BBBBc                     t          |          S r^   )_minusr  s     r9   r  r    s    fQii r8   ceil
__nv_ceilf	__nv_ceil__ocml_ceil_f32__ocml_ceil_f64floor__nv_floorf
__nv_floorfloat16c                *    t          j        |          S r^   r  r  r  s     r9   r  r        ,2DQ2G2G r8   bfloat16c                *    t          j        |          S r^   r  r  s     r9   r  r        <3Ea3H3H r8   __ocml_floor_f32__ocml_floor_f64c                *    t          j        |          S r^   r  r  s     r9   r  r    r  r8   c                *    t          j        |          S r^   r  r  s     r9   r  r    r  r8   exp	__nv_expf__nv_expc                *    t          j        |          S r^   r  r  r  s     r9   r  r        ,2B12E2E r8   c                *    t          j        |          S r^   r  r  s     r9   r  r        <3CA3F3F r8   c                *    t          j        |          S r^   r  r  s     r9   r  r    r  r8   c                *    t          j        |          S r^   r  r  s     r9   r  r    r  r8   c                *    t          j        |          S r^   r  r  s     r9   r  r    r  r8   c                *    t          j        |          S r^   r  r  s     r9   r  r    r  r8   exp2
__nv_exp2f	__nv_exp2c                *    t          j        |          S r^   r  r  r  s     r9   r  r        ,2CA2F2F r8   c                *    t          j        |          S r^   r   r  s     r9   r  r        <3DQ3G3G r8   __ocml_exp2_f32__ocml_exp2_f64c                *    t          j        |          S r^   r   r  s     r9   r  r    r  r8   c                *    t          j        |          S r^   r   r  s     r9   r  r    r  r8   expm1__nv_expm1f
__nv_expm1__ocml_expm1_f32__ocml_expm1_f64log	__nv_logf__nv_logc                *    t          j        |          S r^   r  r  r  s     r9   r  r    r  r8   c                *    t          j        |          S r^   r  r  s     r9   r  r    r  r8   __ocml_log_f32__ocml_log_f64c                *    t          j        |          S r^   r  r  s     r9   r  r    r  r8   c                *    t          j        |          S r^   r  r  s     r9   r  r    r  r8   log1p__nv_log1pf
__nv_log1p__ocml_log1p_f32__ocml_log1p_f64sqrt
__nv_sqrtf	__nv_sqrtc                *    t          j        |          S r^   r  r  r  s     r9   r  r    r  r8   c                *    t          j        |          S r^   r   r  s     r9   r  r    r  r8   __ocml_sqrt_f32__ocml_sqrt_f64c                *    t          j        |          S r^   r   r  s     r9   r  r    r  r8   c                *    t          j        |          S r^   r   r  s     r9   r  r    r  r8   pow
__nv_powif	__nv_powi	__nv_powf__nv_pow__ocml_pown_f32__ocml_pown_f64__ocml_pow_f32__ocml_pow_f64cbrt
__nv_cbrtf	__nv_cbrt__ocml_cbrt_f32__ocml_cbrt_f64rsqrt__nv_rsqrtf
__nv_rsqrt__ocml_rsqrt_f32__ocml_rsqrt_f64sin	__nv_sinf__nv_sinc                *    t          j        |          S r^   r  r9  r  s     r9   r  r    r  r8   c                *    t          j        |          S r^   r=  r  s     r9   r  r    r  r8   __ocml_sin_f32__ocml_sin_f64c                *    t          j        |          S r^   r=  r  s     r9   r  r  !  r  r8   c                *    t          j        |          S r^   r=  r  s     r9   r  r  "  r  r8   cos	__nv_cosf__nv_cosc                *    t          j        |          S r^   r  rC  r  s     r9   r  r  *  r  r8   c                *    t          j        |          S r^   rG  r  s     r9   r  r  +  r  r8   __ocml_cos_f32__ocml_cos_f64c                *    t          j        |          S r^   rG  r  s     r9   r  r  0  r  r8   c                *    t          j        |          S r^   rG  r  s     r9   r  r  1  r  r8   tan	__nv_tanf__nv_tan__ocml_tan_f32__ocml_tan_f64asin
__nv_asinf	__nv_asin__ocml_asin_f32__ocml_asin_f64acos
__nv_acosf	__nv_acos__ocml_acos_f32__ocml_acos_f64atan
__nv_atanf	__nv_atan__ocml_atan_f32__ocml_atan_f64atan2__nv_atan2f
__nv_atan2__ocml_atan2_f32__ocml_atan2_f64sinh
__nv_sinhf	__nv_sinh__ocml_sinh_f32__ocml_sinh_f64cosh
__nv_coshf	__nv_cosh__ocml_cosh_f32__ocml_cosh_f64tanh
__nv_tanhf	__nv_tanh__ocml_tanh_f32__ocml_tanh_f64asinh__nv_asinhf
__nv_asinh__ocml_asinh_f32__ocml_asinh_f64acosh__nv_acoshf
__nv_acosh__ocml_acosh_f32__ocml_acosh_f64atanh__nv_atanhf
__nv_atanh__ocml_atanh_f32__ocml_atanh_f64population_count	__nv_popc__nv_popcllc                *    t          j        |          S r^   r  ctpopr  s     r9   r  r        0B10E0E r8   c                *    t          j        |          S r^   r  r  s     r9   r  r    r  r8   clz__nv_clz
__nv_clzllc                *    t          j        |          S r^   r  ctlzr  s     r9   r  r        0A!0D0D r8   c                *    t          j        |          S r^   r  r  s     r9   r  r    r  r8   	nextafter__nv_nextafterf__nv_nextafter__ocml_nextafter_f32__ocml_nextafter_f64c                    t           j                            t          | j                            rt          d| j                   t          t          | j        d          |           S )Nzunsupported type: r   )r{   r   rz   rq  rc   r   _subr  )r   s    r9   r  r    sZ    &&}QV'<'<== =
;16;;
<
<<	eAFA	"	""r8   c                   t          | j                  }t          |j                  }t          j                            |          r)t          j                            |          rJ || }} ||}}t          j                            |          rt          j        | j        | |          S | j        |j        k    s0J t          | j                  t          |j                  f            t	          |t          j                  rt          j
        | |          S t	          |t          j                  rt          j        | |          S t          d| j         d|j                   Nzunsupported dtypes:  and )rq  rc   r{   r   rz   addptrr(   r   rV  r  addi	FloatTypeaddfr   r   r   x_element_typey_element_types       r9   _addr    s>    ((. ((.&&~66 D%00@@@@@aqA%3^NN&&~66 +QVQ***	
16			CKKQV5			// La###.",// La###
JQVJJ!&JJ
K
KKr8   c                   t          | j                  }t          |j                  }t          j                            |          r(t          j        | j        | t          |                    S t          j                            |          s| j        |j        k    s0J t          | j                  t          |j                  f            t	          |t          j	                  rt          j        | |          S t	          |t          j                  rt          j        | |          S t          d|j                   )Nzunsupported dtype: )rq  rc   r{   r   rz   r  r  r(   r   rV  r  subir  subfr   r  s       r9   r  r    s    ((. ((.&&~66 &QVQq		222!,,^<< &6QVc!&kk3qv;;7.".11 &1%%%	NBL	1	1 &1%%%:!&::;;;r8   c                   | j         |j         k    s0J t          | j                   t          |j                   f            t          | j                   }t          |t          j                  rt          j        | |          S t          |t          j                  rt          j	        | |          S t          d| j          d|j                    Nunsupported types: r  )rc   r(   rq  rz   r   rV  r  mulir  mulfr   )r   r   r  s      r9   ra   ra     s    	
16			CKKQV5			 ((.// $a###.",// $a###G!&GGqvGGHHHr8   r   c                  | j         |j         k    s0J t          | j                   t          |j                   f            t          | j                   }t          |t          j        t          j        f          rt          j        | |          S t          |t          j	                  st          d| j          d|j                    |rt          j        | |          S t          j        | |          S r  )rc   r(   rq  rz   r   F32TypeF64Typer  divfrV  r   divsidivuir   r   r   r  s       r9   r   r     s    	
16			CKKQV5			 ((.RZ 899 $a###	NBN	3	3 K
IAFIIII
J
JJ %q!$$$q!$$$r8   c                  | j         |j         k    s0J t          | j                   t          |j                   f            t          | j                   }t          |t          j                  rBt          j                                        }t          | ||          } t          |||          }t          |t          j        t          j	        f          rt          j        | |          S t          d| j          d|j                    )Nr   r  r  )rc   r(   rq  rz   r   rV  r  r}   _int_float_castr  r  r  r   r  s       r9   _truedivr    s    	
16			CKKQV5			 ((.// :Z^^%%N>&999A>&999ARZ 899 $a###G!&GGqvGGHHHr8   c                  | j         |j         k    s0J t          | j                   t          |j                   f            t          | j                   }t          |t          j                  rt          j        | |          S t          |t          j                  st          d| j          d|j                    |rt          j
        | |          S t          j        | |          S r  )rc   r(   rq  rz   r   r  r  remfrV  r   remsiremuir  s       r9   r   r     s    	
16			CKKQV5			 ((.-- $a###	NBN	3	3 K
IAFIIII
J
JJ %q!$$$q!$$$r8   si_predarith_dialect.CmpIPredicateui_predf_predarith_dialect.CmpFPredicatec                  | j         |j         k    s0J t          | j                   t          |j                   f            t          | j                   }t          |t          j                  rt          j        |r|n|| |          S t          |t          j                  rt          j	        || |          S t          d| j          d|j                    r  )rc   r(   rq  rz   r   rV  r  cmpir  cmpfr   )r   r   r  r  r  r   r  s          r9   _cmpr  )  s     
16			CKKQV5			 ((.// K<ggWaCCC.",// Kfa+++
IAFIIII
J
JJr8   )r  r  r  c                T    t          ||g| j        | j        R  \  }} |||          S r^   )r   rE   rF   )rP   r   r   r   s       r9   signless_ruler  q  s5    !Q66666DAq2a88Or8   c                    | j         \  }}t          ||g| j         | j        R  \  }} |||t          j        |j        t          j                            S Nr   )rE   r   rF   rX   r  r   signedinteger)rP   r   r   r   r   r   s         r9   signed_ruler    sY    IFA!Q66666DAq2a3>&,8IJJKKKKr8   r$  fmthas_placeholdersc               ^    |rt          d          t          j        d| dd|           dS )NzFpl.debug_print() does not support placeholders when lowering to Triton F)hexr$  r7   )r   r{   print_)rP   r  r  r$  s       r9   debug_print_lowering_ruler    sH      
P   JJJJE5555	r8   r  attrir.AttributeNonec                V   t           j                            |           s|| j        j        |<   d S t          j        |           }|d|j         z  }|j        }|j        j        d         |k    }|sd S |j        j        x}r&t          |t          j
                  s||j        |<   d S d S d S )N_argr   )r   BlockArgumentrz   owner
attributes
arg_numberregionr   	operationr{   r   )r  r)   r  rs  r  is_entryr<  s          r9   	_set_attrr    s    			$	$Q	'	' #AGt
F
#
!
!
!!$
)%\ #u,(	 
FK!!b :b*:K+L+L BM$   r8   valuesSequence[int]c           
        | j         \  }t          dt          |j                            t          |          k    sJ t	          |dt
          j                            t          j	        |t          j
                                       |S )Nrt   r   rn  )rE   maxr   ro   r  r   DenseIntElementsAttrr}   r   asarrayrY   rP   r   r  r   s       r9   _multiple_of_ruler    sz    \(6	QFL!!	"	"c&kk	1	1	1	1!!"*V28"D"D"DEE  
 
(r8   c           
         | j         \  }t          |j                  t          |          k    sJ t          |dt          j                            t          j        |t          j	                                       |S )Nztt.contiguityrn  )
rE   r   ro   r  r   r  r}   r   r  rY   r  s       r9   _max_contiguous_ruler    sp    \(6	V\		c&kk	)	)	)	)!!"*V28"D"D"DEE  
 
(r8   c                P    | j         \  }t          t          ||          |          S r^   )rE   r   rU   )rP   r   ro   r   s       r9   _broadcast_to_ruler    s&    l)6	#Av..	6	66r8   c                  |dk    rt          |j        d          S |dk     }|r| }d }|dk    rEt          |d          \  }}|r||nt          ||          }|dk    rt          ||          }|dk    E|J | j        \  }| j        \  }t          ||j        |j                  }|rIt          j	        |j        t          j
                  }t          t          |j        d          ||          S |S )Nr   rt   r   r   )r  rc   divmodra   rE   rF   r  r   rX   r  r  r  )	rP   r   r   is_reciprocalaccmodr   r   r   s	            r9   _integer_pow_ruler    s   !VVa%- 	
A#	AAq\\FAs
 /AA$sA,,c1uu
q!**a 	
A 
\(6}*8c6<00# ^HNC,=>>FU38Q''V<<<<Jr8   funCallable[..., Any]r#  c               ,     r n fddfd}|S )Nc                      | i |fS r^   r7   )r$  kwr  s     r9   r  zlower_fun.<locals>.<lambda>  s    d9Ib9I9I8K r8   rP   rB   c                    t          j        |          }t          j        || j                  \  }}}\   t          j        ||          }t          | g|R d|i}	r|n|d         S )N
call_jaxprr   )rx  ry  r   rz  rE   rV   ClosedJaxpr_closed_call_lowering_rule)
rP   r$  r  wrapped_funrg   r   r  outr   r#  s
           r9   	f_loweredzlower_fun.<locals>.f_lowered  ss    ,r6**K4[#,OOE1fb //E
$S
B4
B
B
BE
B
BC".33A.r8   rP   rB   r7   )r  r#  r  r   s   `` @r9   	lower_funr    sJ     Kss$K$K$K$K"/ / / / / / / 
r8   c                R    t          j        t          j        | |          |          S r^   )rX   minimummaximum)minrm   r  s      r9   r  r    s    S[S!1D1Dc%J%J r8   c                8    ddt          j        |            z   z  S Nrt   )rX   r  )rm   s    r9   r  r    s    a1swr{{?3 r8   F)r#  c                   | j         \  }}t          ||g| j         | j        R  \  }}t          j        |j        t          j                  rt          j        ||          S t          j        |j        t          j	                  st          d|j         d|j                   t          j        |j        t          j                  rt          j        ||          S t          j        ||          S r  )rE   r   rF   rX   r  r   r  r  minnumfr  r   r  minsiminuirP   r   r   r   r   s        r9   _min_lowering_ruler
         <.&&	1	4s|	4cm	4	4	4$!Q^FL#,// ' A&&&	ck	2	2 
@v|@@&,@@   	^FL#"344 %q!$$$q!$$$r8   c                   | j         \  }}t          ||g| j         | j        R  \  }}t          j        |j        t          j                  rt          j        ||          S t          j        |j        t          j	                  st          d|j         d|j                   t          j        |j        t          j                  rt          j        ||          S t          j        ||          S r  )rE   r   rF   rX   r  r   r  r  maxnumfr  r   r  maxsimaxuir	  s        r9   _max_lowering_ruler    r  r8   c                   | j         \  }}t          ||g| j         | j        R  \  }}t          j        |j        t          j                  p#t          j        |j        t          j                  }t          j        |j        t          j                  s$t          j        |j        t          j                  rt          |||          S t          |||          S r  )rE   r   rF   rX   r  r   r  r   r  r  r   )rP   r   r   r   r   r   s         r9   _div_lowering_ruler  !  s    <.&&	1	4s|	4cm	4	4	4$!Q>&,(9:: cnlC%? ?& 	^FL"+.. )#.lBK3 3 ) Aq((((	1a	'	'	''r8   c           
     ^   | j         \  }t          j        |j        t          j                  }t          |j        d          }t          t          t          |||          t          j
        |j                  t          t          |||          t          j
        |j                            S )Nr   r   )rE   rX   r  r   r  r  rc   r  r  _greater_thanbool_
_less_than)rP   r   r   r   zeros        r9   _sign_lowering_ruler  /  s    \(6>&,(9::&	qvq		$	M!T&11139flKKJq$v...	6<HH
 
 r8   c                   t          d||                   }t          |t          j        |          }t	          t          |                    D ]}||k    rt          ||          }t          ||          S r  )_make_ranger  rX   rY   r   r   _expand_dimsr   )rP   r   ro   	dimensioniotar[   s         r9   _iota_lowering_ruler  :  so    	Qi(	)	)$	tSY	&	&$U # #aI~~$""d	4		r8   r  ir.Typec                v    t           j                            |           rt          j        |           j        S | S r^   )r   ry   rz   r   r  s    r9   rq  rq  D  s4    ##A&& q!!..Hr8   startendc                (   || k    rt          d| d|            t          | |          dk    rt          d          t          j        t          j                            || z
  gt          j                            d                    | |          S )Nz)end must be greater than start, but got: z <= l        zstart and end must fit in int32r   )	r   r  r{   
make_ranger   ry   r}   rV  get_signless)r!  r"  s     r9   r  r  K  s    E\\
DCDDUDD   	__
6
7
77		sU{mR^-H-H-L-LMM	
 
 r8   objectc                   t          |           }t          |t          j                  r#t	          j        |t          |                    }nDt          |t          j                  r#t	          j        |t          |                    }nt          t          j
                            |           rt          j        | |          S |S r^   )rq  rz   r   rV  r  constantr   r  r   r   ry   r{   r|   )r  r  r   r  s       r9   r  r  Y  s    q!!,bn-- #L#a&&99FF,-- #L%((;;FF
##A&& Av&&&Mr8   ir.valuec                    t           j                            | j                  rt	          d          |s| S t          j        t           j                            || j                  |           S )Nzcannot splat a tensor)r   ry   rz   rc   	TypeErrorr{   r|   r}   )r   ro   s     r9   _splatr,  h  sa    ##AF++ -
+
,
,,	 H		"-11%@@!	D	DDr8   c                   t           j                            | j                  sQt	          t          j        | j                  j                  }|                    |d           t          | |          S t          j	        | |          S r  )
r   ry   rz   rc   r  ro   insertr,  r{   expand_dims)r   r0  ro   s      r9   r  r  p  sq    			'	'	/	/ $QV,,233E	LLq!U		4	(	((r8   srcdst_typec                   t          j        t          | j                            }t          j        t          |                    }|j        dk    s|j        dk    r&t          j        || t
          j        j                  S |j        |j        k    rt          j
        ||           S |j        |j        k     rt          j        ||           S t          )N   )rounding)r   r  rq  rc   widthr{   fp_to_fpRoundingModeRTNEr  truncfextfr   )r0  r1  src_element_typedst_element_types       r9   _float_float_castr=  x  s    \-"9"9::\-"9"9::q  $4$:a$?$?(-   
 .444#... 0 666h,,,
r8   c                   t          j        t          | j                            }t          j        t          |                    }||k    sJ |j        dk    r%t          | t          | j        d          |          S |j        |j        k    rt          j        ||           S |j        |j        k    rt          j	        ||           S |r |j        dk    rt          j
        ||           S t          j        ||           S )Nrt   r   r   )r   rV  rq  rc   r5  
_not_equalr  r  bitcasttrunciextsiextuir0  r1  r   r;  r<  s        r9   _int_int_castrE    s    ^M#($;$;<<^M($;$;<<	-	-	-	-	-q  c51--f====/555 3/// 0 666#... ."(A--x---x---r8   c                  t          | j                  }t          |t          j        t          j        t          j        t          j        f          st          d|  d|           t          j	        t          |                    }|j
        dk    r%t          | t          | j        d          |          S |rt          j        ||           S t          j        ||           S )Ncannot cast  tp rt   r   r   )rq  rc   rz   r   BF16TypeF16Typer  r  r   rV  r5  r?  r  r  fptosifptouirD  s        r9   _float_int_castrM    s     #38,,	$r{BJ
BJ&W	X	X B
@S@@h@@
A
AA^M($;$;<<q  c51--f==== /#...#...r8   c               |   t          j        t          | j                            }t          |          }t	          |t           j        t           j        t           j        t           j        f          st          d|  d|           |j
        dk    s|st          j        ||           S t          j        ||           S )NrG  rH  rt   )r   rV  rq  rc   rz   rI  rJ  r  r  r   r5  r  uitofpsitofprD  s        r9   r  r    s     ^M#($;$;<<"8,,	bj"*bjI
 
 B @S@@h@@
A
AAq   #...#...r8   src_typejax.typing.DTypeLikec                z    t          | t          |          t          j        |t          j                            S r  )_ir_castr   rX   r  r  )r0  rQ  r1  s      r9   r  r    s;    
 
	!!^Hc&788
 
 
 r8   c                  t           j                            | j                  rct           j                            |          sDt          j        | j                  }t           j                            |j        ||j                  }| j        |k    r| S t          | j                  }t          |          }t          |t           j                  st          |t           j                  rt          d          t          |t           j
        t           j        f          rXt          |t           j                  s>t          t          | t           j                                        d          |d          S t          |t           j                  r*t          |t           j                  rt          | |          S t          |t           j                  r,t          |t           j                  rt#          | ||          S t          |t           j                  r,t          |t           j                  rt%          | ||          S t          |t           j                  r,t          |t           j                  rt'          | ||          S t(          j                            |          rt          |t           j                  r|j        dk    rt)          j        ||           S |j        dk    rft          | t           j                            d          |          }t3          |j        d          }t          t5          |||          ||          S t          |t           j                  r4t(          j                            |          rt)          j        ||           S t(          j                            |          r4t(          j                            |          rt)          j        ||           S t          d|  d|           )	Nz&cannot cast from or to float8_e4m3fnuzFr   @   rt   r   rG  rx   )r   ry   rz   rc   r}   ro   r   rq  Float8E4M3FNUZTyper   rJ  rI  r  rT  r  r=  rV  rE  rM  r  r{   r   r5  
ptr_to_intr%  r  r?  
int_to_ptrr@  )r0  r1  r   rQ  r;  r<  r   r  s           r9   rT  rT    s   ##	h  **844 "38,,H"&& H
 	XJ"38,,"8,, ""788 HJ-= = H F
G
GG 2:r{";<< Z
F F  bjnn&&u555x     ",// ,J5 5 , S(+++ ".11 7j7 7 7 hv6666 ",// 9J5 5 9 38888 ".11 9j7 7 9 38888&&'788 SZ> > S ##"8S111		1	$	$
333B77
G
G
Ga161djD888(6RRRR  0))*:;;0  3///&&  -))*:;;- h,,,>3>>H>>???r8   c               z    | j         \  }t          ||          }||j        k    r|S t          ||j        |          S r^   )rE   rU   r   r  )rP   r   	new_dtyper   shardingr   s         r9   #_convert_element_type_lowering_ruler]    sB     \(6q&!!!&,H	q&,		*	**r8   c                    | j         \  }}}| j        \  }t          |||||          \  }}t          |||||          \  }}t          j        |||          S r^   )rE   rF   r   r  select)rP   predr   r   	pred_avala_avalb_avalr   s           r9   select_n_lowering_rulerd    s`    !l)VV}*84Ivx88'$4Ivx88'$		dAq	)	))r8   c               *   t          |g| j        R  }t          j                            |j                  st          ||          S fdt          t          |                    D             }|D ]}t          ||          }t          ||          S )Nc                    g | ]}|v|	S r7   r7   )rZ   r[   broadcast_dimensionss     r9   r   z3_broadcast_in_dim_lowering_rule.<locals>.<listcomp>#  s$    OOOq:N1N1N1N1N1Nr8   )
rU   rE   r   ry   rz   rc   r   r   r   r  )rP   r   rg  ro   r/  rv   s     `   r9   _broadcast_in_dim_lowering_rulerh    s     q(3<(((!			'	'	/	/ QOOOOE#e**--OOO+  cQAA	1e		r8   c               *    ~t          | |d d           S )N)	new_sizes
dimensions)_reshape_lowering_rule)rP   rm   rk  s      r9   _squeeze_lowering_rulerm  )  s    	Q$4	H	H	HHr8   c                    |st          d          t          j        | j                  }t	          j        t          j                            ||j        |j                  | d          S )Nz cannot reshape to an empty shapeF)allow_reorder)	r   r   ry   rc   r{   reshaper}   r   r   )r   ro   tys      r9   _reshaperr  /  sg    	 9
7
8
88	16"""		eR_bkBB
 
 
 r8   c                  ~|t          d          S t          |g| j        R  }| j        \  }t          j                            |j                  s5t          d |j	        D                       sJ t          ||j	                  S g |j	        }dt          j                            |j                  rt	          j        |j                  j	        x}|k    rbt          |          k     r|         nd }t          |          k     r|         nd }||k    rdz  n|dk    rt          |          }dz  n|dk    r|}	t          fdt          |          D                       }
|                     | j        d                             |	          g| j        d                             |
          g          }t#          t$          j        ||f	          }nt)          ||          S t          j                            |j                  r%t	          j        |j                  j	        x}|k    b|S )
Nz`dimensions` is not supported.c              3  "   K   | ]
}|d k    V  dS rs   r7   )rZ   dim_sizes     r9   r\   z)_reshape_lowering_rule.<locals>.<genexpr>E  s&      <<x1}<<<<<<r8   r   rt   r5  c              3  .   K   | ]\  }}|k    |V  d S r^   r7   )rZ   didr[   s      r9   r\   z)_reshape_lowering_rule.<locals>.<genexpr>Z  s+      EEeb!R1WWWWWWEEr8   ru   )rE   rF   r  )r   rU   rE   rF   r   ry   rz   rc   r   ro   r,  r   r  rh   r   rH   update_reduce_loweringrX   r  rr  )rP   rm   rj  rk  r   	dst_shapea_shaperu  dst_dim_sizein_shape	out_shape
reduce_ctxr[   s               @r9   rl  rl  :  sQ    6777q(3<(((!}*8			'	'	/	/ %<<X^<<<<<<<<!X^$$$
  )!$$QV,,$)!&11777I
E
EW--wqzz4H#$s9~~#5#59Q<<4L<1faa			
qq
!
!
!a1faa	QhEEEEy'9'9EEEEEi;;LO***::;\!_++)+<<=   j 37J
=
=
=aaa###' 	$$QV,,$)!&11777I
E
E( 
(r8   root_ptrr&  BlockInfo | None
nd_indexer	NDIndexerarray_shapec                ^
   ||}d}|}n1|j         j        }t          d |j        D                       }|j        }t	          j        |          }|                                }|j        }	|j        }
|t          |	          d          }g }d}|d gt          |
          z  }n|j
        }t          |
          |z   t          |          k    sJ t          |          t          |          k    sJ t          |
          }t          |||          D ]\  }}}|t          j        u rt          d          }nt!          |          }t#          |t$          j                  r+|j        rWt+          t-          |j        |j        g          t3          t5          d|j                  |j        j        d                    }n|j        dk    rrt5          d|j                  }t+          t-          t          |j                  |j        g          t;          |t=          |j        |j                                      }n"t5          |j        |j        |j        z             }t          |	          |z   }t          |          |z
  dz
  }|dz  }nt#          |t>                    r_|t?          d           k    rtA          d          t5          d|          }t          |	          |z   }t          |          |z
  dz
  }|dz  }ntC          |tE          j#        dtH          j%                            }d}t          |          }tL          j'                            |j                  s!tQ          t          |          dz
  d          }nt          |          }|rFtL          j'                            |j                  s"tS          |dgt          |          z            }ngtU          |          D ]}tW          |d          }tU          |          D ]5}t          tY          |j        dg                     }tW          ||          }6t-          ||          }tM          j-        t]          |j                            }|0t3          ||d          }t+          |t-          ||                    }|j/        d	k    rt          |          }nta          |          }tS          ||          }|1                    t;          ||                     te          j3        t*          |t-          | |                    S )
Nr   c              3  2   K   | ]}|t           j        u V  d S r^   )r_   r`   )rZ   rd   s     r9   r\   z1_compute_pointers_from_indices.<locals>.<genexpr>r  s;        $%[     r8   Fr   rt   zOnly `slice(None)` allowed.r7   ro   r   )4r<   ro   sumr@   pallas_utilsstrides_from_shapeget_indexer_shapeint_indexer_shapeindicesr   r>   iterri   r_   r`   r   r  rz   r   Sliceis_dynamic_startr  r   r!  sizerT  r  rc   stridera   r  slicer   rU   rV   rW   rX   rY   r   ry   r  r,  r   r  r  rV  rq  r5  _i64_constantr   r   reduce)r  r&  r  r  
full_shapenum_mapped_dimsr@   stridesindexer_shaper  r  other_shapebcast_indicesother_shape_idxstart_index_offsetsindexer_iter
dim_stridedim_block_sizestart_offsetindexptr_dim_offsetr  num_left_expand_dimsnum_right_expand_dimsr   ndim
index_typestride_sizes                               r9   rQ  rQ  f  s2    JOKK,2J  )3)?    O (K+J77'..00- 2'c"344667+-/&3w<</$2	W	'3z??	:	:	:	:	 	!	!S__	4	4	4	4g,25{/3 3 F< F<.j., +++Aee<  e%)** '1		 LekEJ<00[EJ//1A%PPP
 
 <!1ej))mEK005:,??uTY5566
 

 %U[%+
2JKK !!233oE!+..@1Doo	E5	!	! 1	%++		!"?@@@"1n55n !233oE!+..@1Doo (
%b#)44 n !+.. ++N,?@@ 1"3}#5#5#91== #K 0 0 <R0;;N<OPP <nqcC4F4F.FGGnn)** 9 9!%na88*++ < <!7>.<<==%nd;;~}==Nn.A B BCCJlJuEEEl
)L-@@ n 2!*--kk!*--km44Knk::;;;;		
M9X}==
 
 r8   c               T   t          j        ||          }t          j                            |j                  st          |          dk    sJ |S t          |          dk    rt          d          |d         }t          j        ||fd d f          \  }}t          | g|R |d d ddS )Nr   rt   %No support for multiple indexers yet.F)r]  eviction_policycache_modifieris_volatile)
r
   tree_unflattenr{   r   rz   rc   r   r   tree_flatten_masked_load_lowering_rule)rP   r>  treerj   r_  indexerr^  r]  s           r9   _get_lowering_ruler    s    %dC00(			*	*38	4	4 x==AJ]]Q
E
F
FFQK'"/wj$0MNN)Y	#	

 
 
 
 
 r8   c                .    i | ]}t          |          |S r7   r(   )rZ   r.  s     r9   
<dictcomp>r    s     HHH3q661HHHr8   c                .    i | ]}t          |          |S r7   r  )rZ   cs     r9   r  r    s     FFF#a&&!FFFr8   )r  r  r  otherr  
str | Noner  r  c                  |t           j        j        }n,|dk    s|dk    rt          |         }nt	          d|           |t           j        j        }n0	 t          |         }n!# t          $ r t	          d|           d w xY wt           j	        
                    | j                  rLt          j	        | j                  }t          j        
                    |j                  rt          d          t!          | j                  }t           j	        
                    |          st	          d|           t          j	        |          }||t	          d          t          j        
                    | j                  sj|3t          j        
                    |j                  rt	          d          |3t          j        
                    |j                  rt	          d	          |j        }t          |t          j                  o
|j        d
k    }|rTt          j                            d          }t)          | t           j	                            ||j                  d          } |t)          ||d          }t          j        | |||||          }	|s|	n.t)          |	t          j                            d
          d          S )N.caz.cgunsupported cache modifier: unsupported eviction policy: -loading from a block pointer is not supportedunsupported pointer type: z"other requires mask to be providedz1other cannot be a block if pointer is not a block0mask cannot be a block if pointer is not a blockrt   r3  Fr   )r?  r  cacheevictr  )r{   CacheModifierNONE_STR_TO_CACHE_MODIFIERr   EvictionPolicyNORMAL_STR_TO_EVICTION_POLICYKeyErrorr   rz   rc   r   ry   rH  r   rq  rV  r5  r%  rT  r}   address_spaceload)
r>  r?  r  r  r  r  rJ  rH  is_int1r  s
             r9   _loadr    s    -2NN.E"9"9+N;NN
DNDD
E
EE /6OO/@oo   
;/
;
; 
 &&sx00 Q%ch//H	%%h&;<< Q OPPP38$$(			*	*8	4	4 >
<(<<
=
==#H--(
4<
9
:
::			'	'	1	1 KR0;;EJGGJKKKB/::49EEIJJJ&,|R^44P9Kq9P' >..q11L
""<1GHH  C UL777E?	  & JffFBN77::5IIIs   A$ $Bc                  |                     |          \  }}}}	|                     | j                  ^ }
}}t          |          dk    rt          d          |d         }t          j                            |j                  st          | j                  dk    sJ |S t          || j	        d         || j        d         j
                  }|0t          t          ||          |                                          }|	0t          t          |	|          |                                          }	t          |||	|||          S )Nrt   r  r   )r?  r  r  r  r  )rP  rE   r   r   r{   r   rz   rc   rQ  rG   ro   r   rU   r  r  )rP   r]  r  r  r  r^  r>  r_  r?  r  r   ra  
other_avalrj   s                 r9   r  r  /  sQ     )229==#xu'11#,??1i]]Q
E
F
FF#			*	*38	4	4 s|!!!!J&	3?1sCLO$9	 	# 
%dI668M8M8O8OPPD

++S-B-B-D-D E 
	#%
 
 
 r8   c               P   t          j        ||          }t          j                            |j                  st          |          dk    sJ |S t          |          dk    rt          d          |d         }t          j        ||f|d f          \  }}t          | g|R |d dS )Nr   rt   r  )r]  r  )
r
   r  r{   r   rz   rc   r   r   r  _masked_swap_lowering_rule)	rP   r>  valuer  rj   r_  r  r^  r]  s	            r9   _swap_lowering_ruler  S  s    %dC00(			*	*38	4	4 x==AJ]]Q
E
F
FFQK'"/wj%0NOO)Y	#	

 
!*D
 
 
 r8   )r  r  r  c                  |t           j        j        }n&|dk    rt          |         }nt	          d|           |t           j        j        }n0	 t          |         }n!# t          $ r t	          d|           d w xY wt           j	        
                    | j                  rLt          j	        | j                  }t          j        
                    |j                  rt          d          t!          | j                  }t           j	        
                    |          st	          d|           t          j	        |          }t          j        
                    | j                  sht          j        
                    |j                  rt	          d          |3t          j        
                    |j                  rt	          d          |j        }t          |t          j                  r_|j        dk    rTt          j                            d	          }t)          | t           j	                            ||j                  d
          } t)          ||d
          }t          j        | ||||          S )Nr  r  r  r  r  z1value cannot be a block if pointer is not a blockr  rt   r3  Fr   )r?  r  r  )r{   r  r  r  r   r  r  r  r  r   rz   rc   r   ry   rH  r   rq  rV  r5  r%  rT  r}   r  store)r>  r  r?  r  r  rJ  rH  s          r9   _storer  b  so    -2NN+N;NN
DNDD
E
EE /6OO/@oo   
;/
;
; 
 &&sx00 Q%ch//H	%%h&;<< Q OPPP38$$(			*	*8	4	4 >
<(<<
=
==#H--(			'	'	1	1 K	%%ej11 LJKKKB/::49EEIJJJ&,bn-- ,2D2I2I>..q11L
""<1GHH  C 5,u
5
5
5%			5t>
 
 
 s   A A<c               B   |                     |          \  }}}}|                     | j                  ^ }}	}
t          |          dk    rt          d          |d         }t	          || j        d         || j        d         j                  }d }|t          ||	          }|Tt          t          ||
          |	                                          }|"t          ||	                                          }t          |||          }t          ||||           |S )Nrt   r  r   )r?  r  )r?  r  )rP  rE   r   r   rQ  rG   ro   rU   r   r  r  r  )rP   r]  r  r^  r>  r_  r  r?  r   r`  ra  rj   r  	old_values                 r9   r  r    s'     )229==#x'11#,??1j)]]Q
E
F
FF#&	3?1sCLO$9	 	# %
UJ//E	%dI668M8M8O8OPPDs446677eCd%000)e$@@@@	r8   c                  t          j        ||          }t          j                            |j                  st          |          dk    sJ |S t          |          dk    rt          d          |d         }t          || j	        d         || j
        d         j                  }t          j        j        }t	          t          |j                  t          j                  rt          j        j        }t%          |||           g S )Nr   rt   r  )r
   r  r{   r   rz   rc   r   r   rQ  rG   rE   ro   rT  rW  rq  r   rV  rU  rL  )rP   r>  r  r  rj   r_  r  r<  s           r9   _addupdate_lowering_ruler    s    %dC00(			*	*38	4	4 x==AJ]]Q
E
F
FFQK'&		oa	l1o		 	# "ej))2>:: 			Bb#u	)r8   c               ,    t          j        ||          S r^   )r{   trans)rP   r   permutations      r9   _transpose_loweringr    s    		![	)	))r8   x_typeir.RankedTensorTypey_typeoptionsr   c                    d S r^   r7   )r  r  r  s      r9   _check_dot_operandsr    s	     	&r8   )
allow_tf32max_num_imprecise_accout_typer  r  r  
int | Noner  ir.Type | Nonec                  |t           j                                        }n,t          |t           j                  rt          d|           t          j        | j                  }t          j        |j                  }t          g |j	        |j	        R  dk     rt          d          |j        |j        k    rt          d|j         d|j                   t          ||t                                 |j        }t          |t           j                  r=|j        dk    rt!          d|           t           j                            d          }nGt          |t           j        t           j        f          rt           j                                        }n|}||k    rt!          d	| d
|           |j	        \  }	}
|j	        \  }
}|0t%          t           j                            |	|g|          d          }|.t          |t           j                  r|j        dk    rt
          d}t(          j        j        }|rt(          j        j        }t)          j        | ||||          S )Nzunsupported output type:    z(all dimensions of x and y must be >= 16 z2x and y must have the same element type, but got: r  r3  zunsupported element type: r   zoutput type z does not match element type r   )r  input_precision)r   r  r}   rz   rI  r   ry   rc   r  ro   r   r   r  r&  rV  r5  r+  r%  r  r  r{   InputPrecisionIEEETF32dot)r   r   r  r  r  r  r  r  r   mr   nr  s                r9   _dotr    sw    z~~HH(BK(( F
D(DD
E
EEqv&&&qv&&&&&,&&&&++
?
@
@@F///
	<	< 	<&,&9	< 	<  
 fffhh///$,bn-- QA<AABBB>..r22LL,R[ 9:: :>>##LLLX
LxLLlLL   
$!Q	$!Q[
#''A==q
A
AC",--  ,2D2I2I -2/ 5 /4O		1%
 
 
 r8   c          
        ~|\  \  \  }\  }}|dk    sJ |dk    rt          j        |d          }|dk    rt          j        |d          }|d}	n|\  }
}|
t          v p|t          v }	| j        \  }|j        x}}|t
          j        k    r.|t
          j        k    rt          j        t
          j                  }t          t          |||	t          |                    ||          S )N)r7   r7   r   )rt   r   rt   T)r  r  )r{   r  _TF32_PRECISIONSrF   r   rX   rY   r  r  r  r  r   )rP   rm   rd   dimension_numbers	precisionpreferred_element_typea_contract_dimb_contract_dim
batch_dimsr  prec_aprec_br   	out_dtype	acc_dtypes                  r9   _dot_general_loweringr    s    7H4(N'~*	x				qF##AqF##AJJNFF++Iv9I/IJ}*8".()i#)	S[ 8 8	#+&&I	


$Y//	   	
 	
 	r8   c                   t          j        |          }|\  }d |j        D             }t          j        ||f          }t	          j        t          j        |           |          \  }}	t          j	        |g ||          \  }
}}\    |	            }~|rt          d          d |D             }t          j        ||          }|dz  } |j        d         j        j        | }t           j                            |          5  t'          |j        |
d g|j        R  }t          j        |           d d d            n# 1 swxY w Y   |                                 t1          |j                  S )Nc                B    g | ]}t          j        d |j                  S )r7   )rV   rW   r   r  s     r9   r   z'_reduction_lowering.<locals>.<listcomp>G  s'    PPP4(&r4:66PPPr8   z(Reductions with constants not supported.c                6    g | ]}t          |j                  S r7   rp  rr  s     r9   r   z'_reduction_lowering.<locals>.<listcomp>S  rt  r8   r   r   )r
   ru  rE   rv  r   rw  rx  ry  r   rz  r   r{   ReduceOpr|  r   r   r   r   r   re   rC   r   reduce_returnr~  r  r  )r   rP   rm   r  r  r0  mapped_avalsr  r  r  r  r   r  r  r  	reduce_opr   r   r  s                      r9   _reduction_loweringr
  D  s   #A&&)'4PP3<PPP,$aV,,'%:l4' (N "$!:.,.." "-FB ^( J
H
I
II@@i@@@-!)T22)!+
,)
A

%
,k
:%	''.. & &&]D+0?  G W%%%	& & & & & & & & & & & & & & &
 	i		s   .D>>EEc                   t          |t                    sJ |s|S t          |          dk    rt          |          t          fd|j        D                       }t          | |                    |          |f          }t          t          |                    }|                    |          }t          fd|D                       }t          |          dk    t          | |||          d         S )Nrt   c              3     K   | ]:}|                     |j        d          |j        dz   d          z             V  ;d S )Nrt   ru   )rz  ro   )rZ   r  r0  s     r9   r\   z#_reduce_lowering.<locals>.<genexpr>f  sc       - - hhQWUdU^agdQhii6H%HhII - - - - - -r8   )rF   ry  rE   c              3  (   K   | ]}|k    |V  d S r^   r7   )rZ   axr0  s     r9   r\   z#_reduce_lowering.<locals>.<genexpr>p  s'      11bDjjjjjj11r8   r   )	rz   rh   r   r  rE   r{  rH   r  r
  )r   rP   rm   r  	dst_avalsr0  s        @r9   r{  r{  `  s    	D%	 	    	 HD		At99D - - - -"|- - - - -IckkIk..	A 	A 	AA 	vayyA
++y+
)
)C1111d11111D 	D		A 
T3	5	5	5a	88r8   c               R   |t           j        k    rt          d          t          |          dk    rt          d          |\  }|j        \  }t          d|j        |                   }t          |j                  dk    rOt          t          |j                            D ]}||k    rt          ||          }t          ||j                  }|
                    ||                    t          j        d                    g          }t          | |||f|          \  }	}
|
S )	Nz`index_type` must be i32.rt   z8`pallas` reduce operations only support one reduce axis.r   rY   rn  r  ry  )rX   rY   r   r   rE   r  ro   r   r  r   rH   rz  r   r
  )r   rP   rm   r  index_dtyper0  rb  r  r[   r   r  s              r9   _argreduce_loweringr    s    CI
0
1
11YY!^^
O
P
PP&4\(6
ad+
,
,%3v|$$%% ' '	
dUA&&eV\**Effmm#)G:L:Lm&M&MNOO#"4q%jtDDD*!W	.r8   c           	         | \  }}|\  }}||k    }||k     }t          j        ||          }t          j        ||t          j        |||                    }	t          j        ||          }
|
|	fS r^   )rX   r   wherer  leftrightvalue1index1value2index2gtlt	index_min	index_ret	value_rets              r9   _reduce_argmax_combiner"    t    .&&.&&""k&&)))iFCIb&)$D$DEE)k&&)))	I	r8   c           	         | \  }}|\  }}||k    }||k     }t          j        ||          }t          j        ||t          j        |||                    }	t          j        ||          }
|
|	fS r^   )rX   r   r  r  s              r9   _reduce_argmin_combiner%    r#  r8   c               Z    |j         rt          t          | j        |j        | j        g|R  S r^   )r  r   re   rC   rg   rG   rP   rg   r$  r   s       r9   _pjit_lowering_ruler(    s>    
\ 
	!	k5;
26
 
 
 r8   c               b    |j         |j        }}|rt          t          | j        || j        g|R  S r^   )rg   r  r   re   rC   rG   )rP   r  r$  r   rg   r  s         r9   r  r    s?    
 "J$5% 
	!#+uco	M	M	M	MMr8   c               4    t          | j        || j        g|R  S r^   )re   rC   rG   r'  s       r9   _remat_lowering_ruler+    s     	!#+uco	M	M	M	MMr8   c                    |S r^   r7   )r   r   s     r9   r  r    s    a r8   	axis_namer   c                   | j         j        j        }||v r$t          | |                    |                    S t          d| d          )Nr5  z
Axis name z not found in grid.)rC   r+   
grid_namesr8  r  LookupError)rP   r-  r/  s      r9   _axis_index_ruler1    sU    {'2**$Sz/?/?	/J/JKKKK????@@@r8   c                    t          |           dk    rdS t          |           dk    rdS | \  }t          |t          j                  S )Nr   Trt   F)r   rz   r   
ReadEffect)ref_effectseffs     r9   _is_read_onlyr6    sL    45&3	C)	*	**r8   c               `   ~|s|dk    rt           t          j        j        rt          nt
          } |d          } ||          }	 |d          }
t          t          || j                  }d | j        D             }t          j
        |ddg|          \  }\   d |j        D             }t          j        ||j                  dd          }t          t          |          }t          t           j        t          t           j        |          |          }t'          ||          \  }}t'          ||          \  }}t)          j        ||	|
|          t-          j        j                  5  j        }fdt5          |          D             }t7          |||          }t9          | j        |d g| j        |g|R  }t7          |||          }t'          ||          \  }}t)          j        |           d d d            n# 1 swxY w Y   t7          ||tA          j!                            S )	Nrt   r   c                D    g | ]}t          |t          j                   S r7   )rz   r   AbstractRefrZ   rm   s     r9   r   z&_for_lowering_rule.<locals>.<listcomp>  s5       /0*Q)
*
**  r8   r7   T)should_dischargec                    g | ]	}|j         
S r7   r  r  s     r9   r   z&_for_lowering_rule.<locals>.<listcomp>  s    +++af+++r8   c                @    g | ]\  }}j         j        |d z            S r   r   r   rZ   r[   r   for_ops      r9   r   z&_for_lowering_rule.<locals>.<listcomp>	  s8       )-Aa!e$  r8   )"r   r   
enable_x64r  r  r   r   rU   rE   r    discharge_stater   r   get_ref_state_effectseffectsr6  operatorand_not_r#   scf_dialectForOpr   r   r   induction_variabler   r"   re   rC   rG   yield_r  results_)rP   rg   which_linearnstepsr  unrollr$  _i_constantlower_boundupper_boundstep	init_argsr;  discharged_jaxprr  state_effects	read_onlyis_loop_argptrsr   non_loop_args	loop_args
loop_indexfor_body_argsloop_body_argsout_dischargedall_outloop_outr@  s                               @r9   _for_lowering_rulerb    s     !
!'!2!8Km+A+F##+	Q$"D#,77) 47L   #2R4";*:";  B ,+el+++(-hFFqrrJ- -//)mS	224D + +Y77'$+KCC-[+tYGG&	%% ! !*J   1:91E1E  M !m]KKN-	  	
 
  N *D.AAG g66KAxx   ! ! ! ! ! ! ! ! ! ! ! ! ! ! !  
[-fo1F1F	G	GGs   -B	HH	Hrt   )rS  
bound_typehas_loop_indexrS  rc  ir.IntegerType | Nonec                  |dk    rt           ||j        dk    rt          |          }nt          |          }t	          j        ||||          t          j                            j	                  5  j
        }	fdt          |          D             }
|r	g ||	|
}ng ||
}t          | j        || j        g|R  }t	          j        |           d d d            n# 1 swxY w Y   t!          j                  S )Nrt   r   c                @    g | ]\  }}j         j        |d z            S r   r>  r?  s      r9   r   z,_lower_jaxpr_to_for_loop.<locals>.<listcomp>.	  s+    NNNdaV[*1q51NNNr8   )r   r5  r   r  rH  rI  r   r   r   r   rJ  r   re   rC   rG   rK  r  rL  )rP   rg   rQ  rR  r  rd  rS  rc  r$  r\  r]  
jaxpr_argsr`  r@  s                @r9   _lower_jaxpr_to_for_loopri  	  sk    
QYY
:+r11DDD[+tTBB&	''44    *JNNNNiooNNNM -8V8Z8-8jj,V,m,j& 
	  G
 w                              
fo		s   6A!C##C'*C'c               `   ~t          |	          |z
  |z
  }
|
rt          |rt          |dk    rt          ~~
~~|j        |j        }}|rt          ~t	          j        |||          \  }}t          t          |	| j                  }	t          j
        |	|g          \  }}	|r3|	^}}	|}t          |t          ||j                            }|}|j        }n=t          d          }t          |          }t          j                            d          }t%          | ||||g|	R |d|d}|r|g|S |S )Nrt   r   r   rd  rS  rc  )r   r   rg   r  r  pattern_match_scan_to_fori_loopr   rU   rE   r   r$   r  rb   rc   r   r   rV  r%  ri  )rP   rg   linearlengthr  rO  
num_consts	num_carry_split_transposer$  num_extensivejaxpr_constsrd  r  lbrQ  ubrR  rc  for_outs                       r9   _scan_lowering_rulerw  =	  s}    d))j(94----'''q[[++mVWU\%,,, 25*iPP % 
tS\	2	2$
|44,&$ 	1IBK	b,vrw//	0	0BKJJ""K''K,,R00J$	5+{FD59D D#!
D D D'  # "'""	.r8   c               x   |rd S t          |j        j        |g          \  }}d |D             }t          |          dk     rd S |d d         \  }	}
|	j        dk    s|	j        t          j        t          j        fvrd S |
j        dk    s|
j        t          j        t          j        fvrd S |d d         \  }}|j        j	        d         }|j
        j        t          j        k    sJ t          |j        j                  dk    rd S |j        j        d         }|j        t          j        k    rd S |j	        |gk    rd S |j        ||gk    rd S t          |j        j        |g          \  }}|d d         \  }}|j        j	        d d         \  }}||urd S t!          |j        j                  D ]w\  }}|j        t          j        u r_|j        d         |u rPt%          |j        d         t&          j                  r+|j        d         j        dk    r|j	        d         |k    r|} nxd S |j        }g |j        d |         |j        |         |j        |dz   d          R }t-          |j	        dd                    }|                    |j        d |         |j        |dz   d          z   ||          }t          |||g          \  }}}|d d         |dd          c\  }}}t          | j        |g          \  }}|                     g |d |dd                    } t3          | ||||g|R dd|j        d	}||g|S )
Nc                    g | ]	}|j         
S r7   r  r  s     r9   r   z2_maybe_pattern_match_fori_loop.<locals>.<listcomp>z	  s    ///a16///r8   r   r7   r   rt   )r  r   r   )rG   Trk  )r$   rg   r   r   ro   r   rX   rY   r  r   r   r  r  r   r	   lt_pr   add_prz   rV   r  r	  rh   rH   rG   ri  rc   )rP   cond_nconsts
cond_jaxprbody_nconsts
body_jaxprr$  r   cond_invarscond_in_avalsa1a2v1v2outvarr'  body_invarsvo1vo2r[   	eqn_indexrg   
new_invarsnew_outvarsbody_constscarryrt  ru  const_block_infosargs_block_infosrv  s                                 r9   _maybe_pattern_match_fori_loopr  o	  s     4j.5~FF.![//;///-!4!&"bX^^rx	39'===4X^^rx	39'===4rr?&"b#A&&		ci	'	'	'	'		1$$4a #]ch4[VH4ZB84j.5~FF.![rr?&"b%bqb)(#ss]]4**/00 	 	fa
}	!!	A"		cjmX%566 	Z]!##{1~$$ie4

%2m|m, 2\*2lQ.//02 2* emABB'((+
--:jyj!EJy1}~~$>>   % %TL,+GHH![%!9eABBi.(2rD(23?4@>)C )C%% !8"3 !8T !8"2122"6!8 	9 	9#$	
 
 
 
 
 
' b	7	r8   c                  t          t          || j                  }t          | g|R ||||d}||S t	          j        |||g          \  }}}	t	          j        | j        ||g          \  }
}}d |D             }d |D             }d |	D             }g |||}t          j        ||          } |j	        j
        j        | }t	          j        |j        ||g          \  }}}g ||}t          j                            |          5  t!          | j        |j        g |
|g|R  \  }t          j        ||j                   d d d            n# 1 swxY w Y    |j        j
        j        | }t	          j        |j        ||g          \  }}}g |||}t	          j        |||g          \  }}}t          j                            |          5  t!          | j        |j        g ||g||R  }g |||}|rt          j        |           d d d            n# 1 swxY w Y   t-          |j                  } | ||z   d          S )N)r|  r~  r}  r  c                    g | ]	}|j         
S r7   rc   r:  s     r9   r   z(_while_lowering_rule.<locals>.<listcomp>	      222af222r8   c                    g | ]	}|j         
S r7   r  r:  s     r9   r   z(_while_lowering_rule.<locals>.<listcomp>	  r  r8   c                    g | ]	}|j         
S r7   r  r:  s     r9   r   z(_while_lowering_rule.<locals>.<listcomp>	  s    '''A'''r8   )r   rU   rE   r  r   r$   rG   rH  WhileOpbeforer   r   r   r   r   r   re   rC   rg   	conditionafterrK  r  rL  )!rP   r|  r}  r~  r  r$  r  cond_constsr  r  cond_const_block_infosbody_const_block_infoscarry_block_infoscond_const_typesbody_const_typescarry_types	all_typeswhile_opbefore_blockcond_consts_r   carry_	cond_argscondafter_blockbody_consts_all_argscond_const_argsbody_const_args
carry_argsra  all_handlesr`  s!                                    r9   _while_lowering_ruler  	  s    
tS\	2	2$ *# A A A<7CPZ5?A A A& M$(O
\<(% %!+{E ocol'CDD D02C 32k22222k222'''''+B B#3BkB) D11(.'.	:, O\" ,6 '&v&)	''55 8 8%5
 5#45 
	  FT $ 67778 8 8 8 8 8 8 8 8 8 8 8 8 8 8 -%,i8+'+\"( ($,f 5|4l4V4(15|,2 2./?J 
''44 
& 
&'5
 5#45 
	
 
  H BOAoAAK &%%%
& 
& 
& 
& 
& 
& 
& 
& 
& 
& 
& 
& 
& 
& 
& "##'	,..	//s$   :EEE>HH Hc               V  
 | j         }d 

fd| j        D             }t          |t          d|j                  d          }t          j        ||d          }t          j        	                    |j
                  5  t          | j        |d         j        |dd          g|R  }t          j        |           d d d            n# 1 swxY w Y   t          j        	                    |j                  5  t!          |          d	k    r;t#          | t%          |t          d|j                            g|R d
|dd          i}	n't          | j        |d         j        |dd          g|R  }	t          j        |	           d d d            n# 1 swxY w Y   t'          |j                  S )Nc                    t          | j                  }| j        s|S t          j                            | j        |          S r^   )r   r   ro   r   ry   r}   )r   r   s     r9   to_typez$_cond_lowering_rule.<locals>.to_type
  s<    $X^44L> ""8><@@@r8   c                &    g | ]} |          S r7   r7   )rZ   r  r  s     r9   r   z'_cond_lowering_rule.<locals>.<listcomp>
  s!    555wws||555r8   r   Fr   T)hasElsert   r   branches)rG   rF   _equalrb   rc   rH  IfOpr   r   r   
then_blockre   rC   rg   rK  
else_blockr   _cond_lowering_ruler  r  rL  )rP   r  r  r$  rG   	out_typesuse_branch0if_opouts0outs1r  s             @r9   r  r  
  s`    +A A A 6555s}555)ul1ej99%HHH+

;	4
@
@
@%	''(899  $ABB 
	  E
 u               
''(899  
8}}q!

ul1ej11
2
2    ABB<	 ee '
+
1+

abb/ 	  e
 u              " 
en		s%   ?<CCC6B
FFFr   c                    t          | t          j                  r| S t          | t          j        t          j        t          t          f          r"t          | t          |j
                            S t          r^   )rz   r   Valuer   r   r   r   r   rb   r   r   r   )r   r   s     r9   rU   rU   4
  s[    28 :H!biS%899 :,TZ88999r8   c                H   t          | t          j        t          j        t          t
          f          rjt          |t          j                  rt	          |           } n+t          |t          j                  sJ t          |           } t          j
        ||           S t          r^   )rz   r   r   r   r   r   r   rV  r  r  r(  r   )r  r  s     r9   rb   rb   <
  s~    BIrz3677 (!R^$$ 
a&&aa2<(((((
((a!!Q'''r8   c                f    t          j        t          j                            d          |           S )Nr   r  r(  r   rV  r%  r  s    r9   r   r   G
  %    		 ; ;B ? ?	C	CCr8   c                f    t          j        t          j                            d          |           S )NrV  r  r  s    r9   r  r  K
  r  r8   r   	jnp.dtypec                    t          j        | t          j                  r't          j                            | j        dz            S t          j	        |           S )Nr3  )
rX   r  r   r  r   rV  r%  itemsizer   dtype_to_ir_typern  s    r9   r   r   O
  sF    ^E2:&& ;>&&u~'9:::		u	%	%%r8   )rP   r'   rQ   rR   )rm   rn   ro   r?   rp   rn   )r   rn   r   rn   r   r   r   r   r   r   rp   rn   )r   r   rp   r   )r+   r*   )rp   r   )
rg   r   r+   r*   r)   r(   r0   r(   rp   rJ   )rP   r'   rg   r   rG   r  rp   r=   )r0  r   rp   rn   r  )r<  r=  r>  rn   r	  rn   r?  r@  rA  rB  rC  rD  rp   rn   )rP   rB   rM  rN  )rP   rB   r0  r   r  r  )r)   r(   r  r  rp   r  )r   rn   rp   rn   )r   rn   r   rn   )r   rn   r   rn   rp   rn   )r   rn   r   rn   r   r  rp   rn   )r   rn   r   rn   r  r  r  r  r  r  r   r  rp   rn   )rP   rB   r$  rn   r  r(   r  r  )r  rn   r)   r(   r  r  rp   r  )rP   rB   r  r  )rP   rB   ro   r  )rP   rB   r   r   )r  r  r#  r  rp   r  )r  r  rp   r  )r!  r   r"  r   rp   rn   )r  r  r  r&  rp   r  )r   r)  ro   r  rp   rn   )r   rn   r0  r   rp   rn   )r0  rn   r1  r  rp   rn   )r0  rn   r1  r  r   r  rp   rn   )r0  rn   rQ  rR  r1  rR  rp   rn   )r   rn   ro   r  rp   rn   )
r  rn   r&  r  r  r  r  r?   rp   rn   )NN)r>  rn   r?  r@  r  r@  r  r  r  r  r  r  rp   rn   r^   )r>  rn   r  rn   r?  r@  r  r  r  r  rp   rn   )r  r  r  r  r  r   )r   rn   r   rn   r  r@  r  r  r  r  r  r  rp   rn   )rP   rB   r-  r   )rp   r  )
rP   rB   rg   r   rd  r  rS  r   rc  re  )r   r&  r   r   rp   rn   )r  r&  r  r  rp   rn   )r  r   rp   rn   )r   r  rp   r  (A  rM   
__future__r   collections.abcr   r   r5   r   r   rE  typingr   r   r   jaxr	   r
   jax._srcr   r   r   r   r   rV   r   r   rx  r   r   r   r   jax._src.interpretersr   r   r   jax._src.lax.control_flowr   jax._src.lib.mlirr   jax._src.lib.mlir.dialectsr   r  r  r   rH  jax._src.lib.tritonr   r{   jax._src.pallasr_   r   r   r  jax._src.stater    r!   spjax._src.utilr"   r#   r$   	jax.numpynumpyrX   r   r%   safe_mapr   
unsafe_mapsafe_zipri   
unsafe_zipr  r*   rR   r   	dataclassr'   r;   rB   rJ   r"  rO   rl   r   r   r   r   r   r   r  re   r   program_id_pr8  num_programs_pr;  re  rf  rg  rh  rL  atomic_rmw_prb  atomic_cas_prl  r  cumsum_pr  not_pr  r  r  r  r  abs_pr  rz  neg_pceil_pfloor_pexp_pexp2_pexpm1_plog_plog1p_psqrt_ppow_pcbrt_prsqrt_psin_pcos_ptan_pasin_pacos_patan_patan2_psinh_pcosh_ptanh_pasinh_pacosh_patanh_ppopulation_count_pclz_pnextafter_pr  r  r  ra   r   r  r   r  r   CmpIPredicateeqCmpFPredicateOEQr  neUNEr?  sltultOLTr  sleuleOLE_less_equalsgtugtOGTr  sgeugeOGE_greater_equalr{  sub_pmul_pand_pandior_porixor_pr  shift_left_pshlishift_right_arithmetic_pshrsishift_right_logical_pshrui	add_any_p_JAX_TO_TRITON_BINARYitemsprimr   r  rem_peq_pne_pgt_pge_prz  le_p_JAX_TO_TRITON_SIGNED_BINARYr  debug_print_pr  r  multiple_of_pr  max_contiguous_pr  broadcast_to_pr  integer_pow_pr  r  clamp_p
logistic_p_JAX_FN_MAPPINGmin_pr
  max_pr  div_pr  sign_pr  iota_pr  rq  r  r  r,  r  r=  rE  rM  r  r  rT  convert_element_type_pr]  
select_n_prd  broadcast_in_dim_prh  	squeeze_prm  rr  	reshape_prl  rQ  get_pr  r  r  r  r  r  load_pr  swap_pr  r  r  addupdate_pr  transpose_pr  r  r  	PrecisionHIGHDEFAULTr  dot_general_pr  r
  r{  r  reduce_max_pr   reduce_min_pr  reduce_sum_pr  r"  argmax_pr%  argmin_ppjit_pr(  closed_call_pcustom_jvp_call_pr  remat_pr+  stop_gradient_paxis_index_pr1  r6  for_prb  ri  scan_prw  r  while_pr  cond_pr  rU   rb   r   r  r   r7   r8   r9   <module>r\     s%   7 6 " " " " " " . . . . . . . .           ) ) ) ) ) ) ) ) ) ) 



             " " " " " "                   % % % % % % ' ' ' ' ' ' & & & & & &       % % % % % %             & & & & & & 4 4 4 4 4 4 . . . . . .             = = = = = = ; ; ; ; ; ; 9 9 9 9 9 9 5 5 5 5 5 5 / / / / / / & & & & & & 1 1 1 1 1 1 $ $ $ $ $ $ # # # # # # + + + + + + % % % % % % ( ( ( ( ( ( $ $ $ $ $ $           WT]]-Z-Z	%'

                                            I          $   0     ; ; ; ;|   V, V, V, V,r9& 9& 9& 9&@) ) ) ) :*++' ' ' ,+' :,--+ + + .-+ !'1'='M*4*A*E    * :*++$. $. $. ,+$.N :*++   ,+(   @ 3<  A A A ! A 39E E E E
 d###       $#8 d###       $#   2 +*		:w//	<11\955[)44	
 		7)@@AA	7)@@AA	9+BBCC	9+BBCC	
   " 39
 
 
 
    KI''KJ$$GYKy99GYKi88

 GYK!2I>>GYK!2I>>

 
 
K K%%GYK	::GYKy99Iyk#G#GHHIzl$H$HII	
 GYK!3Y??GYK!3Y??Iyk#G#GHHIzl$H$HII	
  K8 I##GYKi88GYKY77Iyk#E#EFFIzl$F$FGG	
 Iyk#E#EFFIyk#E#EFFIyk#E#EFFIzl$F$FGG	
  9KV J$$GYKy99GYKi88Iyk#F#FGGIzl$G$GHH	
 GYK!2I>>GYK!2I>>Iyk#F#FGGIzl$G$GHH	
  WKt K%%GYK	::GYKy99

 GYK!3Y??GYK!3Y??

 
 
uKJ I##GYKi88GYKY77Iyk#E#EFFIzl$F$FGG	
 GYK!19==GYK!19==Iyk#E#EFFIzl$F$FGG	
  KKh K%%GYK	::GYKy99

 GYK!3Y??GYK!3Y??

 
 
iK~ J$$GYKy99GYKi88Iyk#F#FGGIzl$G$GHH	
 GYK!2I>>GYK!2I>>Iyk#F#FGGIzl$G$GHH	
  K\ I##GY(,	BBGY(+yAAGY	*KCCGY	*J	BB	
 GY(*;YGGGY(*;YGGGY	*,<iHHGY	*,<iHH	
  ]Kz J$$GYKy99GYKi88

 GYK!2I>>GYK!2I>>

 
 
{KP K%%GYK	::GYKy99

 GYK!3Y??GYK!3Y??

 
 
QKf I##GYKi88GYKY77Iyk#E#EFFIzl$F$FGG	
 GYK!19==GYK!19==Iyk#E#EFFIzl$F$FGG	
  gKD I##GYKi88GYKY77Iyk#E#EFFIzl$F$FGG	
 GYK!19==GYK!19==Iyk#E#EFFIzl$F$FGG	
  EKb I##GYKi88GYKY77

 GYK!19==GYK!19==

 
 
cKx J$$GYKy99GYKi88

 GYK!2I>>GYK!2I>>

 
 
yKN J$$GYKy99GYKi88

 GYK!2I>>GYK!2I>>

 
 
OKd J$$GYKy99GYKi88

 GYK!2I>>GYK!2I>>

 
 
 K%%GY	*M9EEGY	*L)DD

 GY	*,>	JJGY	*,>	JJ

 
 
 J$$GYKy99GYKi88

 GYK!2I>>GYK!2I>>

 
 
 J$$GYKy99GYKi88

 GYK!2I>>GYK!2I>>

 
 
 J$$GYKy99GYKi88

 GYK!2I>>GYK!2I>>

 
 
 K%%GYK	::GYKy99

 GYK!3Y??GYK!3Y??

 
 
 K%%GYK	::GYKy99

 GYK!3Y??GYK!3Y??

 
 
 K%%GYK	::GYKy99

 GYK!3Y??GYK!3Y??

 
 
 00GWI{G44GWI}g66

 Iwi!E!EFFIwi!E!EFF

 
 
 I##GWIz733GWI|W55

 Iwi!D!DEEIwi!D!DEE

 
 
 O))GY	*,=yIIGY	*,<iHH

 GY	*,BINNGY	*,BINN

 
 
A
K K K K K\
# # # #L L L L(< < < <I I I I
% 
% 
% 
%	I 	I 	I 	I
% 
% 
% 
%K K K K& 
	'*'*&*	
 
 
 Y'*'*&*	  
 Y'+'+&*	  
  i'+'+&*	   "	!'+'+&*	   #"'+'+&*	   ItItItI}!HmI}!m( -"5}2t  &++-- . .HD"79      !. ItHfHjHmHnHjHk   -2244 , ,HD"57 L L L L L
 !, :+,,   -,    :+,,   -, :.//   0/ 2$%%7 7 7 &%7
 3$%%   &%6     KJJN33
  %%'' F FHD" )	"u E E E 39% % % %" 39% % % %" 39
( 
( 
( 
( 3:    3:                E E E E) ) ) )   ". . . ."/ / / // / / /	 	 	 	A@ A@ A@ A@H 3-..+ + + /.+ 3>""* * * #"* 3)**	 	 	 +*	 3=!!I I I "!I
    3=!!( ( ( "!(Vh h h hV 28   & IHj.GHHH FFZ-EFFF 
 !!C
 "&"&C C C C C CL :$%%      &% F 29   " !3
 "&"&3 3 3 3 3 3l :$%%   &%0 2>""   #"* 3?##* * * $#*	 	 	 	  A
 (,#A A A A A AH M&(=>  3$%%& & & &%&R       89 9 9 9( +<)*;ck+ + c& ' +<)*;ck+ + c& ' +<)*;cg+ + c& '
   *   '8i&7/' ' cl #
   '8i&7/' ' cl #
 4;     8)**%788N N N 98 +*N =())N N N *)N 2@ g- . 3#$$A A A %$A+ + + + 8>""3H 3H 3H #"3H| (,! ! ! ! ! !H 3:. . . .bM M M M` 3;B0 B0 B0  B0J 3:+ + + +\      D D D DD D D D& & & & & &r8   