ukiKs ddlmZddlmZmZmZddlZddlZddlZddl Z ddl m Z m Z m Z mZddlZddlmZddlmZddlmZddlmZdd lmZdd lmZdd lmZdd lmZdd lmZddlmZddlm Z ddl!m"Z"ddl!m#Z#ddl$m%Z%ddl&m'Z'm(Z(m)Z)m*Z*m+Z+ejXe-cZ-Z.ee/e zdzZ0 d7 d8dZ1Gdde Z2 d9 d:dZ3 d9 d;dZ4ddZ8 d? d@dZ9ddddd dAd Z:ddddd dBd!Z;e*ejxzZ= dCd"Z>dDd#Z? dEd$Z@ed%d%d%d%d%d%d%e)d& dFd'ZAed%d%d%d%d%d%d%e)d& dGd(ZAdddddd)de)d& dHd*ZAdId+ZBdJd,ZCejd-.Gd/d0ejZFeFZGejjeFejjeF dKd1ZKd2ZLd3ZM dLd4ZN dMd5ZOejd6ZQd-eQ_RejeQeQjeKeLejeQ<eMejeQ<ejeOeQejeQ<ejeQeNy)N) annotations)CallableMappingSequenceN)Any TypedDict NotRequiredoverload)core)dispatch)effects)util) xla_bridge) HashableArray) FrozenDict)ad)batching)mlir)Layout)jaxlib) xla_client)ir)Array ArrayLike DeprecatedArgDuckTypedArrayShapec 4tj||||fi|S)aRegisters a foreign function target. Args: name: the name of the target. fn: a ``PyCapsule`` object containing the function pointer, or a ``dict`` where the keys are FFI stage names (e.g. `"execute"`) and the values are ``PyCapsule`` objects containing a pointer to the handler for that stage. platform: the target platform. api_version: the XLA custom call API version to use. Supported versions are: 1 (default) for the typed FFI or 0 for the earlier "custom call" API. kwargs: any extra keyword arguments are passed directly to :func:`~jaxlib.xla_client.register_custom_call_target` for more advanced use cases. )rregister_custom_call_target)namefnplatform api_versionkwargss G/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/_src/ffi.pyregister_ffi_targetr&/s'*  / /b(K :28 ::c&eZdZUdZded<ded<y)TypeRegistrationaA dictionary type for registering FFI types. Attributes: type_id: A ``PyCapsule`` object containing a pointer to the ``XLA_FFI_TypeId``. type_info: An optional ``PyCapsule`` object containing a pointer to the type ``XLA_FFI_TypeInfo``. rtype_idzNotRequired[Any] type_infoN)__name__ __module__ __qualname____doc____annotations__r'r%r)r)Hs , r'r)ctd)zRegisters a custom type ID for a FFI target. Args: name: the name of the type ID. This name must be unique within the process. obj: a ``PyCapsule`` object encapsulating a pointer to the type ID. platform: the target platform. z?register_ffi_type_id is not supported after jaxlib version 381. ValueError)r objr"s r%register_ffi_type_idr6Vs G IIr'c2tj|||S)zRegisters a custom type for a FFI target. Args: name: the name of the type. This name must be unique within the process. type_registration: a ``TypeRegistration`` defining the external type. platform: the target platform. )r")rregister_custom_type)r type_registrationr"s r%register_ffi_typer:es  ( (  r'ctj|tjt j tj|y)z]Registers an FFI target as batch partitionable. Args: name: the name of the target. N)r+register_custom_call_as_batch_partitionablerregister_plugin_callbacks functoolspartial)r s r%*register_ffi_target_as_batch_partitionabler@vs9  88> && NNr'ctjdtj}tjj}tj|_tj tj|f|_||d|dS)aWrap a ctypes function pointer in a PyCapsule. The primary use of this function, and the reason why it lives with in the ``jax.ffi`` submodule, is to wrap function calls from external compiled libraries to be registered as XLA custom calls. Example usage:: import ctypes import jax from jax.lib import xla_client libfoo = ctypes.cdll.LoadLibrary('./foo.so') xla_client.register_custom_call_target( name="bar", fn=jax.ffi.pycapsule(libfoo.bar), platform=PLATFORM, api_version=API_VERSION ) Args: funcptr: A function pointer loaded from a dynamic library using ``ctypes``. Returns: An opaque ``PyCapsule`` object wrapping ``funcptr``. Nr) ctypes CFUNCTYPE py_object pythonapi PyCapsule_Newrestypec_void_pc_char_pargtypes)funcptr destructorbuilders r% pycapsulerNsd6f&6&67*    * *'$$'/oov C' $ 1 ..r'ctjjtjjtj }tjj |dS)zIGet the path to the directory containing header files bundled with jaxlibinclude)ospathdirnameabspathr__file__join) jaxlib_dirs r% include_dirrXs:wwrwwv?@* j) ,,r'ch|tjurdStj|jSNr1)r abstract_token physical_avalshape)avals r% _aval_shaper_s+t***N0B0B40H0N0NNr'c |/ttttt |St |t r)|j td|jdddSt|S)zIConvert a layout to the minor-to-major order used by the custom call API.Nz,The FFI does not support layouts with tiling) tuplereversedrangelenr_ isinstancertilingr4major_to_minor)r^layouts r%_convert_layout_for_loweringrjsg ^ %K$5 678 99&&! }} E FF  2 && =r'Foperand_layoutsresult_layoutsbackend_configskip_ffi_layout_processingc 8 dfd }|S)a<Build a lowering op for an foreign function interface (FFI) target. By default, this lowering rule can use the input and output abstract values to compute the input and output types and shapes for the custom call, assuming row-major layouts. Note that layouts passed to this function as tuples should be in minor-to-major order (as expected by XLA) rather than major-to-minor as used by :func:`~jax.ffi.ffi_call` and ``Layout``. If keyword arguments are passed to the lowering rule, these are treated as attributes, and added to `backend_config`. Args: call_target_name: The name of the custom call target. operand_layouts: A sequence of layouts (dimension orders) for each operand. By default, the operands are assumed to be row-major. result_layouts: A sequence of layouts (dimension orders) for each result. By default, the results are assumed to be row-major. backend_config: Configuration data for the custom call. Any keyword arguments passed to the lowering rule will added to this dictionary. lowering_args: Any other arguments to :func:`mlir.custom_call` will also be passed through if provided as extra arguments to this function. skip_ffi_layout_processing: If true, skip processing of operand and result layout arguments passed to the lowering rule. c *t }|jdd|ddk\rftts tdtxsifi|j Dcic]\}}|t j |c}}|d<n|rtd|dd|d<d|vr0|jDcgc]}t j|c}|d< s tt|j|d<n-t|j Dcgc] }t| c}|d< tt|j|d <n-t|j Dcgc] }t| c}|d <d |vritd |jDsM|jDcgc]4}t jt j|t!|6c}|d <t j" fd |i|Scc}}wcc}wcc}wcc}wcc}w) Nr#z:When api_version > 4, backend_config must be a dictionary.rnzaThe use of ffi_call attributes requires a custom call API version of at least 4; got api_version=. result_typesrlrm result_shapesc3XK|]"}tjt|$ywN)r is_constant_shaper_).0r^s r% zDbuild_ffi_lowering_function.._lowering_op..s'1M6:{4011Ms(*operands)dict setdefaultrfr4itemsr ir_attribute avals_outaval_to_ir_typemaprjavals_inzipall shape_tensoreval_dynamic_shape_as_ivalsr_ custom_call)ctxr{paramsr$kvr^argsrncall_target_name lowering_argsrlrmros r% _lowering_opz1build_ffi_lowering_function.._lowering_ops;- F mQ' m!  #J~t,L HJ J!%""VEK\\^ TTQD$5$5a$8!8 T"Vf   ..4].C-DA GH H"0f V#GJ}}Ut 4 4T :Uf^ %  $' (#,,%  ! CLL/:%  )$ /%  !  #& (#--$  CMM>:$  )$ /$  f$S1M>Amm1M.Mmm!%   D<Build a lowering rule for an foreign function interface (FFI) target. By default, this lowering rule can use the input and output abstract values to compute the input and output types and shapes for the custom call, assuming row-major layouts. Note that layouts passed to this function as tuples should be in minor-to-major order (as expected by XLA) rather than major-to-minor as used by :func:`~jax.ffi.ffi_call` and ``Layout``. If keyword arguments are passed to the lowering rule, these are treated as attributes, and added to `backend_config`. Args: call_target_name: The name of the custom call target. operand_layouts: A sequence of layouts (dimension orders) for each operand. By default, the operands are assumed to be row-major. result_layouts: A sequence of layouts (dimension orders) for each result. By default, the results are assumed to be row-major. backend_config: Configuration data for the custom call. Any keyword arguments passed to the lowering rule will added to this dictionary. lowering_args: Any other arguments to :func:`mlir.custom_call` will also be passed through if provided as extra arguments to this function. skip_ffi_layout_processing: If true, skip processing of operand and result layout arguments passed to the lowering rule. c Rtf d|g|i|}|jS)Nrk)rresults) rr{rresultrnrrrlrmros r% _loweringzffi_lowering.._lowering4sZ ('%%#=        F >>r')rrr{rrrrz'Sequence[ir.Value | Sequence[ir.Value]]r1)rrlrmrnrorrs`````` r% ffi_loweringrs6H  ! .6 BE .   r'c&td|DS)Nc3FK|]}tj|ywrw)r shaped_abstractifyryrs r%rzz _result_avals..Js;at&&q);s!rb)rs r% _result_avalsrHs ;7; ;;r'ct|tjrt|tjryt|ddt|ddk7ryt|ddt|ddk7ryy)NTr]r1Fdtype)rfr AbstractTokengetattr)abs r%_check_compatible_avalsrLsa4%%&:a9K9K+L  Q7B 77  Q7B 77  r'c:tdt||DS)Nc3nK|]-\}}t||t|tr|n|ddd/yw)Nra)rjrfr)ryr^ris r%rzz0_convert_layouts_for_ffi_call..YsA/ $ # Nj&@&dd|/s35)rbr)avalslayoutss r%_convert_layouts_for_ffi_callrVs' / eW- / //r'.)has_side_effect vmap_method input_layoutsoutput_layoutsinput_output_aliasescustom_call_api_versionlegacy_backend_config vectorizedcyrwr1 target_nameresult_shape_dtypesrrrrrrrrs r%ffi_callrcr'cyrwr1rs r%rrtrr'rrc 2 t| ts tdgd} | vrtd| dttr| d t nd t g|f dk\rtdd d  f d } | S) a Call a foreign function interface (FFI) target. See the :ref:`ffi-tutorial` tutorial for more information. Like :func:`~jax.pure_callback`, the behavior of ``ffi_call`` under :func:`~jax.vmap` depends on the value of ``vmap_method``. See the :func:`~jax.pure_callback` documentation for more details about the allowed values and examples of their behavior. 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: target_name: the name of the XLA FFI custom call target that was registered using :func:`~jax.ffi.register_ffi_target`. result_shape_dtypes: an object, or sequence of objects, with ``shape`` and ``dtype`` attributes which are expected to match the shape and dtype of the custom call output or outputs. :class:`~jax.ShapeDtypeStruct` is often used to define the elements of ``result_shape_dtypes``. ``jax.core.abstract_token`` may be used to represent a token-typed output. has_side_effect: boolean specifying whether the custom call has side effects. When ``True``, the FFI call will be executed even when the outputs are not used. vmap_method: string specifying how the FFI call transforms under :func:`~jax.vmap` as described above. input_layouts: a sequence of layouts for each input argument. In each case, the layout can be (a) ``None`` indicating that this input is in default row-major order, (b) a ``Layout`` specifying the axis order, or (c) a sequence of integers specifying the major-to-minor axis ordering. Users who are familiar with XLA layouts should note that this function expects layouts in major-to-minor order instead of the minor-to-major order that XLA uses. For example, a batch of row-major matrices could be specified using the layout ``[0, 1, 2]``, whereas a batch of column-major matrices would have layout ``[0, 2, 1]``. In both of these examples, the leading/batch dimension is the "slowest" axis. The ``input_layouts`` parameter should be used to request the memory layout expected by the FFI call target, and XLA will ensure that the buffers have the correct layouts before the handler is executed. output_layouts: like ``input_layouts``, but specifying the required layouts for the output arrays. input_output_aliases: a dictionary where the keys are input indices and the values are output indices. This mapping indicates which output arrays alias specific input arrays. custom_call_api_version: the version number of the custom call API implemented by the FFI target ``target_name``. The only formally supported version is the typed FFI API with ``custom_call_api_version=4``, but earlier unsupported custom calls can be executed using this argument. legacy_backend_config: for legacy targets implemented using ``custom_call_api_version<4``, attributes are passed using the opaque string representation provided by this argument. This parameter cannot be used with ``custom_call_api_version>=4``. Returns: A function that can be called with the input arrays as positional arguments to execute the FFI handler. Any keyword arguments are passed as named attributes to the FFI handler using XLA's FFI interface. zcThe 'vectorized' argument of jax.ffi.ffi_call was removed in JAX v0.6.0. Use 'vmap_method' instead.) sequentialsequential_unrolled expand_dims broadcast_alllegacy_vectorizedNz.vmap_method must be on of the allowed methods z , but got: TFrrzYThe use of the legacy_backend_config parameter requires custom_call_api_version < 4; got rsc |Dcgc]}tj|}}ttt|}nGt t |k7r$t dt |dt dt|}ttt}nGt t k7r$t dt dt dt}d}tjD]\}}t|t|}}|t |k\r!t d|d|d |d t |d |t k\r!t d|d|d |d t d ||} |} t| | st d|d|d | d | d ||||k7rt d|d|d||d||d |||ffz }tj|}tj| ||| t|d } rt!tr t| S| S| dScc}w)NzThe number of input arguments (z*) must equal the number of input layouts (z).zThe number of outputs (z+) must equal the number of output layouts (r1z+input_output_aliases contains the mapping ':z' with input index z outside the range [0, z' with output index z,' referring to an input with abstract value z/ and an output with a different abstract value rsz$' referring to an input with layout z' and an output with a different layout ) result_avalsrrrrrrrr attributesr)r get_avalrbrrjrer4rsortedr~intrstandard_insert_pvary ffi_call_pbind_wrap_kwargs_hashablerf)rr$xin_avalsstatic_input_layoutsstatic_output_layoutsstatic_input_output_aliasesi_idxo_idxin_avalout_avalrrrrrrmultiple_resultsoutput_layouts_rrrrs r%wrappedzffi_call..wrappeds*./Q a /H/"3'CX#NO ] s8} ,-c(m_=((+M(:';2 ?@ @;8;HJ#C(D(4%67 _ \!2 2%c,&7%89)),_)=(>b BC C9,%5z3u:u CI ;E7!E7K""'(?T 2   C % %;E7!E7K##(')@\"#2'( (5/&&w9;E7!E7K;;B)D88@zDE E  &*?*F F;E7!E7K33G3N2OP7&u-.a12 2 $'88#394  % %t ,Doo !'*,8 73(0 G ' /W~ n QZK0sH?)rrr$r)rfrr4rr)rrrrrrrrrrallowed_vmap_methodsrrrrs````` ``` @@@r%rrsR J .  - ..F,,  89M8NO= " ## #X.$O !45L "5!67L%'O!&;&G  ,,C+DA G HHFFP .r'cg}t|jD]\}}t|tjr|j |t |f=t|tr|j |t|fj t||j ||ft|S#t$r}td|d||d}~wwxYw)Nz*Non-hashable keyword argument to ffi_call z: ) rr~rfnpndarrayappendrr|rhash TypeErrorrb)r$hashable_kwargsrres r%rr8s+-/V\\^$ 'da!RZZ aq!123 At aA/0' Q 1v& '   J82aS ACHI JJs B22 C;C  Cci}|D]J\}}t|tr|j||<&t|tr|j||<F|||<L|Srw)rfrvalr_d)r$unwrapped_kwargsrrs r%_unwrap_kwargs_hashablerJs`%'da!]#EEq Az "DDqq  r'T)frozenceZdZdZy) FfiEffectcy)NFFIr1)selfs r%__str__zFfiEffect.__str__Xs r'N)r,r-r.rr1r'r%rrVsr'rctjdg||rthntj}t d|D|fS)Nrc3K|]h}|tjur|nO|j|jjj rtj n |jjyw))shardingN)r r[updatermeshemptyget_cur_mesh_shardingrs r%rzz)ffi_call_abstract_eval..ise&,,,qxx%&ZZ__%:%:#'"<"<">@A MM&sA.A0)r standard_vma_rule _FfiEffect no_effectsrb)rrr_r s r%ffi_call_abstract_evalrasL /h/+ZL' &%& &(/ //r'c"~td|dNzThe FFI call to `z\` cannot be differentiated. You can use `jax.custom_jvp` or `jax.custom_jvp` to add support.r3rrrs r% ffi_call_jvpro' +'II JJr'c"~td|drr3rs r%ffi_call_transposervrr'c bt||||t|||} | |g| it|S)N)rrlrmoperand_output_aliasesr#rn)rr|r) rrrrrrrrrr{rrules r%ffi_call_loweringr}sD k?&3%3-12F-G"9%:  <$ c DH D 7 C DDr'c ddlm}ddlm}t||D chc]&\}} | tj ur|j | (c} }\t||D cgc]0\} } | tj ur| nt j| | d2} } } tfdD} dvr)jdtddDd<dk(rJjd "td td |Dd <j| | d }n!d k(sd k(rd k(rnd}t| |D cgc],\}} | tj ur|j||fn|.}}} jd tdd Dd <j|| d }ndk(sdk(rf|D cgc]} | tj uc} tj| \}fddk(}fd}|j|d||\}}ntdj ddt|dt#|zfScc} }wcc} } wcc} }wcc} w)Nr) control_flow)laxc3LK|]}tjd|yw)rN)r unmapped_aval)ryr^ axis_sizes r%rzz$ffi_batching_rule..s'G15dAt,Gs!$)rrrc3NK|]}|dntd|Ddzyw)Nc3&K|] }|dz ywNr1ryns r%rzz.ffi_batching_rule...s)@A!a%)@rrryris r%rzz$ffi_batching_rule..s1%0 E)@)@$@4$GG%0#%rrc3|K|]4\}}|tjur|n|dntd|Ddz6yw)Nc3&K|] }|dz ywrr1rs r%rzz.ffi_batching_rule...s,CqQU,Crr)r not_mappedrb)ryrids r%rzz$ffi_batching_rule..sJ&?fa,,,&>4u,CF,C'Cd'J L&?s:<)rrrrrc3NK|]}|dntd|Ddzyw)Nc3&K|] }|dz ywrr1rs r%rzz.ffi_batching_rule...s+BaAE+Brrrrs r%rzz$ffi_batching_rule..s1&1.$e+B6+B&BT&I I&1rrrc\tj|}j|dS)N)rr)r merge_listsr) batched_args merged_args is_batchedr$primrunbatched_argsrs r% _batch_funz%ffi_batching_rule.._batch_funs=$$ZNk TYY #!  r'cd|fSrZr1)rrr#s r%z#ffi_batching_rule..sb*Q-(r'r1)unrollzvmap is only supported for the z primitive when vmap_method is one of 'sequential', 'sequential_unrolled', 'expand_dims', 'broadcast_all', or 'legacy_vectorized'. Got vmap_method=rsr) jax._src.laxrrrrrr]moveaxisrbgetr broadcastrpartition_listscanNotImplementedErrorr re)r!rdimsrrr$rrrrargdimnew_argsbatched_result_avalsoutvalssizer bcast_argsrr&grr#r r r"s` ``` @@@@r%ffi_batching_ruler7s('*41tq!H/// 1*)>At_N19cH///cS!,-N(NG9EGG >> jj!".$%0-.%0 0F ''zz/". %&?vo6=&?!?f_dii )  G m#{o'E#69AD$') Aq&'(*=*=%= a$!1D)J)zz/". %&1/&1!1f_dii'  G l"k5J&J8<=1!8...=J#'#6#6z8#L NL1 1F(A""1b,v"FJAw  )$))5E8C~Q H II wG , ,,C1N>)>s+I&$5I,1I2I8r)cpur) r strr!rr"r9r#rr$rrNone)r8)r r9r5rr"r9rr:)r r9r9r)r"r9rr:)r r9rr:)rr9)r^core.AbstractValuerrrw)r^r;riFfiLayoutOptionsrz Sequence[int])rr9rl!Sequence[FfiLayoutOptions] | Nonermr=rn'Mapping[str, ir.Attribute] | str | NoneroboolrrrzCallable[..., ir.Operation])rr9rlr=rmr=rnr>ror?rrrzmlir.LoweringRule)rSequence[ResultMetadata]rtuple[core.AbstractValue, ...])rr;rr;rr?)rzSequence[core.AbstractValue]rzSequence[FfiLayoutOptions]rztuple[Sequence[int], ...])rr9rResultMetadatarr?r str | Nonerr=r4FfiLayoutOptions | Sequence[FfiLayoutOptions] | Nonerdict[int, int] | NonerrrrCrbool | None | DeprecatedArgrzCallable[..., Array])rr9rr@rr?rrCrr=rrDrrErrrrCrrFrzCallable[..., Sequence[Array]])rr9rz)ResultMetadata | Sequence[ResultMetadata]rr?rrCrr=rrDrrErrrrCrrFrz&Callable[..., Array | Sequence[Array]])r$dict[str, Any]rSequence[tuple[str, Any]])r$rHrrG)rrArr?)rrr{rrr9rr?rSequence[Sequence[int]]rrIrzSequence[tuple[int, int]]rrrrCrrHrzSequence[ir.Value])rrCrzSequence[core.ShapedArray]r$r)Z __future__rcollections.abcrrrrB dataclassesr>rQtypingrrr r numpyrjax._srcr r r rrjax._src.hashable_arrayrjax._src.frozen_dictrjax._src.interpretersrrrjax._src.layoutr jax._src.librrjax._src.lib.mlirrjax._src.typingrrrrrsafe_mapr unsafe_maprr<r&r)r6r:r@rNrXr_rjrrrrBrrrrrr dataclassEffectrrlowerable_effectsadd_typecontrol_flow_allowed_effectsrrrrr7 Primitiverr simple_impldef_effectful_abstract_evalprimitive_jvpsprimitive_transposesr?primitive_batchersregister_loweringr1r'r%res#77  881+$*&"# $$--ZC=6)D0  : : :: :  :  :2 y " I  I  I I I$ ' " /D- O :>   &6 BO  :>8<>B', RR7R6 R < R !% RR!Rp:>8<>B', 22726 2 < 2 !% 222j $"4"44<5< / ' / ' /,E /  !7:KN25#&(+.;o  '     5 I 0 ! & ,       !7:KN25#&(+.;o  1     5 I 0 ! & , $   (""7;KO26#$(,.;okkBk k  k 5 kIk0k!k&k,k,kf $ d#$ [ ""9-$$--i8 /0 / /JJE !EEE E + E , E4E!E&E*EE,M-  M- - M-M-`T^^J ' " Z &&'=> ,*&8 #*;)*;*;z+#J'z#45r'