
    uki+                       U 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	m
Z
mZ ddlmZ ddlmZ ddlmZ dd	lmZ dd
lmZ ddlmZ ddlmZ ddlmZ ddlmZ ddlmZ ddlmZ ddlmZ ddlmZ 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&m'Z'm(Z(m)Z) ddl*m+Z+m,Z, ddl-Z. ej^                  e0      Z1 ejd                  d      Z3de3_4        ejj                  jm                  e3       ejn                  e8cZ8Z9ejt                  e;cZ;Z< ejz                  d       G d d             Z>	 	 	 	 	 	 dDdZ?e3j                   ej                  ej                  e3             e3j                  	 	 	 	 	 	 dDd        ZDd! ZEeEej                  e3<   d" ZGeGej                  e3<    ej                  ej                  e3      ej                  e3<   dEd#ZK	 	 dFd$ZL	 	 	 	 dGd%ZM ej                  e3eMd&'       d( ZOd e,       dd)	 	 	 	 	 	 	 	 	 	 	 	 	 dHd*ZP ejd                  d+      ZQdeQ_4        ejj                  jm                  eQ        G d, d-ej                        ZS G d. d/ej                        ZT eS       ZU eT       ZVej                  j                  eS       ej                  j                  eT       ej                  j                  eS       ej                  j                  eT       ej                  j                  eT       ej                  j                  eT       	 	 	 	 	 	 dId0Z\eQj                   ej                  ej                  eQ             eQj                  	 	 	 	 	 	 dId1       Z^d2 Z_e_ej                  eQ<   d3 Z`e`ej                  eQ<   d4 Zaeaej                  eQ<   d5 Zb ej                  eQebd&'       dd&d6	 	 	 	 	 	 	 	 	 	 	 dJd7ZcdKd8Zddd9	 	 	 	 	 	 	 	 	 	 	 dLd:Zedd9	 	 	 	 	 	 	 	 	 	 	 dMd;ZfdNd<Zgi Zhd=eid><   dOd?Zjejehej                  <   d@ ehej                  <   dd9	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dPdAZmdd&ddB	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dQdCZny)RzModule for JAX callbacks.    )annotations)CallableSequenceN)Anycast)api)config)core)dispatch)dtypes)effects)ffi)pickle_util)sharding_impls)	tree_util)util)
xla_bridge)ad)batching)mlir)
xla_client)ir)hlo)SdyArraySdyArrayListSdyDimSingleDeviceSharding)ArrayDeprecatedArgpure_callbackT)frozenc                  .    e Zd ZU dZded<   ded<   ddZy)	_FlatCallbacka/  A Python function callable with flat arguments and results.

  An instance of this class is used as a parameter for the callback primitives.
  We prefer it to an anonymous flattened function because it produces
  equal objects when we call the same Python function with the same argument
  structure.
  Callable[..., Any]callback_funcztree_util.PyTreeDefin_treec                    t        j                  | j                  |      \  }}t        j                   | j                  |i |      S N)r   tree_unflattenr&   tree_leavesr%   )self	flat_argsargskwargss       L/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/_src/callback.py__call__z_FlatCallback.__call__D   s@    ++DLL)DLD&  !3!3!3T!DV!DEE    N)r,   r   returnzSequence[Array])__name__
__module____qualname____doc____annotations__r0    r1   r/   r#   r#   8   s     $#Fr1   r#   c                   ~~~ 	 t        j                  d      ^}}t        j                  ||      }t        j                  |      5  	 t        j                  t        j                   ||       cd d d        S # t        $ r}t        d      |d }~ww xY w# t        $ r t        j                  d        w xY w# 1 sw Y   y xY w)Ncpubackendzjax.pure_callback failed to find a local CPU device to place the inputs on. Make sure "cpu" is listed in --jax_platforms or the JAX_PLATFORMS environment variable.zjax.pure_callback failedxblocal_devicesRuntimeErrorr   
device_putr	   default_devicer   tree_mapnpasarrayBaseExceptionlogger	exception)result_avalscallbackshardingvmap_methodr-   
cpu_device_es           r/   pure_callback_implrP   I   s     \%%e4NJ 
j	)$Z( 

HdO<  
 
	/ 	  12 5   A: B:
&B:	BBB B77B::Cc                    ~~ ~~|S r(   r8   )rJ   rI   rK   rL   avalss        r/   pure_callback_abstract_evalrT   f   s     Xx	r1   c                     ~ ~t        d      )NzgPure callbacks do not support JVP. Please use `jax.custom_jvp` to use callbacks while taking gradients.
ValueErrorr-   r.   s     r/   pure_callback_jvp_rulerY   r       
FM	N Nr1   c                     ~ ~t        d      )NzmPure callbacks do not support transpose. Please use `jax.custom_vjp` to use callbacks while taking gradients.rV   rX   s     r/   pure_callback_transpose_ruler\   |   rZ   r1   c                    dg}| r4| D cg c])  }t        |t        j                        s|j                  + }}t	        |D cg c]  }t        dt        g d      g|z  d      ! c}      S c c}w c c}w )zGReturns an SdyArrayList with `max(1, len(avals))` replicated shardings.r   r8   F)axesis_open
mesh_shapedim_shardingslogical_device_ids)
isinstancer
   ShapedArrayndimr   r   r   )rS   ndimsxrf   s       r/   !_get_sdy_array_list_for_callbacksri      s    #%
"FjD4D4D&EQVVFEF	 .3	4 &* R784?!4 
5 5 G4s   A.A.$A3c                ,   t        | t        j                        r| j                  t	        | j
                  j                        k7  rt        d      |t        d      t        j                  j                  rt        |      }|S t        j                         }t        j                  j                  j                  |_        |S t        | t        j"                        r|qt        |t$              st        dt!        |             t'        t)        |j*                              }| j,                  }|t/        d      	 |j1                  |      }nd}t        j                  j                  r3t7        d	t9        |            }t;        |t=        d
g |f      gz        }|S t        j                         }t        j                  j                  j>                  |_        d	g|_         |g|_!        |S y # t2        $ r}t5        d| d| d      |d }~ww xY w)Nz~callbacks are only supported in spmd computations when all mesh axes are partitioned manually (no partial automatic sharding).zEcallbacks do not support specifying sharding inside spmd computationsz:pure_callback only supports SingleDeviceSharding, but got z9Please file a bug at https://github.com/jax-ml/jax/issuesz6Sharding provided to pure_callback specifies a device z' that is not in the device assignment ()r      r8   r`   )"rd   r   SPMDAxisContextmanual_axes	frozensetmesh
axis_namesNotImplementedErrorr	   use_shardy_partitionervalueri   xc
OpShardingTypeMANUALtypeShardingContextr   nextiter
device_setdevice_assignmentAssertionErrorindex
IndexErrorrW   maxlenr   r   MAXIMALtile_assignment_dimensionstile_assignment_devices)	axis_contextrK   	avals_outop_shardingdevicer~   device_indexrO   num_sdy_shardingss	            r/   _callback_op_shardingr      s#    n<<= 9\->->-I-I#JJL    $$**5i@k  MMOk++22kn<<="67!X!
 	
 D,,-.f&88		"GI 	I.(..v6 l
 $$** aY0 !2
".263 "3 4k 	 MMOk++33k01sk,-9Nk) 
;  .x "#1&' -.	..s   G2 2	H;HHc                   fd}t        | j                  j                  || j                        }t	        | |d t        |      | j                  | j                  dd|	      \  }}}|S )Nc                 2    t        t        | d d      S )N)rJ   rK   )tuplerP   )r,   rJ   paramss    r/   	_callbackz)pure_callback_lowering.<locals>._callback   s+    	
 		
 r1   Fhas_side_effectreturns_tokenrK   )r   module_contextr   r   emit_python_callbacklistavals_in)	ctxrJ   rK   r-   r   r   r   resultrN   s	    `  `    r/   pure_callback_loweringr      sj     &	%%x@+%	

4j	ll	mm
,&!Q 
-r1   F)	cacheablec                    t        j                  | j                        }t        j                  |      |k7  rt	        d      y )NzQresult_shape_dtypes cannot specify 64-bit types when `jax_enable_x64` is disabled)rD   dtyper   canonicalize_dtyperW   )shape_dtypedts     r/   _check_shape_dtyper      s@    	xx!!""r"b(
[] ] )r1   )rK   
vectorizedrL   c                  t        |t              st        d      g d}||vrt        d| d|       t        j                  ||f      \  }}	t        j
                  t        |       t        j
                  d |      }
t        j                  |
      \  }}t        j                  |t        | |	      t        |      ||d}t        j                  ||      S )a  Calls a pure Python callback. Works under :func:`jit`/:func:`~vmap`/etc.

  For more explanation, see `External Callbacks`_.

  ``pure_callback`` enables calling a Python function in JIT-ed JAX functions.
  The input ``callback`` will be passed JAX arrays placed on a local CPU, and
  it should also return JAX arrays on CPU.

  The callback is treated as functionally pure, meaning it has no side-effects
  and its output value depends only on its argument values. As a consequence, it
  is safe to be called multiple times (e.g. when transformed by :func:`~vmap` or
  :func:`~pmap`), or not to be called at all when e.g. the output of a
  `jit`-decorated function has no data dependence on its value. Pure callbacks
  may also be reordered if data-dependence allows.

  .. warning::

     In the context of JAX transformations, Python exceptions should be
     considered side-effects: this means that intentionally raising an error
     within a `pure_callback` breaks the API contract, and the behavior of
     the resulting program is undefined.

  When `vmap`-ed the behavior will depend on the value of the ``vmap_method``.

  * Calling :func:`~jax.vmap` on a callback without an explicit ``vmap_method``
    raises a ``NotImplementedError``.
  * ``vmap_method="sequential"`` uses :func:`~jax.lax.map` to loop over
    the batched arguments, calling ``callback`` once for each batch element.
  * ``vmap_method="sequential_unrolled"`` is like ``sequential``, but the loop
    is unrolled.
  * ``vmap_method="expand_dims"`` calls ``callback`` with new axes of size ``1``
    added as the leading dimension unbatched inputs.
  * ``vmap_method="broadcast_all"`` behaves like ``expand_dims``, but the
    inputs are tiled to the expected batched shape.

  If necessary, the legacy behavior provided by the removed ``vectorized=True``
  argument can be recovered using ``vmap_method="legacy_vectorized"``.

  The current default behavior is to use ``vmap_method="sequential"`` when
  not specified, but this behavior is deprecated, and in the future, the
  default will be to raise a ``NotImplementedError`` unless ``vmap_method`` is
  explicitly specified.

  Args:
    callback: function to execute on the host. The callback is assumed to be a pure
      function (i.e. one without side-effects): if an impure function is passed, it
      may behave in unexpected ways, particularly under transformation. The callable
      will be passed PyTrees of arrays as arguments, and should return a PyTree of
      arrays that matches ``result_shape_dtypes``.
    result_shape_dtypes: pytree whose leaves have ``shape`` and ``dtype`` attributes,
      whose structure matches the expected output of the callback function at runtime.
      :class:`jax.ShapeDtypeStruct` is often used to define leaf values.
    *args: arguments to be passed to the callback function
    sharding: optional sharding that specifies the device from which the callback should
      be invoked.
    vmap_method: string specifying how the callback transforms under
      :func:`~jax.vmap` as described above.
    **kwargs: keyword arguments to be passed to the callback function

  Returns:
    result: a pytree of :class:`jax.Array` objects whose structure matches that of
      ``result_shape_dtypes``.

  See Also:
    - :func:`jax.experimental.io_callback`: callback designed for impure functions.
    - :func:`jax.debug.callback`: callback designed for general-purpose debugging.
    - :func:`jax.debug.print`: callback designed for printing.

  Examples:
    The behavior of ``pure_callback`` under :func:`~jax.vmap` is controlled by
    the ``vmap_method`` argument as described above. It is useful to consider
    some explicit examples that demonstrate the semantics. For example,
    consider the following function:

    >>> def callback(x, y):
    ...   print(jnp.shape(x), jnp.shape(y))
    ...   return x + y

    >>> def fun(x, y, *, vmap_method):
    ...   shape = jnp.broadcast_shapes(jnp.shape(x), jnp.shape(y))
    ...   dtype = jnp.result_type(x, y)
    ...   out_type = jax.ShapeDtypeStruct(shape, dtype)
    ...   return jax.pure_callback(callback, out_type, x, y,
    ...                            vmap_method=vmap_method)

    Calling this with ``vmap_method="expand_dims"`` adds a new axis of size ``1``
    to ``y``:

    >>> from functools import partial
    >>> x = jnp.arange(4)
    >>> y = 1.0
    >>> jax.vmap(partial(fun, vmap_method="expand_dims"), in_axes=(0, None))(x, y)
    (4,) (1,)
    Array([1., 2., 3., 4.], dtype=float32)

    Whereas, ``vmap_method="broadcast_all"`` adds an axis of size ``4`` to
    ``y``:

    >>> jax.vmap(partial(fun, vmap_method="broadcast_all"),
    ...          in_axes=(0, None))(x, y)
    (4,) (4,)
    Array([1., 2., 3., 4.], dtype=float32)

  .. _External Callbacks: https://docs.jax.dev/en/latest/external-callbacks.html
  zdThe 'vectorized' argument of jax.pure_callback was removed in JAX v0.6.0. Use 'vmap_method' instead.)
sequentialsequential_unrolledexpand_dimsbroadcast_alllegacy_vectorizedNz.vmap_method must be on of the allowed methods z, but got: c                V    t        j                  | j                  | j                        S r(   r
   re   shaper   rh   s    r/   <lambda>zpure_callback.<locals>.<lambda>  s      !''2 r1   )rJ   rI   rK   rL   )rd   r   rW   r   tree_flattenrC   r   pure_callback_pbindr#   r   r)   )rJ   result_shape_dtypesrK   r   rL   r-   r.   allowed_vmap_methodsr,   r&   rI   flat_result_avalsout_treeout_flats                 r/   r    r      s    f 
J	.
	-. .F,,

89M8N O=	"# # !--tVn=)W')<=##24GI, ) 6 6| DX!!Xw/*+( 
	!	!(H	55r1   io_callbackc                      e Zd Zd Zy)IOEffectc                     y)NIOr8   rN   s    r/   r   zIOEffect.<lambda>      r1   Nr3   r4   r5   __str__r8   r1   r/   r   r     s    'r1   r   c                      e Zd Zd Zy)OrderedIOEffectc                     y)N	OrderedIOr8   r   s    r/   r   zOrderedIOEffect.<lambda>  r   r1   Nr   r8   r1   r/   r   r     s    !'r1   r   c                   ~ ~~	 t        j                  d      ^}}t        j                  ||      }t        j                  |      5  	 t        j                  t        j                   ||       cd d d        S # t        $ r}t        d      |d }~ww xY w# t        $ r t        j                  d        w xY w# 1 sw Y   y xY w)Nr:   r;   zjax.io_callback failed to find a local CPU device to place the inputs on. Make sure "cpu" is listed in --jax_platforms or the JAX_PLATFORMS environment variable.zjax.io_callback failedr=   )rI   rJ   rK   orderedr-   rM   rN   rO   s           r/   io_callback_implr     s     Hg%%e4NJ 
j	)$Z( 

HdO<  
 
	/ 	  /0 rQ   c                .    ~~~ |rt         nt        }||hfS r(   )_OrderedIOEffect	_IOEffect)rJ   rI   rK   r   rS   effects         r/   io_callback_abstract_evalr     s"     Xx&I&	x	r1   c                     ~ ~t        d      )Nz IO callbacks do not support JVP.rV   rX   s     r/   io_callback_jvp_ruler     s    
F566r1   c                     ~ ~t        d      )Nz&IO callbacks do not support transpose.rV   rX   s     r/   io_callback_transpose_ruler     s    
F;<<r1   c           	        ddl m} |rt        d      |D cg c]  }|t        j                  u c}t        | |      D 	cg c]0  \  }}	|	t        j                  u r|nt        j                  ||	d      2 }
}}	t        j                  |
      \  }fd} |||      }|dt        |      z  fS c c}w c c}	}w )Nr   )mapz"Cannot `vmap` ordered IO callback.c                b    t        j                  |       }t        j                  |ddS )NF)rJ   rK   rI   r   )r   merge_listsio_callback_pr   )batched_argsmergedrJ   
is_batchedrI   rK   unbatched_argss     r/   
_batch_funz-io_callback_batching_rule.<locals>._batch_fun  s7    j.,GFv8+7H Hr1   )r   )
jax._src.lax.control_flow.loopsr   rW   r   
not_mappedzipmoveaxisr   partition_listr   )r-   dimsrJ   rI   rK   r   lax_mapdargdimnew_argsr   r   out_valsr   r   s     ```         @@r/   io_callback_batching_ruler     s     =
9
::6:;,,,;*=@t_N19c H///cS!,- N( N!%!4!4Z!J.,H H Z.(	4#h-'	'' <Ns   B;5C c                  fd}t        | j                  j                  || j                        }r| j                  j                  t              }t        | ||t        |      | j                  | j                  dd|	      \  }	}}
| j                  t        j                  t        |i             |	S t        | |d t        |      | j                  | j                  dd|	      \  }	}
}
|	S )Nc            	     4    t        t        | d d      S )N)rJ   rK   r   )r   r   )r,   rJ   r   r   s    r/   r   z'io_callback_lowering.<locals>._callback  s.    		

 	
 r1   Tr   F)r   r   r   r   	tokens_ingetr   r   r   r   set_tokens_outr   TokenSet)r   rJ   rK   r   r-   r   r   r   tokenr   rN   s    ` ` `     r/   io_callback_loweringr     s    	 &	%%x@+MM./E+T

FE1 t}}&6%>?@ 
- (T

LFAq 
-r1   )rK   r   c               6   t        j                  ||f      \  }}t        j                  t        |       t        j                  |      \  }}	t	        d |      }
t        j                  |t        | |      t        |
      ||d}t        j                  |	|      S )a"  Calls an impure Python callback.

  For more explanation, see `External Callbacks`_.

  Args:
    callback: function to execute on the host. It is assumed to be an impure function.
      If ``callback`` is pure, using :func:`jax.pure_callback` instead may lead to
      more efficient execution.
    result_shape_dtypes: pytree whose leaves have ``shape`` and ``dtype`` attributes,
      whose structure matches the expected output of the callback function at runtime.
      :class:`jax.ShapeDtypeStruct` is often used to define leaf values.
    *args: arguments to be passed to the callback function
    sharding: optional sharding that specifies the device from which the callback should
      be invoked.
    ordered: boolean specifying whether sequential calls to callback must be ordered.
    **kwargs: keyword arguments to be passed to the callback function

  Returns:
    result: a pytree of :class:`jax.Array` objects whose structure matches that of
      ``result_shape_dtypes``.

  See Also:
    - :func:`jax.pure_callback`: callback designed for pure functions.
    - :func:`jax.debug.callback`: callback designed for general-purpose debugging.
    - :func:`jax.debug.print`: callback designed for printing.

  .. _External Callbacks: https://docs.jax.dev/en/latest/notebooks/external_callbacks.html
  c                V    t        j                  | j                  | j                        S r(   r   r   s    r/   r   zio_callback.<locals>.<lambda>C  s    D$4$4QWWagg$F r1   )rJ   rI   rK   r   )
r   r   rC   r   r   r   r   r#   r   r)   )rJ   r   rK   r   r-   r.   r,   r&   flat_shape_dtypesr   r   r   s               r/   r   r     s    H !--tVn=)W')<= ) 6 67J KXF+-Xw/*+( 
	!	!(H	55r1   c                &    t        d | D              S )Nc              3  &   K   | ]	  }|d k(    yw)r   Nr8   ).0r   s     r/   	<genexpr>z!is_empty_shape.<locals>.<genexpr>P  s     Q!Vs   )any)ss    r/   is_empty_shaper   O  s    	Q	r1   rK   c          
        t         j                  j                  | t        j                        }t        j
                  |g||t        j                  j                  d            }t        j                  j                  t        t        j                  j                  t        |            t        j                  j                  t        |                        |j                  d<   |t        j                  j                  r[t!        |t"              sJ t%        |j&                        dk\  sJ t#        t)        dg |j&                  d   j*                        g      }t        j,                  ||       |j.                  S )	NTis_host_transfer_xla_host_transfer_handler_name_xla_host_transfer_rendezvousmhlo.frontend_attributesrl   r8   r   r`   )r   ChannelHandler   r   SEND_TO_HOST_TYPESendOpr   BoolAttrDictAttrdict
StringAttrstr
attributesr	   rs   rt   rd   r   r   	shardingsr   rc   set_shardingr   )channelr   operandnamerK   channel_handlesend_ops          r/   send_to_hostr  S  s.    $$(($2H2HI.JJy%)+)>@'35;;??
*,--*;*;CI*F(*(9(9#d)(DF4G'/0 $$** ,///##$)))
2!)!3!3A!6!I!IKL Mh 	gx(	r1   c          
        t         j                  j                  | t        j                        }t        j
                  t        j                  |      t         j                  j                         g||t        j                  j                  d            }t        j                  j                  t        t        j                  j                  t        |            t        j                  j                  t        |                        |j                  d<   |t        j                   j"                  rit%        |t&              sJ t)        |j*                        dk\  sJ t'        |j*                  d   t-        dg |j*                  d   j.                        g      }t        j0                  ||       |j2                  \  }}||fS )	NTr   r   r  rl   r   r8   r`   )r   r  r   r   RECV_FROM_HOST_TYPERecvOpaval_to_ir_type	TokenTyper   r  r  r  r  r	  r
  r	   rs   rt   rd   r   r   r  r   rc   r  results)r  r   out_avalr  rK   r  recv_opr   s           r/   receive_from_hostr  s  sj    $$(($2J2JK.JJ,,X6))+-.3^)+)>@' 46;;??
*,--*;*;CI*F(*(9(9#d)(DF4G'/0 $$**,///##$))) 


Q

2!)!3!3A!6!I!IKL Mh
 	gx(//-&%	r1   c                    	 t        t        |          |       S # t        $ r}t        dt        |              |d }~ww xY w)NzNo xla_shape_handler for type: )_xla_shape_handlersry   KeyError	TypeError)avalerrs     r/   _aval_to_xla_shaper"    sG    MtDz*400	 M
5d4j\B
CLMs    	?:?z9dict[type[core.AbstractValue], Callable[[Any], xc.Shape]]r  c                    t        j                  |       } | j                  t        j                  k(  rt        j                  d      n| j                  }t        j                  j                  || j                        S )Nbool)
r
   physical_avalr   r   float0rD   ru   Shapearray_shaper   )r   r   s     r/   _make_array_shaper)    sP    			D	!$"jjFMM9"((6
tzz%			eTZZ	00r1   c                >    t         j                  j                         S r(   )ru   r'  token_shaper   s    r/   r   r     s    BHH4H4H4J r1   c	                  |xs t        j                         }|}g }|s|fd}|j                  j                         }t	        j
                  dt        j                        }t        j                  t        j                  dt        j                              }g |t        |      }t        ||||j                  |
      }|j                  |       nL|D ]G  }|j                  j                         }t        ||||j                  |
      }|j                  |       I g }g }|	r|s|fd}t	        j
                  dt        j                        }t        |      g}|j                  j                         }t        ||||j                  |
      \  }}|j                  |       n||D ]w  }|j                  j                         }t!        |t        j
                        sJ t        ||||j                  |
      \  }}|j                  |       |j                  |       y | j#                  |||||t$        j&                        }|j                  j)                  |       ||fS )Nc                     ~         S r(   r8   )r-   callback_without_argss    r/   _wrapped_callbackz4_emit_tpu_python_callback.<locals>._wrapped_callback  s    
"$$r1   )rl   rl   r   c                      |   y)N)g        r8   )r-   callback_without_resultss    r/   r/  z4_emit_tpu_python_callback.<locals>._wrapped_callback  s    %r1   r8   )r   create_tokenr   new_channelr
   re   rD   float32r   ir_constantzerosr"  r  r3   appendr  rd   ,make_python_callback_from_host_send_and_recvr   dumpsadd_host_callback)r<   r   rJ   r   operandsoperand_avalsoperand_shapesrI   result_shapesr   rK   r/  send_channelssend_channeldummy_send_avaldummy_send_valr  r  recv_channelsoutputsdummy_recv_avalrN   result_avaloutifrt_callbackr.  r1  s                            @@r/   _emit_tpu_python_callbackrI    sF    
%3##%%-	
 .% %%113L&&tRZZ8O%%bhhq"**&=>NK~K'9/'JKNunh>O>O"*,E& $""..0g7E7H4E4E$,.e7#	$ -'<  1 &&r2::6O'89M  ,,.G ):):XOHE1!# $""..0gT%5%5666$We[%-%6%6KjeSnnS7#$ FF[&&(- &&}5	%r1   )r   partitionedrK   c                  t        | j                  j                        dkD  rt        d      | j                  j                  d   dvrt	        d d      |r dvrt        d d      rt	        d	      t        t        j                  | j                  j                               }
D cg c]  }t        |       }}|D cg c]  }t        |       }}fd
}dk(  rt        j                  t        |      D cg c]  \  }}t        |j                        s||f  c}}      \  }}t        |
| |||||||||	      \  }}t!        |      }D cg c]`  }t        |j                        r>t#        j$                  t'        j(                  |j                  |j*                              n
t-        |      b }}||dfS dv rdnd}|rdnd}d| d| d}|r|fd}|g|}t.        j0                  j2                  rN|	Lt        | j4                        dkD  r4t7        |	t8              r$t9        t;        dg d      g|	j<                        }	t?        j@                  | tB        jD                  g| jF                  tB        jD                  g| j4                        } |}| j                  jI                  |       t'        jJ                  t        | j                  jL                        dz
        } tO        jP                  ||      | g|dt'        jJ                  |      i}|	t#        jR                  ||	       |jT                  }|r|^}}|||fS c c}w c c}w c c}}w c c}w )a  Emits MLIR that calls back to a provided Python function.

  Args:
    ctx: The lowering context.
    callback: The Python callback function.
    token: The token to use for the callback.
    operands: The operands to the callback.
    operand_avals: The abstract values of the operands.
    result_avals: The abstract values of the results.
    has_side_effect: Whether the callback has side effects.
    returns_token: Whether the callback should return a token.
    partitioned: If True, then `callback` is called on local shards only. If
      False, then `callback` is called on all shards.
    sharding: The sharding of the callback.

  Returns:
    A tuple of MLIR result values, a new token (if any), and the host callback
    object.
  rl   z+multi-platform lowering for python_callbackr   >   r:   tpucudarocmz&`EmitPythonCallback` not supported on z	 backend.>   r:   rM  rN  z(Partitioned callback not implemented on z6Partitioned callback not supported with return values.c            	     H    |  }t        |      t              k7  r-t        dj                  t              t        |                  t        d |D              }t	        t        |            D ]  \  }\  }}|j                  |j                  k7  r(t        d| d|j                   d|j                         |j                  |j                  k7  sdt        d| d|j                   d|j                          dk(  rt        d t        |      D              }|S |S )	NzDMismatched number of outputs from callback. Expected: {}, Actual: {}c              3  l   K   | ],  }t        j                  t        j                  |             . y wr(   )r   canonicalize_valuerD   rE   )r   as     r/   r   zBemit_python_callback.<locals>._wrapped_callback.<locals>.<genexpr>+  s#     P!V..rzz!}=Ps   24z)Incorrect output shape for return value #z: Expected: z
, Actual: z)Incorrect output dtype for return value #rL  c              3  P   K   | ]  \  }}t        |j                        s|   y wr(   )r   r   )r   out_valrF  s      r/   r   zBemit_python_callback.<locals>._wrapped_callback.<locals>.<genexpr>;  s*      !4"g{ 1 12 !4s   $&)r   r@   formatr   	enumerater   r   r   )	r-   r   irT  r  non_empty_out_valsrJ   platformrI   s	         r/   r/  z/emit_python_callback.<locals>._wrapped_callback$  sC   H
8}L))%%+VC,=s8}%MO O PxPPH"+C,,G"H DGX	(..	(7s ;!(
7==/CD 	D 
(..	(7s ;!(
7==/CD 	DD 5
 ! !4&)(L&A!4 4  or1   rL  )r   rK   )r   N>   rM  rN  gpur:   _partitioned xla_ffi_python_r   c                    | g | S r(   r8   )r   r-   callback_without_tokens     r/   r/  z/emit_python_callback.<locals>._wrapped_callbackY  s    4,d344r1   r8   r`   )r   r   )r   r   )+r   r   	platformsrr   rW   r   ru   Clientget_backendr"  r   unzip2r   r   r   rI  r|   r   r5  rD   r6  r   r{   r	   rs   rt   r   rd   r   r   r  dataclassesreplacer
   abstract_tokenr   r:  uint64host_callbacksr   build_ffi_lowering_functionr  r  )r   rJ   r   r;  r<  rI   r   r   rJ  rK   r<   r   r>  r=  r/  r   non_empty_result_avalsnon_empty_result_shapesnon_empty_outputsnon_empty_outputs_iterrF  rD  r   	partitioncall_target_namerH  r   r   r  r`  rY  s    `   `                       @@r/   r   r     s   @ 				%	%&*
K
LL))!,(33

0
)DF F..4XJi
HJ JOPPBIIs'9'9'E'E'GH'8DE%d+E-E9FG&t,G.G> 6:kk|];C+D%djj) 
uC+ 7,33  9'%- 7#h	 8u
 ""34 ()  +++, 	+"3"3;;L;LMN267M2N	O)G ) E4 005e& +n)yk&C
.5!!H%%++ "x.
 
!#%
   h 

%%55&&77C $-&&}5
))C**99:Q>
?%3**% , , ))E*,&
 fh'NN'
OEG	%	&&I FGHC+)s   6MM##M(
"A%M.)rJ   r#   rK   SingleDeviceSharding | NonerL   
str | None)rS   Sequence[core.ShapedArray]r2   r   )rK   rq  )rJ   r#   rK   rq  )rJ   r$   r   r   r-   r   rK   rq  r   zbool | None | DeprecatedArgrL   rr  r.   r   )rJ   r#   rK   rq  r   r$  )rJ   r$   r   r   r-   r   rK   rq  r   r$  r.   r   )r   z
core.Shaper2   r$  )r  intr   hlo.TokenTyper  r   r  r	  rK   #SdyArrayList | xc.OpSharding | Noner2   zir.Value)r  rt  r   ru  r  core.ShapedArrayr  r	  rK   rv  r2   ztuple[ir.Value, ir.Value])r   zcore.AbstractValuer2   xc.Shape)r   rw  r2   rx  )r<   z	xc.Clientr   mlir.LoweringRuleContextr   
Any | Noner;  Sequence[ir.Value]r<  rs  r=  Sequence[xc.Shape]rI   rs  r>  r|  r   r$  rK   rv  r2   ztuple[Sequence[ir.Value], Any])r   ry  r   rz  r;  r{  r<  rs  rI   rs  r   r$  r   r$  rJ  r$  rK   rv  r2   z(tuple[Sequence[mlir.IrValues], Any, Any])or6   
__future__r   collections.abcr   r   re  	functoolsloggingtypingr   r   jax._srcr   r	   r
   r   r   r   r   r   r   r   r   r   r>   jax._src.interpretersr   r   r   jax._src.libr   ru   jax._src.lib.mlirr   jax._src.lib.mlir.dialectsr   jax._src.sharding_implsr   r   r   r   jax._src.typingr   r   numpyrD   	getLoggerr3   rG   	Primitiver   multiple_results%prim_requires_devices_during_loweringaddsafe_mapr   
unsafe_mapsafe_zipr   
unsafe_zip	dataclassr#   rP   def_implpartialapply_primitivedef_abstract_evalrT   rY   primitive_jvpsr\   primitive_transposesffi_batching_ruleprimitive_batchersri   r   r   register_loweringr   r    r   Effectr   r   r   r   lowerable_effectsadd_typecontrol_flow_allowed_effectsordered_effectsshardable_ordered_effectsr   def_effectful_abstract_evalr   r   r   r   r   r   r   r  r  r"  r  r7   r)  re   AbstractTokenrI  r   r8   r1   r/   <module>r     s;     " .              #   % $ * & )   * X X 0 			8	$ !$..1#'    . . 2 2? C--Z--Z d#F F $F   *	
 2   ***8+C+C+:< = "" *	
  #N &<  / "N ,H   ( 0Ay/@/@?0  O ,	5B7BJ'3N>   (>% P] -1.;o"J6 J6J6 J6 *	J6
 ,J6 J6 J6^ }-!%   . . 2 2= Aw~~ "gnn " J	"$     " "8 ,    " "? 3  $ $ - -h 7  $ $ - -o >        1  ! ! * *? ;  *	
 2   (y(()A)A)68 9 **	 	  *		 
 	  +	 7 $8  -  = *D   &(  .G  M *(X   }&:e L -106 0606 06 *	06
 06 06f  59  	 2 L 59""" " 	" 2" "JM 79  4 91 ): D$$ %*J D&& ' 59EE	!E 	E
 !E .E 'E -E &E E 2E $Eb 48Q'	!Q' Q' !	Q'
 .Q' -Q' Q' Q' Q' 2Q' .Q'r1   