VpfW dZddlmZddlmZddlZddlmZddlZddl Z ddl m Z ddl m Z ddl m Z dd l mZdd l mZdd l mZdd l mZdd l mZddl mZddlmZddlmZddlmZddlmZddl m!Z!ddl"m#Z#dZ$GddZ%ej&Z'dZ(dZ)dZ*dZ+dZ,dZ-dZ.dZ/d Z0e j1d!Z2d"e2_3ej45e2d#Z6d$Z7e28e6e29e7d%Z:ej;Gd&d!Ze2e=ej?e(e-e/e0d"*ej@eej?e(e-e/e0d"+dS)-zrThe implementation of custom partitioning APIs. It was moved out of ``jax.experimental`` to avoid import cycles. ) annotations)partialN)Any) tree_util)api_util)core)custom_api_util)dispatch) linear_util)mesh)sharding_impls) xla_bridge)mlir) partial_eval) xla_client)ir)hlo)UnexpectedTracerErrorctj|j|i|}||jrt d|jS)Nz4keyword arguments could not be resolved to positions)inspect signaturebindapply_defaultskwargs TypeErrorargs)funrrbas \/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/_src/custom_partitioning.py_resolve_kwargsr -sV"w"D3F33"Y J K KK 7Nc eZdZdZdZdZdS)_ShardingCallbackInfoc ||_||_||_||_||_||_||_||_| |_dSN) propagate_user_sharding partitionto_mesh_pspec_shardingin_treeout_treeinfer_sharding_from_operandsmodule_contextr static_args) selfr&r'r(r)r*r+r,r r-s r__init__z_ShardingCallbackInfo.__init__8sN$;D DN"8DDLDM(DD%(DDI"Dr!c t|||t|Sr%)_to_jax_sharded_shaper(len dimensions)r.sshardings runflatten_arg_shapez)_ShardingCallbackInfo.unflatten_arg_shapeEs9  4 & &xQ\\^^1D1D E E  r!cljfdt||DS)NcBg|]\}}||Sr6).0r4r5r.s r z>_ShardingCallbackInfo.unflatten_arg_shapes..Ls=   8  $ $Q 1 1   r!)r) unflattenzip)r. arg_shapes arg_shardingss` runflatten_arg_shapesz*_ShardingCallbackInfo.unflatten_arg_shapesJsM < ! !    ":}==     r!N)__name__ __module__ __qualname__r/r6rAr9r!rr#r#6sA # # # r!r#CustomSPMDPartitioningcttj||Sr%)r ShapedArrayr3 numpy_dtype)r4s r _to_jax_shaperIXs% !,,..!--// : ::r!cxtj|||S)N)r5)jaxShapeDtypeStructr3rH)r4r5s rr1r1\s3 llnnammoo   r!cz|r tj||S|dS)Nr)is_tuplexc HloShardingtuple_sharding)shaperesult_shardingss r_pack_result_shardingrTbs8 ^^ > ( (0@ A AA A r!c\dt|||DS)Nc rg|]4\}}t|t|5Sr9)_to_hlo_shardingr2r3)r;r5rRs rr<z%_flatten_sharding..jsJ    (ExU%5%5%7%7!8!899   r!)r> flatten_up_to)tree shardingsshapess r_flatten_shardingr\is@     Y ' '   r!ct|j|S|r)|}|}n|f}|f}jfdt||D}jgjj |R}tj||}t||S)NcBg|]\}}||Sr9r:r;r4r5infos rr<z@_custom_partitioning_propagate_user_sharding..~s=a  " "1h / /r!) _sharding_callbacksr&rN tuple_shapestuple_elementsr*r=r>r-r r\rT) user_shardingrRbackend_string user_shapesuser_shardings user_shaperesult_shardingrSr`s @r,_custom_partitioning_propagate_user_shardingrjrs  ^ ,$ !)  ^^&$$&&K"1133NN(K#%N}&& n==* 1D0 $./' m_k33 u&6 7 77r!ct|tjjst d||S)Nz/Custom Partitioning rules must return Sharding.) isinstancerKr5Sharding ValueError_to_xla_hlo_sharding)r5num_dimensionss rrWrWs= Hcl3 4 4H F G GG  & &~ 6 66r!c t||r)|}|}n|f}|f}jgjj||j fdt||DR\}}}}j } tj||}tj ||}dt||D} dt||D} tj|t!|j| } | j| kr4t)dt+| jdt+| t-j|} t1j|j5t5j| d| j| j| t?|j }dddn #1swxYwYtC||}t5j"|||fS) NcBg|]\}}||Sr9r:r_s rr<z2_custom_partitioning_partition..s=   !X&&q(33   r!cXg|]'\}}t||(Sr9rItiler;r5r4s rr<z2_custom_partitioning_partition..s@ (AHMM!$$%%r!cXg|]'\}}t||(Sr9rtrvs rr<z2_custom_partitioning_partition..s@ (AHMM!$$%%r!)axis_envzMismatch in result shapes. z vs tmp_xla_computation)name platformsbackend_or_name axis_context)#rarNrbrcr'r-r rAr*r=r>r,r\r)rK make_jaxprlistrRitems out_avalsrnreprr SPMDAxisContextrextend_axis_env_ndrbuild_mlir_module_helperr{r| extend_manual frozenset axis_namesrTmodule_to_bytecode)r?r@ result_shaperire result_shapesrSr lower_fnr, tiled_args tiled_results closed_jaxprr}moduler`s @r_custom_partitioning_partitionrs ^ ,$* --//M&5577!OM')3A4> 4  4 i 4  M:: 4 m    !$]4D!E!E    4 4 40$/=&.& m_m55#DL-LL-]J77*-}==-M4 8H8H8J8J3K3KLLL,},, *  & ' ' ' 'm)<)<)< >   /55, tz//1122  * " *&6!// $/0J0JKK F*,8HII/  ( (- HHsA H11H58H5cft|}|r|}n|f}|jg|j|j||||jd|DR}t|j||}t||S)Nc,g|]}t|Sr9)rIr;r4s rr<zE_custom_partitioning_infer_sharding_from_operands..s GGGA}Q//GGGr!) rarNrbr+r-r rAr*r=r\rT)r?r@rrer`rrirSs r1_custom_partitioning_infer_sharding_from_operandsrs ^ ,$$ --//MM!OM5D5  i  M:: mGGGGGHH / ' m_m55 |-= > >>r!custom_partitioningTc~~~~~~~|jSr%)r) callr)r*r&r'r+decode_shardingsr-avalss r"_custom_partitioning_abstract_evalrs x0)"$4k r!c>~~~~~~~tj||Sr%)r jaxpr_as_fun) rr)r*r&r'r+rr-rs r_custom_partitioning_implrs.x0)"$4k  4 $ ''r!cztdtj|DrtddS)Nc3JK|]}t|tjVdSr%)rlrTracer)r;leafs r z%_check_for_tracers..s.LL4D$+ & &LLLLLLr!zFound a JAX Tracer object passed as an argument to acustom_partitioning function in a position indicated as static bystatic_argnums. )anyr tree_leavesr)xs r_check_for_tracersrsMLL93H3K3KLLLLL    r!cDeZdZUdZd dZejZded< d dZ d Z dS) raPInserts a CustomCallOp into the XLA graph with custom SPMD lowering rules. .. code-block:: python @custom_partitioning def f(*args): return ... def propagate_user_sharding(mesh, user_shape): '''Update the sharding of the op from a user's shape.sharding.''' user_sharding = jax.tree.map(lambda x: x.sharding, user_shape) def partition(mesh, arg_shapes, result_shape): def lower_fn(*args): ... builds computation on per-device shapes ... result_shardings = jax.tree.map(lambda x: x.sharding, result_shape) arg_shardings = jax.tree.map(lambda x: x.sharding, arg_shapes) # result_sharding and arg_shardings may optionally be modified and the # partitioner will insert collectives to reshape. return mesh, lower_fn, result_sharding, arg_shardings def infer_sharding_from_operands(mesh, arg_shapes, shape): '''Compute the result sharding from the sharding of the operands.''' arg_shardings = jax.tree.map(lambda x: x.sharding, arg_shapes) f.def_partition(partition, propagate_user_sharding, infer_sharding_from_operands) The args to ``def_partition`` are as follows: * ``propagate_user_sharding``: Callable which takes the sharding of a user (in the dag) and returns a suggestion for a new `NamedSharding`. The default implementation is just to return the suggested sharding. * ``partition``: Callable which takes the SPMD suggested partition shapes and partition specs and returns the mesh, a per-shard lowering function, and the final input and output sharding specs (the SPMD partitioner will repartition the inputs to match). The mesh is returned to allow configuring axis_names for collectives when no mesh is provided. * ``infer_sharding_from_operands``: Callable which computes an output ``NamedSharding`` from the ``NamedSharding`` chosen for each argument. * ``decode_shardings``: When set to True, convert input ``GSPMDSharding``s to ``NamedSharding`` if possible. This may not be possible if the user does not provide a contextual mesh. Positional arguments can be specified as static using static_argnums. JAX uses :code:`inspect.signature(fun)` to resolve these positional arguments. Examples: As an example, assume we want to enhance the existing ``jax.numpy.fft.fft``. This function computes the discrete Fourier transform of an N-dimensional input along the last dimension, and is batched along the first N-1 dimensions. By default, however, it will ignore the sharding of the input and gather the input on all devices. However, since ``jax.numpy.fft.fft`` is batched along the first N-1 dimensions, this is unnecessary. We will create a new ``my_fft`` op that, instead, does not alter the sharding along the first `N-1` dimensions, and only gathers the input along the last dimension if needed. .. code-block:: python import jax from jax.sharding import NamedSharding from jax.experimental.custom_partitioning import custom_partitioning from jax.experimental.pjit import pjit from jax.sharding import PartitionSpec as P from jax.sharding import Mesh from jax.numpy.fft import fft import regex as re import numpy as np # Pattern to detect all-gather or dynamic-slice in the generated HLO _PATTERN = '(dynamic-slice|all-gather)' # For an N-D input, keeps sharding along the first N-1 dimensions # but replicate along the last dimension def supported_sharding(sharding, shape): rank = len(shape.shape) max_shared_dims = min(len(sharding.spec), rank-1) names = tuple(sharding.spec[:max_shared_dims]) + tuple(None for _ in range(rank - max_shared_dims)) return NamedSharding(sharding.mesh, P(*names)) def partition(mesh, arg_shapes, result_shape): result_shardings = jax.tree.map(lambda x: x.sharding, result_shape) arg_shardings = jax.tree.map(lambda x: x.sharding, arg_shapes) return mesh, fft, supported_sharding(arg_shardings[0], arg_shapes[0]), (supported_sharding(arg_shardings[0], arg_shapes[0]),) def infer_sharding_from_operands(mesh, arg_shapes, result_shape): arg_shardings = jax.tree.map(lambda x: x.sharding, arg_shapes) return supported_sharding(arg_shardings[0], arg_shapes[0]) @custom_partitioning def my_fft(x): return fft(x) my_fft.def_partition( infer_sharding_from_operands=infer_sharding_from_operands, partition=partition) Now create a 2D array sharded along the first axis, pass it through ``my_fft`` and notice how it is still sharded as expected, and identical to the output of ``fft``. However, inspecting the HLO (using ``lower(x).compile().runtime_executable().hlo_modules()``) reveals that ``my_fft`` does not create any all-gather or dynamic-slice, while ``fft`` does. .. code-block:: with Mesh(np.array(jax.devices()), ('x',)): x = np.asarray(np.random.randn(32*1024, 1024), dtype=np.complex64) y = pjit(lambda x: x, in_shardings=None, out_shardings=P('x'))(x) pjit_my_fft = pjit(my_fft, in_shardings=P('x'), out_shardings=P('x')) pjit_fft = pjit(fft, in_shardings=P('x'), out_shardings=P('x')) print(pjit_my_fft(y)) print(pjit_fft(y)) # dynamic-slice or all-gather are not present in the HLO for my_fft, because x is a 2D array assert(re.search(_PATTERN, pjit_my_fft.lower(x).compile().runtime_executable().hlo_modules()[0].to_string()) is None) # dynamic-slice or all-gather are present in the HLO for fft assert(re.search(_PATTERN, pjit_fft.lower(x).compile().runtime_executable().hlo_modules()[0].to_string()) is not None) .. code-block:: # my_fft [[-38.840824 +0.j -40.649452 +11.845365j ... -1.6937828 +0.8402481j 15.999859 -4.0156755j]] # jax.numpy.fft.fft [[-38.840824 +0.j -40.649452 +11.845365j ... -1.6937828 +0.8402481j 15.999859 -4.0156755j]] Because of the logic in ``supported_sharding``, ``my_fft`` also works on 1-dimensional arrays. However, in this case, the HLO of ``my_fft`` does show a dynamic-slice, since the last dimension is the dimension along which FFTs are calculated and needs to be replicated on all devices before the computation can be done. .. code-block:: with Mesh(np.array(jax.devices()), ('x',)): x = np.asarray(np.random.randn(32*1024*1024), dtype=np.complex64) y = pjit(lambda x: x, in_shardings=None, out_shardings=P('x'))(x) pjit_my_fft = pjit(my_fft, in_shardings=P('x'), out_shardings=P('x')) pjit_fft = pjit(fft, in_shardings=P('x'), out_shardings=P('x')) print(pjit_my_fft(y)) print(pjit_fft(y)) # dynamic-slice or all-gather are present in the HLO for my_fft, because x is a 1D array assert(re.search(_PATTERN, pjit_my_fft.lower(x).compile().runtime_executable().hlo_modules()[0].to_string()) is None) # dynamic-slice or all-gather are present in the HLO for fft assert(re.search(_PATTERN, pjit_fft.lower(x).compile().runtime_executable().hlo_modules()[0].to_string()) is not None) .. code-block:: # my_fft [ 7.217285 +0.j -3012.4937 +4287.635j -405.83594 +3042.984j ... 1422.4502 +7271.4297j -405.84033 -3042.983j -3012.4963 -4287.6343j] # jax.numpy.fft.fft [ 7.217285 +0.j -3012.4937 +4287.635j -405.83594 +3042.984j ... 1422.4502 +7271.4297j -405.84033 -3042.983j -3012.4963 -4287.6343j] r9cL||_d|_||_d|_d|_dSr%)rr'static_argnumsr&r+)r.rrs rr/zcustom_partitioning.__init__s.DHDN(D#'D (,D%%%r!r __getattr__NTc>||_||_||_||_|Sr%)r'r&r+r)r.r'r+r&rs r def_partitionz!custom_partitioning.def_partitions'DN#:D (DD%,D r!c t|j||jrt|jt fdt Dfdt tD}tj tj |j|d\}}fd|jD}t|ng}tj |j}}tj|\}}tj||\} } d|D} t!j|j|| dd} t!j| | | \} }}\t|rJt'jt!j| d}t-jg||R||j|j|j|j|| |d }tj| |S) Nc3.K|]\}}|vr|n|VdSr%r9)r;irrs rrz/custom_partitioning.__call__..s5OOtq!^++11OOOOOOr!cg|]}|v| Sr9r9)r;rrs rr<z0custom_partitioning.__call__..s#LLL1A^4K4KQ4K4K4Kr!F)require_static_args_hashablec g|] }| Sr9r9)r;rrs rr<z0custom_partitioning.__call__..s:::T!W:::r!cZg|](}tjtj|)Sr9)rraise_to_shapedget_aval)r;rs rr<z0custom_partitioning.__call__..s-JJJ1$T]1%5%566JJJr!rr9)rr'r&r+rr)r*r-)r rrsettuple enumerateranger2rargnums_partiallu wrap_initrr tree_flattenflatten_fun_nokwargspe debug_infotrace_to_jaxpr_dynamicr ClosedJaxprconvert_constvars_jaxprcustom_partitioning_prr'r&r+rtree_unflatten)r.rr dyn_argnumsf_dyn_argsr- args_flatr)flat_funr*in_avalsdebugjaxpr_consts closed_callout_flatrs ` @r__call__zcustom_partitioning.__call__s 48T6 2 2D 24.//n OOOOyOOO O OdLLLLc$ii 0 0LLLk- ,tx  ', lb( ;:::d&9:::k%%%%k\$(++T(b"/99Iw!6r7CCHhJJ JJJH M$(GXu/ 1 1E4XxOOE1fb6{{"2#=e#D#DbIIK$)      . $ <%)%F.   H  #HHJJ 9 99r!)r9)NT) rBrCrD__doc__r/r forward_attrr__annotations__rrr9r!rrrswbbH----%1+1111DH%:%:%:%:%:r!ctxmlir.LoweringRuleContextc tjjj|jj} t | tjr^t| j t| j j kr/tjtj|d|g| RSt | tjr| jt'dn)t | tjr | j jndrt+dkr/tjtj|d|g| RSd fd } t-||| ||||j| } t/t1| } | t2t5| d<|j| d |jD}t;j|t?| t@j!"tFt@j$"d tj%d t@j&"gt@j!"| dd }|j'S)NT)multiple_resultsz9Please file a bug at https://github.com/google/jax/issues hlo_shardingxc.HloSharding | Nonec@||SjssJtj|Stj|d}t jjg|d|t|z zR}t j |S)Nrr%) emptyr _op_sharding_to_pos_shardingparse_flatten_op_shardingget_partition_specrKr5 PartitionSpecr2 NamedSharding)rndimpspecrdevicesr s rr(zB_custom_partitioning_lowering_rule..to_mesh_pspec_shardings  zP)P   8w O OO  4d   1133 L & P PD3u::. s#BBBa$&q))BBBr!F)call_target_namehas_side_effect api_versioncalled_computationsbackend_configoperand_layoutsresult_layouts)rr)(mesh_libthread_resourcesenv physical_meshr,r}rlr rr manual_axesr rr lower_funrrShardingContextdevice_assignmentAssertionError_flat_devices_tupler2r#stridrabytes add_keepaliverr CustomCallOprr StringAttrget_CUSTOM_PARTITIONING_CALL_NAMEBoolAttri32_attr ArrayAttrresults)rrr)r*r&r'r+rr-valuesr}r(sharding_callback_infokey result_typesoutrr s ` @@r"_custom_partitioning_lowering_rulers  " & 4$#0,~=>>X , "##s<+<+G'H'HHH I4>$+D11D I I I# WPV W W WW n<==,G  E G GG, >??3GGG FCLLA%% 84> $$ 8 8 88; F>D F F FF 3 3 3 3 3 3 3 311H'("C$6kKK B% & &''#,BeC(()""#9:::BB4>BBB,  6ll}(()GHHkooe,,-"",**2..]&&s++   # r!)+can_side_effecting_have_replicated_sharding)rzprop_user_shardingr'r+r)rr)Ar __future__r functoolsrrtypingrweakrefrKrjax._srcrrr r r rr rr rxbjax._src.interpretersrrr jax._src.librrOjax._src.lib.mlirrjax._src.lib.mlir.dialectsr jax.errorsrr r#WeakValueDictionaryrarrIr1rTr\rjrWrr Primitiverr%prim_requires_devices_during_loweringaddrrdef_abstract_evaldef_implrregister_custom_decorator_typerrregister_lowering register_custom_call_partitionerregister_plugin_callbacksr9r!rr&s #""""" $$$$$$&&&&&&%%%%%%######%%%%%%&&&&&&444444)))))) ******,,,,,,:2g133!9;;; 8882777 3I3I3Il???&''<==)-&.223HIII(((''(JKKK8999/[:[:[:[:[:[:[:0/[:|9999v,9;;;$#"0"504  G + +G0%V48      r!