uki$o HdZddlmZddlmZddlZddlmZddlm Z ddl Z ddl Z ddl mZddl mZdd l mZdd l mZdd l mZdd l mZdd l mZddl mZddl mZddl mZddl mZddl mZddlmZmZm Z ddl!m"Z"ddl!m#Z$ddl%m&Z'ddl(m)Z)ddl*m+Z+ddl,m-Z-dZ.GddZ/e j`Z1dZ2dZ3dZ4d Z5d!Z6d"Z7d#Z8d$Z9d%Z:ejvd&ZsH$;D DN"8DDLDM(DD%(DDI"Dr(c ht||j|t|jSr,)_to_jax_sharded_shaper/len dimensions)r5sshardings r&unflatten_arg_shapez)_ShardingCallbackInfo.unflatten_arg_shapeKs,  4 & &xQ\\^1D E r(c |jjt||Dcgc]\}}|j||c}}Scc}}wr,)r0 unflattenzipr=)r5 arg_shapes arg_shardingsr;r<s r&unflatten_arg_shapesz*_ShardingCallbackInfo.unflatten_arg_shapesPsN << ! ! #:}= 8  $ $Q 1   sA N)__name__ __module__ __qualname__r6r=rCr(r&r*r*<s # r(r*CustomSPMDPartitioningcftj|j|jSr,)r ShapedArrayr: numpy_dtype)r;s r& _to_jax_shaperL^s!  !,,.!--/ ::r(cjtj|j|j|S)N)r<)rShapeDtypeStructr:rK)r;r<s r&r8r8bs(  llnammo r(cl|jr tjj||S|dS)Nr)is_tuplexc HloShardingtuple_sharding)shaperesult_shardingss r&_pack_result_shardingrVhs/ ^^ >> ( (0@ AA A r(c t|j||Dcgc](\}}t|t|j *c}}Scc}}wr,)r@ flatten_up_to_to_hlo_shardingr9r:)tree shardingsshapesr<rTs r&_flatten_shardingr]osR!   Y '  (ExU%5%5%7!89  s-Ac t|}|j|S|jr!|j}|j }n|f}|f}|j j t||Dcgc]\}}|j||c}}}|jg|j|j|} t|j | |} t|| Scc}}wr,) _sharding_callbacksr-rP tuple_shapestuple_elementsr1r?r@r=r4rr]rV) user_shardingrTbackend_stringinfo user_shapesuser_shardingsr;r< user_shaperesult_shardingrUs r&,_custom_partitioning_propagate_user_shardingrixs ^ ,$ !!)  ^^$$&K"113N(K#%N}}&&!n=a  " "1h /* 1D00 $./' mm_k3 u&6 77s3C$ cZt|ts td|j|S)Nz/Custom Partitioning rules must return Sharding.) isinstancer ValueError_to_xla_hlo_sharding)r<num_dimensionss r&rYrYs) Hh ' F GG  & &~ 66r(c Zt|}|jr!|j}|j}n|f}|f}|jg|j |j |j|||jjt||D cgc]\}} |j|| c} }\} } }}|j} t|j||}t|j||}t||D cgc]\} }t| j!|!} } }t||D cgc]\} }t| j!|!}} }t#j$| t'| j(j+|jj| }|j,Dcgc]}|j(|j.fc}|Dcgc]}|j(|j.fc}k7r-t1dt3|j,dt3|t5j6| }t9j:| j(j+5t=j>|d| j@| jB|jEtG| jH}dddtK||}t=jL||fScc} }wcc}} wcc}} wcc}wcc}w#1swYIxYw)N)axis_envzMismatch in result shapes. z vs tmp_xla_computation)name platformsbackend axis_context)'r_rPr`rar.r4rrCr1r?r@r=r3r]r0rLtiler make_jaxprlistrTitems out_avalsdtyperlreprrSPMDAxisContextr extend_axis_env_ndrbuild_mlir_module_helperrsrt extend_manual frozenset axis_namesrVmodule_to_bytecode)rArB result_shaperhrcrd result_shapesrUr;r<rlower_fnr3 tiled_args tiled_results closed_jaxprotrumodules r&_custom_partitioning_partitionrs ^ ,$ --/M&557!OM')3A4>> 4  4 ii 4  M: 4 mm"%]4D!E !X&&q(3  40$/=&&.& mm_m5#DLL-L-]J7 (AHMM!$%* -}= (AHMM!$%-M4 8H8H8J3KL ||j),$0#9#9:a:#01a12   && 'm)< >  //5, tzz//12  * * " **&&!// $//0JK F*,8HI/  (- HHM ;1 s+L$L $LL5LAL!!L*c t|}|jr|j}n|f}|jg|j|j |j |||jj|Dcgc] }t|c}}t|j||}t||Scc}wr,) r_rPr`r2r4rrCr1r?rLr]rV) rArBrrcrdrr;rhrUs r&1_custom_partitioning_infer_sharding_from_operandsrs ^ ,$ --/M!OM5D55  ii  M: mmGA}Q/GH / ' mm_m5 |-= >> Hs>B=custom_partitioningTc *~~~~~~~~|jSr,)rz) callr0r1r-r.r2decode_shardings sharding_ruler4avalss r&"_custom_partitioning_abstract_evalrs# x0)"$4m r(c B~~~~~~~~tj|| Sr,)r jaxpr_as_fun) rr0r1r-r.r2rrr4r#s r&_custom_partitioning_implrs0x0)"$4k=   4 $ ''r(cxtdtj|Drtjdy)Nc3PK|]}t|tj ywr,)rkr Tracer).0leafs r& z%_check_for_tracers..sL4D$++ &Ls$&zFound a JAX Tracer object passed as an argument to acustom_partitioning function in a position indicated as static bystatic_argnums. )anyr tree_leavesr UnexpectedTracerError)xs r&_check_for_tracersrs9L93H3H3KLL  & &  Mr(cZeZdZUdZd dZej Zded< d dddddZ d Z y) ra'Inserts 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=infer_sharding_from_operands, sharding_rule='i j -> 'i j') 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 value is None. A trivial implementation is just to return the input 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. * ``sharding_rule``: an SdyShardingRule object, an Einsum-like notation string that describes the sharding rule, or a Callable that produces either of these. We call the index labels in Einsum notation factors in our sharding rule. We borrow the idea from the einops.rearrange string , to use a space separator between factors and allow multiple letters factor names. By default, a factor corresponds to a passthrough/elementwise dimension. Factors corresponding to other dimensions can be specified via keyword arguments described below. See `jax-shardy-guide `_ for more details and examples. * ``reduction_factors``: A tuple of strings, specifying the reduction factors for a string `sharding_rule`. A reduction factor corresponds to a dimension that appears in operands but not in the result, such as the contracting dimensions in a matmul operation. If a reduction factor is sharded, the result would need to be all-reduced along the same axes. * ``need_replication_factors``: A tuple of strings, specifying the need_replication factors for a string `sharding_rule`. A need_replication factor corresponds to a dimension that shouldn't be sharded to support the implementation. * ``permutation_factors``: A tuple of strings, specifying the permutation factors for a string `sharding_rule`. A permutation factor corresponds to a dimension that would trigger collective permute if it is sharded. * ``factor_sizes``: A dictionary of variable keyword arguments, specifying the sizes of the factors that are only used in compound factors in a string `sharding_rule`. When config.use_shardy_partitioner.value is True, `sharding_rule` is used; otherwise, `propagate_user_sharding` and `infer_sharding_from_operands` are used. 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) # Use Einsum-like notation to specify the sharding rule. my_fft.def_partition( infer_sharding_from_operands=infer_sharding_from_operands, partition=partition, sharding_rule='...i -> ...i') # Use SdyShardingRule object to specify the sharding rule. my_fft.def_partition( infer_sharding_from_operands=infer_sharding_from_operands, partition=partition, sharding_rule=SdyShardingRule(operand_mappings=((BATCHING, 'i'),), result_mappings=((BATCHING, 'i'),)))) 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] rGcX||_d|_||_d|_d|_d|_yr,)r$r.static_argnumsr-r2r)r5r$rs r&r6zcustom_partitioning.__init__s0DHDN(D#'D (,D%Dr(r __getattr__Nreduction_factorsneed_replication_factorspermutation_factorsc T||_||_||_||_| t |t st |t rT| } t|dkDr|| d<t|dkDr|| d<t|dkDr|| d<| rtd| ||_ |St|f|||d| |_ |S)NrrrrzUnknown keyword arguments: r) r.r-r2rrkrrr9rlrr) r5r.r2r-rrrrr factor_sizessharding_rule_dicts r& def_partitionz!custom_partitioning.def_partitions DN#:D (DD%,DM8!D=/2'  ! #2C./ % & *9Q56 !A %4G01 67I6JKLL(d  4 -#;1   d r(c (t|j|tjd|ji|j}|jrt |jt fdtDttDcgc] }|vs| }}tjtj|j||d\}}t fd|jD}t|n%d}tj|j|}}tj|\} } tj || \} } | D cgc]} t#j$| }} t&j(j*j,}t#j.|j0j35t5j6| |\}}}dddtrJt#j8t5j:d}d}d}d}t<j>j@r<|jB#|jD |jF tId |jB}n|jD}|jF}tKjLg|| ||jN|||jP|| | |d }tjR| |Scc}wcc} w#1swYxYw) Nr)rc34K|]\}}|vr|n|ywr,rG)rirrs r&rz/custom_partitioning.__call__..s!Otq!^+12Os) debug_infoF)require_static_args_hashablec3(K|] }| ywr,rG)rrr#s r&rz/custom_partitioning.__call__.. s?a$q'?srGzShardy is used, but sharding propagation callbacks instead of sharding_rule are provided. Need to provide sharding_rule to migrate to Shardy.) rr.r-r2rrr0r1r4)*r'r$rrrsettuple enumerateranger9argnums_partiallu wrap_initrr tree_flattenflatten_fun_nokwargsr get_avalmesh_libthread_resourcesenv physical_meshr~rTrypetrace_to_jaxpr_dynamic ClosedJaxprconvert_constvars_jaxprr use_shardy_partitionervaluerr-r2NotImplementedErrorcustom_partitioning_prr.rtree_unflatten)r5r#r!debugr dyn_argnumsf_dyn_argsr4 args_flatr0flat_funr1rin_avalsrjaxpr_consts closed_callr-r2rout_flatrs ` @r&__call__zcustom_partitioning.__call__s 488T6 2D    5txx $b/3/B/B DE 4../n OyO Od %c$i 0L1A^4KQLkL-- ,,txxE 2  ', lb( ?4+>+>??k%k\\$((u=t(b"//9Iw!66r7CHh*34Q a 4H4  $ $ ( ( 6 6D !1!1!3 4G228XFeQG6{?""2#=#=e#DbIK"#' M $$**    $  ' ' 3  - - 9! !  ((m $ < <%)%F%F"$))     .. 7%A..# H  # #HJ 99gM5GGs K=&K=+LLL)rG)NNTN) rDrErF__doc__r6r forward_attrr__annotations__rrrGr(r&rrs@HT%11+s1BFCG"&=?-/R<;:r(c |jj} t| tjrft | j t | jjk(r1tjtj|d|g| Stjjj t| tj"rt| j$ t'd| j(} | tj*t-j.j1| j2| jn3t| tjr| jj4ndrt7dk(r1tjtj|d|g| Sd fd } t9||| ||||j| }t;t=|}|t>tA|d<|jjC||jDDcgc]}tjF|}}tIjJ|tM| tNjPjStTtNjVjSdtjXdtNjZjSgtNjPjS|dd }||j\Dcgc]}tjF|}}t_|rr|g| ||}t|tLt`frt7|dk(r|\}}ni}t|t:r tc|fi|}nt|tds tgd ti||||jjd <|jlScc}wcc}w) NT)multiple_resultsz9Please file a bug at https://github.com/jax-ml/jax/issuesc||SjssJtj|Stj|d}tjg|d|t |z z}tj |S)Nrr,)emptyr GSPMDShardingparse_flatten_op_sharding PartitionSpecr9 NamedSharding) hlo_shardingndimpspecrdevicesrs r&r/zB_custom_partitioning_lowering_rule..to_mesh_pspec_shardingYs  zz)    ) )'< @@  4 4d  E  ( ( R% R7dSZ>O3P RE  ' 'e 44r(utf8F)call_target_namehas_side_effect api_versioncalled_computationsbackend_configoperand_layoutsresult_layoutszgsharding_rule callable must produce either an SdyShardingRule object or an Einsum-like notation string.zsdy.sharding_rule)rzxc.HloSharding | None)7r3rurkrr}r manual_axesrrr lower_funr rrrrrShardingContextdevice_assignmentAssertionError abstract_meshMeshnparrayreshape axis_sizes_flat_devices_tupler9r*stridr_bytes add_keepaliverzaval_to_ir_typer CustomCallOprxr StringAttrget_CUSTOM_PARTITIONING_CALL_NAMEBoolAttri32_attr ArrayAttrrcallablerrrrlr attributesresults)ctxrr0r1r-r.r2rrr4valuesruamr/sharding_callback_infokeyr; result_typesout value_typesrrrs ` @@r&"_custom_partitioning_lowering_ruler:sI##00,~==> , " "#s<+<+<+G+G'HH I4>>$++D1D I# WPV WW  " " & & 4 4$ n<<=,,G  E GG  # #B ~ ]]288G,44R]]C==*d, > >?33GG CLA% 84>> $$ 88; F>D FF 511H'("C$6$6kK B% &'#,BeC()""#9:37>>Ba$&&q)B,B  6l}}(()GHkkooe,--",,**2.]]&&s+ #48MMBq4''*BKB #R[R$R R\Rm MD%= 1c-6HA6M,9) ) M3 '0UBTU -9%& &+D[,+0CNN&' 7CCs 5OO)+can_side_effecting_have_replicated_sharding)rrprop_user_shardingr.r2r)rzmlir.LoweringRuleContext)Kr __future__r functoolsrrtypingrcollections.abcrweakrefnumpyrjax._srcrrr r r r r rrrrrrrxb*jax._src.custom_partitioning_sharding_rulerrrjax._src.interpretersrrr jax._src.librrQjax._src.lib.mlirrjax._src.lib.mlir.dialectsrjax._src.shardingrr'r*WeakValueDictionaryr_rrLr8rVr]rirYrr Primitiverr%prim_requires_devices_during_loweringaddrrdef_abstract_evaldef_implrregister_custom_decorator_typerrregister_lowering register_custom_call_partitionerregister_plugin_callbacksrGr(r&r2s #$$&%#%{{&4) *&:2g113!9; 827 4In?&''<=)-&..223HI (''(JK89//n:n:0n:b N`,9;$##"0"504   ++ +G0%V48  r(