
    Vpfn                    (	   d Z ddlmZ ddlZddlZddlZddl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* ej+        Z+e#j,        Z,e#j-        Z-ej.        e/cZ/Z0ej1        e2cZ2Z3 ej4        d          Z5dxdZ6e5j7        dyd            Z8e5j9        d             Z: ej4        d          Z;dzd!Z<e;j7        dyd"            Z=e;j9        d#             Z> G d$ d%ej?                  Z@ ej4        d&          ZAd{d(ZB  e"jC        eA          eB           eAjD        d{d)            ZEdd*d|d-ZFdd*d}d.ZGdd*d}d/ZHdd*d}d0ZIdd*d}d1ZJdd*d}d2ZKdd*d}d3ZLdd*d}d4ZM ej4        d5          ZNeNjD        d6             ZOd7 ZP e"jC        eN          d8             ZQ ej4        d9          ZReRS                    d:             e'jT        eRd;            d< ZUeRj9        d=             ZV ej4        d>          ZWeWS                    d?             e'jT        eWd@            d~dDZXeWj9        dE             ZY ej4        dF          ZZeZjD        dG             Z[dH Z\e\ej]        eZ<   dI Z^e^ej_        eZ<   dJ Z`	 ddLZa e+eadMN          Zb e"jC        eZ          dO             Zc ej4        dP          ZdedjD        dQ             ZedR Zfefej]        ed<   dS Zgegej_        ed<    e"jC        ed          dT             ZhdddddKdUddVZidddWdXddYZjdddZdd\Zk	 	 dddbZl G dc ddejm                  Zn en            Zoejp        q                    en           ejr        q                    en           ejs        q                    en           ejt        q                    en            ej4        de          ZudMeu_v        ddjZwddkZxeujS        ddn            ZyeujD        ddp            Zzdq Z{ ej+        e{eu          ej|        eu<    ej+        e'jT        eu          dr             Z} ej4        ds          Z~dMe~_v        ddvZe~jD        dw             ZdS )zPallas-specific JAX primitives.    )annotationsN)AnyCallable)lax)	tree_util)ad_util)api_util)core)dtypes)effects)linear_util)pretty_printer)state)util)ad)batching)partial_eval)	discharge)indexing)
primitives)mlir
program_idaxisintreturn	jax.Arrayc                8    t                               |           S )aH  Returns the kernel execution position along the given axis of the grid.

  For example, with a 2D `grid` in the kernel execution corresponding to the
  grid coordinates `(1, 2)`,
  `program_id(axis=0)` returns `1` and `program_id(axis=1)` returns `2`.

  Args:
    axis: the axis of the grid along which to count the program.
  r   )program_id_pbindr   s    Z/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/_src/pallas/primitives.pyr   r   7   s     
				%	%%    c                    t          j                    }|r||          j        S t          j                    }|                    |           }t
          j                            t          |           S Nr   )	pallas_corecurrent_grid_envindex
axis_framesizejax_core	Primitiver    r   )r   grid_envframe_s       r!   program_id_bindr/   C   sb    )++(  D>

 
"
"% jj!			 	 D	 	9	99r"   c                 @    t          j        dt          j                  S N r*   ShapedArrayjnpint32r.   s    r!   _program_id_abstract_evalr8   N       		b#)	,	,,r"   num_programsint | jax.Arrayc                8    t                               |           S )z2Returns the size of the grid along the given axis.r   )num_programs_pr    r   s    r!   r:   r:   T   s    			$		'	''r"   c                   t          j                    }|r||          j        S t          j                    }|                    |           }|t           j        u r&t
          j                            t          |           S |S r$   )	r%   r&   r)   r(   dynamic_grid_dimr*   r+   r    r=   )r   r,   r-   r)   s       r!   _num_programs_bindr@   X   su     )++( D>

 
"
"%	D		$	[)))"">"===	+r"   c                 @    t          j        dt          j                  S r1   r3   r7   s    r!   _num_programs_abstract_evalrB   e   r9   r"   c                  *    e Zd ZdZdZdZdZdZdZdZ	dS )	AtomicOpTypexchgaddmaxminandorxorN)
__name__
__module____qualname__XCHGADDMAXMINANDORXORr2   r"   r!   rD   rD   i   s1        	$####"###r"   rD   
atomic_rmwatomic_typec                  ~|                     |          \  }}}}t          |          dk    rt          d          |d         }	|t          |t          j        k    rd }
nI|t          j        k    rt          j        }
n,|t          j        k    rt          j	        }
nt          |          t          d |	j        D                       r|	j        }d |D             }d |D             }t          d |D                       }t          j        |||	          }t          d
 |D                       }||         } |
||          }t          j        |||          }t          d |D                       }||         }nbt          d |	j        D                       r=||	j                 }|j        |	j                                      |
||                    }nt          |fdt          |           dz
  z  z   |fS )N   zOnly one indexer is supported.r   c                    | |z   S Nr2   )xys     r!   <lambda>z,_atomic_rmw_discharge_rule.<locals>.<lambda>   s
    !a% r"   c              3  P   K   | ]!}t          |t                    p|j         V  "d S r[   
isinstanceSliceshape.0ss     r!   	<genexpr>z-_atomic_rmw_discharge_rule.<locals>.<genexpr>   5      DD1*Q


-ag+DDDDDDr"   c                P    g | ]#}t          |t                     o
|j        d k    $S )r2   r`   rd   s     r!   
<listcomp>z._atomic_rmw_discharge_rule.<locals>.<listcomp>   s0    OOO!z!U+++=2OOOr"   c                J    g | ] }t          |t                    r|j        n|!S r2   ra   rb   startrd   s     r!   rj   z._atomic_rmw_discharge_rule.<locals>.<listcomp>   -    LLLqz!U33:AGGLLLr"   c              3  R   K   | ]"}t          |t                    r|j        nd V  #dS rY   Nra   rb   r)   rd   s     r!   rg   z-_atomic_rmw_discharge_rule.<locals>.<genexpr>   7      OO!*Q"6"6=AOOOOOOr"   slice_sizesc              3  <   K   | ]}|rd nt          d           V  d S r[   slicere   scalars     r!   rg   z-_atomic_rmw_discharge_rule.<locals>.<genexpr>   s1      RRF7E$KKRRRRRRr"   start_indicesc              3  <   K   | ]}|rd nt          d          V  dS r   Nrv   rx   s     r!   rg   z-_atomic_rmw_discharge_rule.<locals>.<genexpr>   1      OOV4tOOOOOOr"   c              3  B   K   | ]}t          |t                     V  d S r[   ra   rb   rd   s     r!   rg   z-_atomic_rmw_discharge_rule.<locals>.<genexpr>   /      
9
9z!U###
9
9
9
9
9
9r"   r[   )	unflattenlenNotImplementedErrorrD   rP   rQ   r5   maximumrR   minimumallindicestupler   dynamic_slicedynamic_update_sliceatset)in_avals	out_avals	args_treerW   	args_flatrefindexersvalmaskidxmonoidr   scalar_dimsslice_startsrt   out_onesval_indexerx_newout_indexerouts                       r!   _atomic_rmw_discharge_ruler   u   s"    &00;;#xd]]Q
>
?
??#	
L$$$FFl&&&[FFl&&&[FF
k
*
**DDDDDDD kGOOwOOOKLLGLLLLOOwOOOOOK lLLLHRRkRRRRRK
k
C
&h

C$S#\JJJEOO;OOOOOK
;
CC

9
9S[
9
9
999 
ck
CF3;##FF3$4$455EE

Gs8}}q01	13	66r"   c                   |                      |          \  }}}}|j        t          j        d          k    r(|t          j        k    rt          d|j         d          |j        t          j        d          t          j        d          t          j        d          t          j        hv r t          d|j         d|j         d          t          |d	| iS )
Nfloat16z`atomic_z` does not support f16.boolint8int16z` does not support .r   )	r   dtyper5   rD   rP   
ValueErrorvaluebfloat16_swap_abstract_eval)r   rW   
avals_flatr   r.   s        r!   _atomic_abstract_evalr      s    $$Z00,#q!QY#)I&&&&;,:J+J+J
J 1JJJ
K
KKY	i	i	i	l	   E;$EEEEE   
j	>I	>	>>r"   )r   r   
Any | Nonec                   t          j        | |d          \  }}t          j        ||||f          \  }}t	          j        |||dS )NrV   )r   rW   )spget_ref_and_indexersr   tree_flattenatomic_rmw_pr    )	x_ref_or_viewr   r   r   rW   x_refr   r   r   s	            r!   _atomic_rmwr      sY    +M3MM/%"/#t0LMM)Y		I;
 
 
 r"   c               >    t          | |||t          j                  S )zAtomically exchanges the given value with the value at the given index.

  Args:
    x_ref_or_view: The ref to operate on.
    idx: The indexer to use.
    mask: TO BE DOCUMENTED.

  Returns:
    The value at the given index prior to the aupdate.
  r   rW   )r   rD   rO   r   r   r   r   s       r!   atomic_xchgr      s)     
S#Dl6G
 
 
 r"   c               >    t          | |||t          j                  S )zAtomically computes ``x_ref_or_view[idx] += val``.

  Args:
    x_ref_or_view: The ref to operate on.
    idx: The indexer to use.
    mask: TO BE DOCUMENTED.

  Returns:
    The value at the given index prior to the atomic operation.
  r   )r   rD   rP   r   s       r!   
atomic_addr      )     
S#Dl6F
 
 
 r"   c               >    t          | |||t          j                  S )a  Atomically computes ``x_ref_or_view[idx] = max(x_ref_or_view[idx], val)``.

  Args:
    x_ref_or_view: The ref to operate on.
    idx: The indexer to use.
    mask: TO BE DOCUMENTED.

  Returns:
    The value at the given index prior to the atomic operation.
  r   )r   rD   rQ   r   s       r!   
atomic_maxr      r   r"   c               >    t          | |||t          j                  S )a  Atomically computes ``x_ref_or_view[idx] = min(x_ref_or_view[idx], val)``.

  Args:
    x_ref_or_view: The ref to operate on.
    idx: The indexer to use.
    mask: TO BE DOCUMENTED.

  Returns:
    The value at the given index prior to the atomic operation.
  r   )r   rD   rR   r   s       r!   
atomic_minr      r   r"   c               >    t          | |||t          j                  S )zAtomically computes ``x_ref_or_view[idx] &= val``.

  Args:
    x_ref_or_view: The ref to operate on.
    idx: The indexer to use.
    mask: TO BE DOCUMENTED.

  Returns:
    The value at the given index prior to the atomic operation.
  r   )r   rD   rS   r   s       r!   
atomic_andr      r   r"   c               >    t          | |||t          j                  S )zAtomically computes ``x_ref_or_view[idx] |= val``.

  Args:
    x_ref_or_view: The ref to operate on.
    idx: The indexer to use.
    mask: TO BE DOCUMENTED.

  Returns:
    The value at the given index prior to the atomic operation.
  r   )r   rD   rT   r   s       r!   	atomic_orr   
  s(     
S#Dlo
 
 
 r"   c               >    t          | |||t          j                  S )zAtomically computes ``x_ref_or_view[idx] ^= val``.

  Args:
    x_ref_or_view: The ref to operate on.
    idx: The indexer to use.
    mask: TO BE DOCUMENTED.

  Returns:
    The value at the given index prior to the atomic operation.
  r   )r   rD   rU   r   s       r!   
atomic_xorr     r   r"   
atomic_casc                L   |j         |j         k    s|j        |j        k    rt          d          | j        rt          d          |j        rt          d          |j        rt          d          t          j        |j        |j                   t          j        d          hfS )Nz1cmp and val must have identical dtypes and shapeszref must be scalar.zcmp must be scalar.zval must be scalar.r   )r   rc   r   r*   r4   r   WriteEffect)ref_avalcmp_avalval_avals      r!   _atomic_cas_abstract_evalr   +  s    ^x~%%8>)I)I
H
I
II^ ,
*
+
++^ ,
*
+
++^ ,
*
+
++		hnhn	=	=@QRS@T@T?U	UUr"   c                :    t                               | ||          S )a  Performs an atomic compare-and-swap of the value in the ref with the
  given value.

  Args:
    ref: The ref to operate on.
    cmp: The expected value to compare against.
    val: The value to swap in.

  Returns:
    The value at the given index prior to the atomic operation.
  )atomic_cas_pr    )r   cmpr   s      r!   r   r   8  s     
		3S	)	))r"   c                H    ~ ~t          j        ||k    ||          }|d d f|fS r[   )r5   where)r   r   r   r   r   new_vals         r!   _atomic_cas_discharge_ruler   F  s0    	IcSj#s++'
4		##r"   max_contiguousc                    | S r[   r2   r\   r.   s     r!   r^   r^   N  s     r"   c                    |gS r[   r2   r.   r\   __s      r!   r^   r^   O  s    QC r"   c                j    t          |t                    s|g}t                              | |          S N)values)ra   listmax_contiguous_pr    r\   r   s     r!   r   r   Q  s4    	FD	!	! XF			q		0	00r"   c                    | S r[   r2   avalr.   s     r!   _max_contiguous_abstract_evalr   V      	+r"   multiple_ofc                    | S r[   r2   r   s     r!   r^   r^   \  s    a r"   c                    |gS r[   r2   r   s      r!   r^   r^   ]  s    ! r"   r\   r   list[int] | intc                j    t          |t                    s|g}t                              | |          S r   )ra   r   multiple_of_pr    r   s     r!   r   r   _  s4    	FD	!	! XF			Af		-	--r"   c                    | S r[   r2   r   s     r!   _multiple_of_abstract_evalr   d  r   r"   masked_loadc                    |                      |          \  }}}}t          j        |d                                         |j                  t          j        d          hfS )Nr   )r   r*   r4   get_indexer_shaper   r   
ReadEffect)r   r   r.   r   r   s        r!   _load_abstract_evalr   k  s[    !++J77#xA8B<99;;SYGG
 r"   c           	     l   | j         \  }t          j        | j        d         | j                  \  }}}}t          j        |g||j                  }|t          j	        d          t          j        |||          g}	|Q|	t          j	        d          t          j	        d          t          j	        t          j        ||                    gz  }	|Q|	t          j	        d          t          j	        d          t          j	        t          j        ||                    gz  }	t          j        |	          S )Nr   print_shapes <-  mask=zother=)outvarsr   tree_unflattenparamsinvarsr*   pp_varsr   pptextr   pp_ref_indexerspp_varconcat)
eqncontextsettingsr]   r\   r   r   otherlhsresults
             r!   _load_pp_ruler
  t  s   
{"!'6sz+7N7:zC C!XtU 	!gH4IJJJ#	gfoo!X..&
 




g..// F
 



w//00 F
 
6		r"   c           	     :   |                     |           \  }}}}|                     |          \  }}	}	}
|
t          j        |
          }
t          j        t          j        ||||f          d|i|t          j        t          j        ||||
f          d|i|fS Nr   )r   r   instantiateload_pr    r   tree_leaves)primalstangentsr   r   
ref_primalr   r   other_primalref_tangentr.   other_tangents              r!   	_load_jvpr    s    -6-@-@-I-I**hl%.%8%8%B%B"+q!]'66Mk *hl!KLL  
 k +x}!MNN  
 r"   c                   t          j        |t           j                  r t          j        | t           j        |          S t          j        |t           j                  r-t          j        | t          j        |          j        |          S t          j        |t           j                  rt          j        | d|          S t          |          )NF)
r5   
issubdtypefloatingfullnanintegeriinforH   r   r   rc   r   s     r!   uninitialized_valuer    s    ^E3<(( )8E37E***
~eS[)) )8E39U++/777
~eSX&& )8E5%(((E"""r"   Fc                    t          d |D                       }|rt          d |D                       }t          d| j                  }t          j        | ||          } | S )a  
  DynamicSlice and DynamicUpdateSlice adjust the start index in cases where the
  requested slice overruns the bounds of the array. This pads the array with
  uninitialised values such that the requested slice will never overrun.

  For example, if arr is [1.,2.,3.,4.] and a slice of size 4, start index 2 is
  requested then the result will be [3.,4.,NaN,NaN] after padding, rather than
  [1.,2.,3.,4.] from the unpadded array

  unpad=True performs the inverse operation
  c              3      K   | ]	}d |d fV  
dS r}   r2   )re   
slice_sizes     r!   rg   z?_pad_values_to_avoid_dynamic_slice_oob_shift.<locals>.<genexpr>  s)      JJ
!Z+JJJJJJr"   c              3  .   K   | ]\  }}}| | | fV  d S r[   r2   )re   lowhighinteriors       r!   rg   z?_pad_values_to_avoid_dynamic_slice_oob_shift.<locals>.<genexpr>  sQ       H H4T8 !D4%(3 H H H H H Hr"   r2   r  )padding_configpadding_value)r   r  r   r   pad)r   rt   unpadr'  r(  s        r!   ,_pad_values_to_avoid_dynamic_slice_oob_shiftr+    s     JJkJJJJJ.
 H H H8FH H H H HN%BekBBB-
'%!/ -/ / /% 
,r"   T)r*  c               T   ~|                     |          \  }}}}t          |          dk    rt          d          |d         }	t          d |	j        D                       r|	j        D ]1}
t          |
t                    r|
j        dk    rt          d          2|	j        }d |D             }d |D             }t          d |D                       }t          ||          }t          j        t          j                  t          j        |fd	|D             |
          }t          d |D                       }||         }n3t          d |	j        D                       r||	j                 }nt          ||t          j        |||          }dt          |           z  |fS )NrY   -Only one indexer supported in discharge rule.r   c              3  P   K   | ]!}t          |t                    p|j         V  "d S r[   r`   rd   s     r!   rg   z'_load_discharge_rule.<locals>.<genexpr>  rh   r"   Unimplemented stride support.c                J    g | ] }t          |t                     o|j         !S r2   r`   rd   s     r!   rj   z(_load_discharge_rule.<locals>.<listcomp>  s.    MMMz!U+++;AGMMMr"   c                J    g | ] }t          |t                    r|j        n|!S r2   rl   rd   s     r!   rj   z(_load_discharge_rule.<locals>.<listcomp>  rn   r"   c              3  R   K   | ]"}t          |t                    r|j        nd V  #dS rp   rq   rd   s     r!   rg   z'_load_discharge_rule.<locals>.<genexpr>  rr   r"   c                :    g | ]}t          j        |          S r2   )r5   astype)re   rf   	idx_dtypes     r!   rj   z(_load_discharge_rule.<locals>.<listcomp>  s%    666Asz!Y666r"   rs   c              3  <   K   | ]}|rd nt          d          V  dS r}   rv   rx   s     r!   rg   z'_load_discharge_rule.<locals>.<genexpr>  r~   r"   c              3  B   K   | ]}t          |t                     V  d S r[   r   rd   s     r!   rg   z'_load_discharge_rule.<locals>.<genexpr>  r   r"   r[   )r   r   r   r   r   ra   rb   strider   r+  r   canonicalize_dtyper5   int64r   r   r   )r   r   r   r   r.   r   r   r   r  r   rf   r   r   r   rt   r   r   r   r5  s                     @r!   _load_discharge_ruler;    s   (229==#xu]]Q
M
N
NN#DDDDDDD [ C C	Au		 C!(Q,,!"ABBBkGMMWMMMKLLGLLLLOOwOOOOOK 7sK
H
HC)#)44I 	6666666  H
 OO;OOOOOK
;
CC

9
9S[
9
9
999 
ck
CC
	%+
)D#u
%
%C	3x==	 #	%%r"   masked_swapc           	        |                      |          \  }}}}|d                                         }||j        k    r#t          d|j         d|j         d| d          |j        |j        k    r t          d|j         d|j         d          t          j        ||j                  t          j        d          hfS )	Nr   z%Invalid shape for `swap`. Ref shape: z. Value shape: z. Indices: z. z%Invalid dtype for `swap`. Ref dtype: z. Value dtype: r   )	r   r   rc   r   r   r*   r4   r   r   )r   r   r.   r   r   r   expected_output_shapes          r!   r   r     s   #--j99#xa"2,88::ci''
	;	 	; 	;		; 	;.6	; 	; 	;   	Y#)
	&	 	& 	&		& 	& 	&  
 0#)<<
 r"   c                4   | j         \  }| j        d                             | j                  \  }}}}t	          j        |||          }t          |t          j                  rNt          j
        |t          j        d          t          j        t          j        ||                    g          S t          j        |g||j                  }|t          j        d          |t          j        d          |t          j        d          t          j        t          j        ||                    g}	|Q|	t          j        d          t          j        d          t          j        t          j        ||                    gz  }	t          j
        |	          S )Nr   r   r   z, r   r   )r   r   r   r   r   r  ra   r*   DropVarr   r  r   r  r   r   )
r  r  r  r]   r\   r   r   r   x_ir	  s
             r!   _swap_pp_rulerB    sa    {"!:k2<<SZHH!XsD
7Ax00#8#$$ B9
g!>!>??A B B B sG(2GHHH!gdmm	gfoo	gdmmghoc7++,,& 




g..// F
 
6		r"   c          	     6   |                     |           \  }}}}|                     |          \  }}	}
}	t          j        |
          }
t          j        t          j        ||||f          d|i|t          j        t          j        |||
|f          d|i|fS r  )r   r   r  swap_pr    r   r  )r  r  r   r   r  r   
val_primalr   r  r.   val_tangents              r!   	_swap_jvprG  &  s    +4+>+>w+G+G(*h
D#,#6#6x#@#@ +q+q#K00+k *h
D!IJJ  
 k +xd!KLL  
 r"   c               ^   ~|                     |          \  }}}}t          |          dk    rt          d          |d         }	t          d |	j        D                       r'|	j        D ]1}
t          |
t                    r|
j        dk    rt          d          2|	j        }d t          |          D             }d |D             }t          d |D                       }t          ||          }t          j        |||	          }t          j        ||          }|.|}t          j        |||          }t          j        |||          }t          j        ||          }t          j        |||
          }t%          ||          }nt          d |	j        D                       rc||	j                 }|.|}t          j        |||          }t          j        |||          }|j        |	j                                     |          }nt          |fdt          |           dz
  z  z   |fS )NrY   r-  r   c              3  P   K   | ]!}t          |t                    p|j         V  "d S r[   r`   rd   s     r!   rg   z'_swap_discharge_rule.<locals>.<genexpr>B  rh   r"   r/  c                P    g | ]#\  }}t          |t                    |j        !|$S r2   r`   )re   irf   s      r!   rj   z(_swap_discharge_rule.<locals>.<listcomp>H  sH       Aq!U## -.G	  r"   c                J    g | ] }t          |t                    r|j        n|!S r2   rl   rd   s     r!   rj   z(_swap_discharge_rule.<locals>.<listcomp>M  rn   r"   c              3  R   K   | ]"}t          |t                    r|j        nd V  #dS rp   rq   rd   s     r!   rg   z'_swap_discharge_rule.<locals>.<genexpr>N  rr   r"   rs   rz   c              3  B   K   | ]}t          |t                     V  d S r[   r   rd   s     r!   rg   z'_swap_discharge_rule.<locals>.<genexpr>\  r   r"   r[   )r   r   r   r   r   ra   rb   r8  	enumerater   r+  r   r   r5   squeezer   expand_dimsr   ._unpad_values_to_avoid_dynamic_slice_oob_shiftr   r   )r   r   r   r   r.   r   r   r   r   r   rf   r   r   r   rt   r   out_r   s                     r!   _swap_discharge_rulerT  ;  sQ   &00;;#xd]]Q
M
N
NN#DDDDDDD "[ C C	Au		 C!(Q,,!"ABBBkG g&&  K
 MLGLLLLOOwOOOOOK 7sK
H
HC

C;
G
G
GC
+c;
'
'CdIdC%%cIdC&&c
/#{
+
+C$S#\JJJE:5+NNEE

9
9S[
9
9
999 
ck
CdIdC%%cIdC&&cF3;##C((EE

Gs8}}q01	13	66r"   )r   r  cache_modifiereviction_policyvolatilec                   t          j        | |d          \  }}t          j        ||||f          \  }	}
t	          j        |	|
|||dS )av  Returns an array loaded from the given index.

  If neither ``mask`` nor ``other`` is specified, this function has the same
  semantics as ``x_ref_or_view[idx]`` in JAX.

  Args:
    x_ref_or_view: The ref to load from.
    idx: The indexer to use.
    mask: An optional boolean mask specifying which indices to load.
      If mask is ``False`` and ``other`` is not given, no assumptions can
      be made about the value in the resulting array.
    other: An optional value to use for indices where mask is ``False``.
    cache_modifier: TO BE DOCUMENTED.
    eviction_policy: TO BE DOCUMENTED.
    volatile: TO BE DOCUMENTED.
  load)r   rU  rV  is_volatile)r   r   r   r   r  r    )r   r   r   r  rU  rV  rW  r   r   r   r   s              r!   rY  rY  h  sb    $ +M3GG/%"/$0NOO)Y	#%
 
 
 r"   swapr   rV  _function_namec                   t          j        | ||          \  }}t          j        ||||f          \  }}	t	          j        ||	|dS )zSwaps the value at the given index and returns the old value.

  See :func:`~jax.experimental.pallas.load` for the meaning of the arguments.

  Returns:
    The value stored in the ref prior to the swap.
  )r   rV  )r   r   r   r   rD  r    )
r   r   r   r   rV  r]  r   r   r   r   s
             r!   r[  r[    sZ     +M3OO/%"/#t0LMM)Y	I
 
 
 r"   )r   rV  Nonec               0    t          | ||||d          }dS )ztStores a value at the given index.

  See :func:`~jax.experimental.pallas.load` for the meaning of the arguments.
  storer\  N)r[  )r   r   r   r   rV  r.   s         r!   ra  ra    s)    
 =#s!# # #!!!r"   trans_ar   trans_b
allow_tf32bool | Nonec                :   | j         dk    s|j         dk    rt          d          |rdnd}|sdnd}|5|t          d          |rt          j        j        nt          j        j        }t          j                            | ||f|ffdf|t          j	                  S )N   z`a` and `b` must be 2D arrays.r   rY   z5Only one of allow_tf32 and precision can be specified)r2   r2   )dimension_numbers	precisionpreferred_element_type)
ndimr   r   	PrecisionHIGHHIGHESTjaxdot_generalr5   float32)abrb  rc  rd  ri  lhs_contract_dimrhs_contract_dims           r!   dotrv    s    fkkqv{{
5
6
66!(QQq%,QQ1NOOO&0K""cm6KI			+-0@/BCXN [ 
 
 
 r"   c                      e Zd Zd ZdS )PrintEffectc                    dS )NPrintr2   )selfs    r!   r^   zPrintEffect.<lambda>  s     r"   N)rL   rM   rN   __str__r2   r"   r!   rx  rx    s          '''r"   rx  debug_printfmtstrargsjax.ArrayLikec                    d}| rGt          t          t          j                                        |                               ^}}}|du}t          j        || |dS )a  Prints scalar values from inside a Pallas kernel.

  Args:
    fmt: A format string to be included in the output. The restrictions on the
      format string depend on the backend:

      * On GPU, when using Triton, ``fmt`` must not contain any placeholders
        (``{...}``), since it is always printed before any of the values.
      * On GPU, when using the experimental Mosaic GPU backend, ``fmt`` must
        contain a placeholder for each value to be printed. Format specs and
        conversions are not supported.
      * In TPU, if ``fmt`` contains placeholders, all values must be 32-bit
        integers. If there are no placeholders, the values are printed after
        the format string.
    *args: The scalar values to print.
  FN)r~  has_placeholders)nextiterstring	Formatterparsedebug_print_pr    )r~  r  r  r.   
field_names        r!   r}  r}    si    "  .T&"2"4"4":":3"?"?@@AAAzA!-		Ts=M	N	N	NNr"   c           
     <   d}t          j                                        |           D ]2\  }}}}||dz  }|s|rt          d          |rt          d          3t	          |          |k    r-t          d| d|dk    rdnd d	t	          |                     d S )
Nr   rY   zDThe format string should not contain any format specs or conversionszDThe format string should not reference arguments by position or namezThe format string expects z	 argument rf   z
, but got )r  r  r  r   r   	TypeError)r~  r  n_placeholdersr.   fieldspec
conversions          r!   check_debug_print_formatr    s    .$*$4$6$6$<$<S$A$A 
 
 ajn z 
P    
P  
 	YY.  
	M^ 	M 	M'1,,22#	M 	MADT	M 	M   ! r"   r   r  c                T    |rt           | j        |            nt          | g|R   dS r1   )printformat)r~  r  r  s      r!   debug_print_implr    s?     	*#*d
	#	r"   avalsc                j    ~ ~t          d |D                       rt          d          g t          hfS )Nc              3  $   K   | ]}|j         V  d S r[   rc   )re   r   s     r!   rg   z,debug_print_abstract_eval.<locals>.<genexpr>  s$      &&&&&&&&r"   z Only scalar values are supported)anyr   debug_print_effect)r~  r  r  s      r!   debug_print_abstract_evalr    sE    	&&&&&&& 9
7
8
88	 !	!!r"   c                \   t          d t          | |          D                       }d }g }t          |          D ]J}t          t	          j        ||          ||           }|                    t          j        |i |           Kd t          | D             }|dt          |          z  fS )z3Unrolls the print primitive across the mapped axis.c              3  :   K   | ]\  }}||j         |         V  d S r[   r  )re   r\   rK  s      r!   rg   z,debug_print_batching_rule.<locals>.<genexpr>  s,      KK$!QQ]171:]]]]KKr"   c                R    |t           j        u r|S t          j        || |d          S )NF)r   keepdims)r   
not_mappedr   index_in_dim)rK  dimargs      r!   get_arg_at_dimz1debug_print_batching_rule.<locals>.get_arg_at_dim  s/    
h!!!jCu====r"   c                6    g | ]}t          j        |          S r2   )r5   stack)re   xss     r!   rj   z-debug_print_batching_rule.<locals>.<listcomp>  s     	-	-	-B#)B--	-	-	-r"   )r   )
r  ziprangemap	functoolspartialappendr  r    r   )r  dimsr   	axis_sizer  outsrK  args_idxs           r!   debug_print_batching_ruler    s    KK#dD//KKKKK)> > > 
$ 9 9a9$^Q77tDDHKK"H7778888	-	-#t*	-	-	-$	tc$ii	r"   c           	         t          j        | t          j        t          j        fi |d t          |          | j        | j        d          \  }}}|S )NT)has_side_effect)	r   emit_python_callbackr  r  r  implr   avals_in	avals_out)ctxr  r   r	  r.   s        r!   debug_print_lowering_ruler    s[    *	*55f55

4jj	l	m  ,&!Q 
-r"   
run_scopedfCallable[..., Any]c                2   t          j        ||f          \  }}t          j        t	          j        |           |          \  }}d |D             }t          j        ||          \  }}	}
\   t          j	        |
d|i}t          j
         |            |          S )zCall the function with allocated references.

  Args:
    f: The function that generatest the jaxpr.
    *types: The types of the function's positional arguments.
    **kw_types: The types of the function's keyword arguments.
  c                6    g | ]}|                                 S r2   )get_aval)re   ts     r!   rj   zrun_scoped.<locals>.<listcomp>4  s     
,
,
,A1::<<
,
,
,r"   jaxpr)r   r   r	   flatten_funlu	wrap_initpetrace_to_jaxpr_dynamicrun_scoped_pr    r   )r  typeskw_types
flat_typesin_treeflat_funout_tree_thunkr  r  r.   constsr   s               r!   r  r  )  s     ".x/@AA*g%1",q//7KK(N
,
,
,
,
,% 28UCC%FB6///#		!.."2"2C	8	88r"   c                R     ~ fd j         D             }d  j        D             |fS )Nc                    h | ];}t          |t          j                  r|j        t	          j                  k    9|<S r2   )ra   r   JaxprInputEffectinput_indexr   	constvars)re   effr  s     r!   	<setcomp>z,_run_scoped_abstract_eval.<locals>.<setcomp>E  sS       

S'2
3
3	
 oU_!5!555	 
 655r"   c                    g | ]	}|j         
S r2   )r   )re   vs     r!   rj   z-_run_scoped_abstract_eval.<locals>.<listcomp>M  s    	(	(	(Q!&	(	(	(r"   )r   r   )r  r  nonlocal_effectss   `  r!   _run_scoped_abstract_evalr  ?  sO    
      
)	(%-	(	(	(*:	::r"   )r   r   r   r   )r   r   )r   r   r   r;   )rW   rD   )r   r   rW   rD   )r   r   )r\   r   r   r   r   r   )F)r   r   )r   r_  )FFNN)rb  r   rc  r   rd  re  )r~  r  r  r  )r  r   r~  r  r  r   )r  r   r~  r  r  r   )r  r  r   r   )__doc__
__future__r   enumr  r  typingr   r   ro  r   r   jax._srcr   r	   r
   r*   r   r   r   r  r   r   r   r   jax._src.interpretersr   r   r   r  jax._src.pallasr%   jax._src.stater   state_discharger   r   r   jax.interpretersr   	jax.numpynumpyr5   r  rb   	NDIndexersafe_mapr  
unsafe_mapsafe_zipr  
unsafe_zipr+   r   r   def_custom_bindr/   def_abstract_evalr8   r=   r:   r@   rB   EnumrD   r   r   register_discharge_ruledef_effectful_abstract_evalr   r   r   r   r   r   r   r   r   r   r   r   r   r   def_implregister_loweringr   r   r   r   r   r  r   r
  pp_eqn_rulesr  primitive_jvpsr  r+  rR  r;  rD  r   rB  rG  rT  rY  r[  ra  rv  Effectrx  r  lowerable_effectsadd_typecontrol_flow_allowed_effectsremat_allowed_effects"custom_derivatives_allowed_effectsr  multiple_resultsr}  r  r  r  r  primitive_batchersr  r  r  r  r2   r"   r!   <module>r     s	   & % " " " " " "                       



                         % % % % % %             & & & & & & ) ) ) ) ) )             $ $ $ $ $ $ * * * * * * 4 4 4 4 4 4 / / / / / / 7 7 7 7 7 7 # # # # # # + + + + + + ! ! ! ! ! !      

	-Z-Z!x!,//
& 
& 
& 
& : : : : - -  - $#N33( ( ( ( 
 
 
  
 !- - "!-    49    "x!,//&7 &7 &7 &7R 6 ' ' 5 56P Q Q Q )? ? ? *)?  @D       @D        ?C        ?C        ?C        ?C        >B        ?C       "x!,//)	V 	V *)	V* * * )(66$ $ 76$
 &8%&677    ** + + +  ')?)? @ @ @1 1 1
 #  $# #"=11   '' ( ( (  }&<&< = = =. . . .
    !  
	M	*	* #  $#  2 !. f   & & & # # # 7<   0 29.d2< 2< 2< . )(00"& "& 10"&J 
	M	*	* #  $#&  8 !. f   $ & &  )(00)7 )7 10)7X &*d     8 +/      ,0 # # # # # # 6;26    &! ! ! ! !'. ! ! ! ![]]    " "; / / /  $ - -k : : :   & &{ 3 3 3  * 3 3K @ @ @ #"=11!% O O O O0   .     *" " " +*"     & .?Y->}. . M *
 4)=99
 
 :9
 "x!,// $ 9 9 9 9, ); ; *); ; ;r"   