ukiddlmZddlZddlZddlmZmZddlZddlZddl Z ddl Z ddl m Z m Z mZddlmZddlmZddlmZddlmZdd lmZdd lmZdd lmZdd lmZdd lmZddlmZddlmZddlmZ ddl!m"Z"ddl#m$Z$m%Z%m&Z&m'Z'm(Z(m)Z)m*Z*m+Z+m,Z,m-Z-m.Z.m/Z/m0Z0m1Z1ddl2m3Z3m4Z4m5Z5ddl6m7Z7ddl8m9Z9m:Z:m;Z;ddlZidXd?Zjd@ZkdYdAZldBZmdZdCZnd[dDZoej d\ d]dEZp d^dFZqej d_ d`dGZr d\ddH dadIZsGdJdKZtdbdLZuejdcdMZwy)d) annotationsN)MappingSequence)Any NamedTuplecast)config)core)meshsharding)sharding_specs) tree_util)util)source_info_util) xla_bridge) mesh_utils) deprecations) xla_client)sdy)SdyArraySdyDimUnspecifiedValueAUTO flatten_spec NamedSharding_check_unique_resources UNSPECIFIED ArrayMappingArrayMappingOrAutoOrUnspecifiedget_array_mappingarray_mapping_to_axis_resources"named_sharding_to_xla_hlo_sharding"modify_sdy_sharding_wrt_axis_types)are_hlo_shardings_equalget_num_ways_dim_shardedis_hlo_sharding_replicated) PartitionSpec)safe_zip use_cpp_classuse_cpp_method.c`td|DsJttd|DS)Nc3ZK|]#}t|ts|jdu%ywN) isinstanceslicestep.0vs R/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/_src/sharding_impls.py zhashed_index..As! ?*Q*>QVVt^ ?s++c3pK|].}t|tr|j|jfn|0ywr.)r/r0startstopr2s r5r6zhashed_index..Bs+PAu)=QWWaff%1DPs46)allhashtuplexs r5 hashed_indexr?>s. ?Q ? ?? ? ePaPP QQiF)max_sizetrace_context_in_keyc |j}tj}i}||j D]'\}}t |}||}||xxdz cc<|||<)|S#t$rtd|ddwxYw)Nz,Cannot calculate replica ids from sharding: zf. Please create a device to index mapping for your sharding from which replica ids will be calculated.)devices_indices_mapAttributeError ValueError collectionsCounteritemsr?) r global_shapedevice_indices_map_fnindex_to_replicaoutdeviceindexh_index replica_ids r5device_replica_id_maprSEs-$88&1%8%8%: #,\:@@Bmfe5!G!'*JW"CK  * -  6xjA" " #)---s A&&Bc eZdZUded<ddZy) SdyArrayListzSequence[SdyArray] shardingsctjj|jDcgc]}|j c}Scc}wr.)rTensorShardingPerValueAttrgetrVbuild)selfr s r5rZzSdyArrayList.build]s7  ) ) - -*...9h 9 ;;9sAN)returnzsdy.TensorShardingPerValueAttr)__name__ __module__ __qualname____annotations__rZr@r5rUrUYs ;r@rUct||SN memory_kind)SingleDeviceSharding)rOres r5 _unpickle_single_device_shardingrges f+ >>r@ceZdZUdZded<ded<eddddZd Zd Zd Z d Z e dd Z e ddZ e ddZddZddZe ddZddZd dZe d!dZe d!dZd"dZy)#rfzA :class:`Sharding` that places its data on a single device. Args: device: A single :py:class:`Device`. Examples: >>> single_device_sharding = jax.sharding.SingleDeviceSharding( ... jax.devices()[0]) Device_device str | None _memory_kindNrdc ||_||_yr.)rjrl)r[rOres r5__init__zSingleDeviceSharding.__init__ysDL#Dr@c>t|j|jffSr.)rgrjrlr[s r5 __reduce__zSingleDeviceSharding.__reduce__~s ,t||T=N=N.O PPr@c`|jdnd|j}d|j|dS)N, memory_kind=zSingleDeviceSharding(device=))rlrjr[mems r5__repr__zSingleDeviceSharding.__repr__s:!!)"@Q@Q?R/SC )$,,)9#a @@r@c~t|ds&t|j|jf|_|jSN_hash)hasattrr;rjrer{rps r5__hash__zSingleDeviceSharding.__hash__s1 4 !t'7'789dj ::r@ct|tsy||ury|j|jk(xr|j|jk(SNFT)r/rfrjrer[others r5__eq__zSingleDeviceSharding.__eq__sH e1 2  u}  LLEMM ) 2    1 1 13r@c,t|jSr.len device_setrps r5 num_devicesz SingleDeviceSharding.num_devices t r@c|jhSr.rjrps r5rzSingleDeviceSharding.device_sets LL>r@c|jSr.rlrps r5rez SingleDeviceSharding.memory_kind   r@c0t|j|Src)rfrjr[kinds r5with_memory_kindz%SingleDeviceSharding.with_memory_kinds  $ ??r@cJ|jtdft|ziSr.)rjr0rr[rKs r5rEz(SingleDeviceSharding.devices_indices_maps! LL5;.3|+<< ==r@c|jfSr.rrps r5_device_assignmentz'SingleDeviceSharding._device_assignments LL?r@ctSr.replicated_hlo_shardingr[num_dimensionss r5_to_xla_hlo_shardingz)SingleDeviceSharding._to_xla_hlo_shardings ""r@clt|Dcgc]}tgd}}td|Scc}w)NF)axesis_open) mesh_shape dim_shardings)rangerr)r[r_sdy_dim_shardings r5_to_sdy_shardingz%SingleDeviceSharding._to_sdy_shardings?!&~!68B688 t3C DD8s1cy)NTrarps r5is_fully_replicatedz(SingleDeviceSharding.is_fully_replicateds r@ctj|jj|jjk(Sr.)xb process_indexrjclientrps r5is_fully_addressablez)SingleDeviceSharding.is_fully_addressables,  DLL// 0DLL4N4N NNr@cyr.rar[ aval_shapes r5check_compatible_avalz*SingleDeviceSharding.check_compatible_aval r@)rOrirerkr\intr\z set[Device]r\rk)rstrr\rfrKShaper\zMapping[Device, Index]r\XLADeviceAssignmentrrr\xc.HloShardingrrr\rr\boolrrr\None)r]r^r___doc__r`r+rnrqrxr}rpropertyrrrerrErrrrrrrar@r5rfrfis  /BF$$QA 3        @>  #E    O O r@rfz jax.shardingc|j|tj||j}t t |j j|Sr.) shard_shaperspec_to_indices sharding_specdictr)devicesflat)r[rKindicess r5!pmap_sharding_devices_indices_maprsF<  * *<9K9K L' ht||(('2 33r@ceZdZUded<ded<ded<e ddZdZd Zd Zd Z d Z dd Z e d ddZ ed dZej"d!dZd"dZej"d#dZed$dZd%dZd&dZd'dZej"d(dZej"d(dZd)dZd*dZy)+ PmapShardingz np.ndarrayrsharding_specs.ShardingSpecr xc.DeviceList_internal_device_listcFtj||_||_yr.)npasarrayrr)r[rrs r5rnzPmapSharding.__init__s::g&DL&Dr@cHt||j|jffSr.)typerrrps r5rqzPmapSharding.__reduce__s Jt'9'9: ;;r@ct|tsy||ury|j|jk(xrH|jj|jjk(xr|j |j k(Sr)r/rrrshaperrs r5rzPmapSharding.__eq__so e\ *  u}    %"5"5 5 F LL  %--"5"5 5 F  & &%*E*E EGr@c~t|ds&t|j|jf|_|jSrz)r|r;rrr{rps r5r}zPmapSharding.__hash__s3 4 !33T5G5GHIdj ::r@c |jjDcgc]}|j}}d|jd|d|jjdjj d|jj d Scc}w)NPmapSharding(sharding_spec=z , device_ids=z, device_platform=rz, device_shape=ru)rridrplatformupperr)r[d device_idss r5__str__zPmapSharding.__str__s $ 1 121!$$2J2)$*<*<)=>m#||003<<BBDEF LL../q 233sB c<d|jd|jdS)Nrz , devices=ru)rrrps r5rxzPmapSharding.__repr__s*)$*<*<)=>||nA '(r@c ||k(Sr.ra)r[rndims r5is_equivalent_tozPmapSharding.is_equivalent_tos 5=r@Nc |J| tdt|}|tj|t j |||dSt j t||}d}|jD]~}t|tjr|J|j}.t|tjsI|Jt|jdk(r|jd}utd|+tjtj d|}ntj|}|||S)aCreates a :class:`PmapSharding` which matches the default placement used by :func:`jax.pmap`. Args: shape: The shape of the input array. sharded_dim: Dimension the input array is sharded on. Defaults to 0. devices: Optional sequence of devices to use. If omitted, the implicit device order used by pmap is used, which is the order of :func:`jax.local_devices`. Nz*One of sharded_dim or devices must be set.rDrz3Multiple chunks in Chunked dimension not supported.)rGrrarrayrpmap_sharding_speccreate_pmap_sharding_specr<r r/ UnstackedsizeChunkedchunksNotImplementedErrorr local_devices) clsr sharded_dimrnreprnum_ways_shardeds pmap_devicess r5defaultzPmapSharding.defaults@ EFF \d '"  + +D$t DFF #<< e k#M  # # E A~// 0'''66 a// 0''' qxx=A XXa[ #CE E E!#"*:*:*<=N>N*O!PlXXg&l |] ++r@c,t|jSr.rrps r5rzPmapSharding.num_devicesrr@c@t|jjSr.)setrrrps r5rzPmapSharding.device_set#s t||  !!r@ct||Sr.)rrs r5rEz PmapSharding.devices_indices_map's ,T< @@r@c@t|jjSr.)r<rrrps r5rzPmapSharding._device_assignment*s "" ##r@c> |jjS#YyxYwr.)rdefault_memory_kindrps r5rezPmapSharding.memory_kind.s#  ' ' ; ;; sctd)Nzpmap does not support memories.rrs r5rzPmapSharding.with_memory_kind5s ? @@r@ctd)Nzpmap doesn't use OpSharding.rrs r5rz!PmapSharding._to_xla_hlo_sharding8s < ==r@ctd)Nzpmap doesn't use SdyArray.rrs r5rzPmapSharding._to_sdy_sharding;s : ;;r@c|jjD].}t|tjtj fs.yyr)rr r/rrr)r[rs r5rz PmapSharding.is_fully_replicated>s@    ( ( A00.2H2HI J r@c.|jjSr.rrrps r5rz!PmapSharding.is_fully_addressableE  % % : ::r@cyr.rars r5rz"PmapSharding.check_compatible_avalIrr@c (d}d}t|jjD]\}}t|tj r&|}|j }tj||}not|tjsa|}t|jdk(sJ|j|jd}tj||d}n||S|||k7r-td|d||d|dt|jS)NrDrzkThe sharded dimension must be equal to the number of devices passed to PmapSharding. Got sharded dimension z with value z in shape z and the number of devices=) enumeraterr r/rrrr tuple_deleterrr tuple_updaterGr)r[rKrsharded_dim_sizeir sharded_shapes r5rzPmapSharding.shard_shapeLs0K$,,556 1 A~// 0 66)), D  a// 0 188}!+188+!88A;)), QG    K $44  CCN-P$[12*\NK##&t'>'>#?"@ B CC r@)rzSequence[Device] | np.ndarrayrr)r[rrrrrr\r)rN)rrrz int | NonerSequence[xc.Device] | Noner\rrrrrr)rrrrrr)rKrr\r)r]r^r_r`r+rnrqrr}rrxr classmethodrrr functoolscached_propertyrrErrerrrrrrrrar@r5rrs; ,,&&'9'' <G 3( %) ;<48),1),=I),),V     ""A $$   A><   ;; r@rct|||Src) GSPMDSharding)r op_shardingres r5_unpickle_gspmd_shardingrhs w  EEr@czeZdZUded<ded<ded<ded<edd  dd Zd Zejd Z d Z dZ dZ ddZ eddZejddZeddZd dZed!dZd"dZd#dZejd$dZejd$dZedd d%dZy)&rr_devicesr _hlo_shardingrkrlrNrdct|tjr|ntjt||_t|tj rtj j|n||_||_ yr.) r/xc DeviceListr<r OpSharding HloSharding from_protorrl)r[rrres r5rnzGSPMDSharding.__init__rsd!+7BMM BW]]5>2 M( R]]C..33K@% $Dr@cpt|j|jj|jffSr.)rrrto_protorlrps r5rqzGSPMDSharding.__reduce__}s4 $ ]]D..7794;L;L M OOr@cb|jrttSt|jSr.)rr;rrrps r5_hlo_sharding_hashz GSPMDSharding._hlo_sharding_hashs(  ) ** "" ##r@ct|tsy||uryt|j|jxr4|j|jk(xr|j |j k(Sr)r/rr%rrerrs r5rzGSPMDSharding.__eq__si e] +  u}  #D$6$68K8K L J  E$5$55 J**e.I.IIKr@ct|ds1t|j|j|jf|_|j Srz)r|r;rr!rer{rps r5r}zGSPMDSharding.__hash__s@ 4 !33T5L5L((*+dj ::r@c`|jdnd|j}d|j|dS)NrsrtzGSPMDSharding(ru)rlrrvs r5rxzGSPMDSharding.__repr__s<!!)"@Q@Q?R/SC D..1#a 88r@c t|j\}}t|t|kr&td|dt|dt|y)Nz Sharding z+ is only valid for values of rank at least z%, but was applied to a value of rank )r&rrrG)r[rnum_ways_dim_shardedrs r5rz#GSPMDSharding.check_compatible_avalsd6t7I7IJ! :122  dVF % & ''L _    3r@c,t|jSr.)rrrps r5rzGSPMDSharding.num_devicess t)) **r@c,t|jSr.)rrrps r5rzGSPMDSharding.device_sets t}} r@c|jSr.rrps r5rezGSPMDSharding.memory_kindrr@cFt|j|j|Src)rrrrs r5rzGSPMDSharding.with_memory_kinds (:(: MMr@c,t|jSr.)r<rrps r5rz GSPMDSharding._device_assignments  r@c|jSr.)rrs r5rz"GSPMDSharding._to_xla_hlo_shardings   r@c|jjrtd|jd|jjr9t j dd}t |tj|S|jjr|jjstd|jdt|jj}tdtt|D}t j ||}t||j|Std|jd)NzCannot convert GSPMDSharding z into SdyArray.rac3&K|] }d| yw)_axis_Nrar3r s r5r6z1GSPMDSharding._to_sdy_sharding..sF!6!F)rtuple_elements TypeError is_replicatedmesh_lib AbstractMeshrr(ris_tiledis_tile_assignment_iotar<get_axis_sizesrr!_gspmd_to_named_sharding_via_mesh)r[r empty_mesh axis_sizes axis_namesr s r5rzGSPMDSharding._to_sdy_shardingsD ((*  )$*<*<)=_ M OO   ) ) +((R0j :} 7 H H    $ $ &    7 7 9+D,>,>+? OQ Q++::<=jFuS_/EFFj  " ":z :d .tT : K K   )$*<*<)=_ M OOr@c,t|jSr.)r'rrps r5rz!GSPMDSharding.is_fully_replicateds %d&8&8 99r@c.|jjSr.rrps r5rz"GSPMDSharding.is_fully_addressablerr@c ||t|Srcr)rdevice_assignmentres r5get_replicatedzGSPMDSharding.get_replicateds  "9& ((r@)rz Sequence[Device] | xc.DeviceListrxc.OpSharding | xc.HloShardingrerkrrrr)rrr\rrrrr)rerk)r]r^r_r`r+rnrqrrr!rr}rxrrrrrerrrrrrrrBrar@r5rrks+ &&-1$:$)$$O $$ K 9  + +   N    O* :: ;;JN((r@rc&tj|d\}}|d}g}|D]T}t|ttfs||j |.t|t jrt|trtd|d|dt|tr(|jjrtd|d|d|s=t|tr-tj|jvrtd|d ||j |t|tst!|d ||s#tj|vrtd|d |t#|||j |Wtj$||S) Nc |duSr.rar=s r5z(prepare_axis_resources..s T r@)is_leafz leaf specificationszOne of z got sharding z which is not allowed.z got an empty NamedSharding: z2Unconstrained dims are not allowed when passed to z: z= are expected to be PartitionSpec instances or None, but got )r tree_flattenr/rrappend jshardingShardingrrGrr emptyr( UNCONSTRAINEDspecr3rtree_unflatten)axis_resourcesarg_nameallow_unconstrained_dimsentriestreedefwhat new_entriesentrys r5prepare_axis_resourcesrXs++13'7 :) *$+ e%*D12em E9-- . E< (74&ug>$$% % E= )ejj.>.>74&(EeWM112 2&:e]+K  % % 3@ Kw   } -4&!DDI7LM M %-*E*E*N@ Kw  eX.3 6  ! !'; 77r@c0eZdZUdZded<ded<ded<y) AxisEnvz5Represents a pmap mesh (only along the replica axes).rnrepsztuple[Any, ...]namesztuple[int, ...]sizesNr]r^r_rr`rar@r5rZrZs= *  r@rZT)frozenc\eZdZUdZded<eZded<edZedZ d dZ y ) SPMDAxisContextzA hardware axis context for parallel computations that use the GSPMD partitioner. This includes the mesh that will later by used to execute this computation, as well as a set of mesh axes that are currently lowered in the MANUAL sharding mode. mesh_lib.Meshr frozenset[MeshAxisName] manual_axesc|jSr.)unsafe_axis_envrps r5axis_envzSPMDAxisContext.axis_envs   r@ct|jj|jjt |jj j S)Nr[r\r])rZr rr=r<rvaluesrps r5rfzSPMDAxisContext.unsafe_axis_envs@ iinnii""DIIOO**,- //r@cHt|j|j|zSr.)rar rd)r[rs r5 extend_manualzSPMDAxisContext.extend_manual$s 499d&6&6&= >>r@N)rrcr\ra) r]r^r_rr` frozensetrdrrgrfrlrar@r5rara sG )2+&4     / / ?r@raceZdZUdZded<y)ReplicaAxisContextzA hardware axis context for parallel computations that are partitioned by JAX. Unlike in the SPMDAxisContext, this means that JAX might need to emit calls to explicit collectives. rZrgNr^rar@r5roro(s r@rocNeZdZUdZded<dZded<dZded<d Zed Z y) ShardingContextzA hardware axis context for parallel computations that use the sharding interface. This context also uses the GSPMD partitioner. rrNztuple[xc.Device, ...] | NonerAzmesh_lib.AbstractMesh | None abstract_meshc|jAt|jtsJ|jt |jk(sJyyr.)rAr/r<rrrps r5 __post_init__zShardingContext.__post_init__=sH ) .. 66 6   T%;%;!< << <*r@ctdddS)NrDrari)rZrps r5rgzShardingContext.axis_envCs "B //r@) r]r^r_rr`rArrrtrrgrar@r5rqrq2s? 481804--4=  0 0r@rqcptj|ddddddtj|zS)z5Returns an array of strides for major-to-minor sizes.N)rcumprodr)r]s r5strides_for_sizesryMs0 E$B$K 2 &"**U*; ;;r@c|jDcic]\}}|dk7s ||}}}tj|jtj}t |}t |t|}t|||Dcic] \}}}||f| }}}}|D cgc]} ||  c} Scc}}wcc}}}wcc} w)aiRecovers the ordering of axis names based on a device assignment. The device assignments that this function can convert into axis orders are of the form:: np.arange(np.prod(named_sizes.values())).transpose(...).flatten() for some transposition ``...``. This is satisfied by all OpSharding assignments generated from partition specs. Arguments: named_sizes: A dictionary mapping axis names to their sizes. assignment: A permutation of integers between 0 and the product of all named sizes. Returns: A major-to-minor list of axis names that corresponds to the given assignment. rDdtype) rJrfromiterrjint64ryexplode_superdimsunflatten_superdimszip) named_sizes assignmentnamerr]stridesdimsstride dim_to_namers r5unflatten_arrayrQs&/:.?.?.AO dTQYtO+O ++k((*"(( ;% e $' 5"5j"A B$@CE7T_@`aa*<$$%a+a"& 'Q+a. '' Pb 's B5B5B;& Cc\d}tj|tj}||ddk(g}|jdkDrd|d}t t |D]}||||zk7sndz }|}|j ||f|dkDsJ|dd|}|jdkDrd|S)aUnflatten a list of dimension sizes and their strides that generates assignment. If this function succeeds for a given ``assignment``, then the following property should be satisfied:: dims_with_strides = unflatten_superdims(assignment) base_array = np.arange(map(fst, sorted(dims_with_strides, key=snd, reverse=True))) assignment == base_array.transpose(argsort(dims_with_strides, key=snd, reverse=True)).flatten() That is, the returned dimensions list all sizes of the base array (with strides indicating their initial order). The order of dimensions in the list corresponds to the permutation that applied to the base array generates the assignment. c|rytd)NzKFailed to convert OpSharding into a ShardingSpec. Please open a bug report!r)conds r5checkz"unflatten_superdims..checkys V : ;;r@r{rrDN)rrr~rrrrI)rrflat_assignmentrrr rs r5rrks;JJz:/a $q Q F 3' (   q6z )5 1fa DKKv !8O8%ff-O q  +r@c|t|t|Dcic]\}}|| }}}tt|}g}|D]r\}}||}g}||kDr9|dkDsJ||zdk(sJ|j ||f||z}||z}||}||kDr9||k(sJ|j ||f|t|z }t|Scc}}w)aExplode superdims to fit a known shape. The unflattening process might mistakenly generate too few too large dimensions. For example, ``unflatten_superdims(np.arange(n))`` always returns ``[(n, 1)]``. This function takes a list of such contiguous super-dimensions and splits them into smaller dimensions such that:: set(map(fst, explode_superdims(sizes, dims))) == set(sizes) rDr)rrylistreversedrI)r]rrrstrides_to_sizes final_dims target_sizenew_dimss r5rrs8;5BSTYBZ7[\|tVfdl\\ htn $* %ldF"6*KH   1__ K 1 $$ $oo{F+, {d f$V,k   ;   OOT6N#(8$$J % !]s B8c t|tjrtjj |}|j r4g}|j D]}|j t|||S|jr tgS|jr|jdk(r tgS|jr+|j}t|j|j}t!|}|j#}g}|D]o} g} | dkDrLt%|} || } | | zdk7rt'd|d|d| d| d | | z} | j)| | dkDrL|j)t+| qt-|j/dkDr t1d|j3r|dd }|r#|d d k(r|j5|r |d d k(rt|gSt7d ) NrDrzshape=z! is incompatible with mesh_shape=z : dim_size=z is not divisible by axis_size=.z5Unhandled HloSharding type. Please open a bug report!rwraz4Unhandled OpSharding type. Please open a bug report!)r/rrrrr2extendparse_flatten_op_shardingr4r( is_maximalrr7rrtile_assignment_devicesitertile_assignment_dimensionsnextrGrIr<rsubgroup_typesrreplicate_on_last_tile_dimpopAssertionError) hlo_shardingr rNrrmesh_axis_order mesh_axisr partitionsdim_sizedim_partitionsaxis axis_sizes r5rrs bmm,>>,,\:L  "!C  ( ( *5 jj*1d345 J!!# O  TYY!^ O J% L88:O_%I  3 3 5EJ /n qLIt$ i 1 $:zm< ; A?   Yd# qLn-. / < & & ()A-  A ..0cr?j B2-nn B2- : & '' O PPr@cN|jJ|j|jfSr.)r1r8r9)rs r5_slice_as_tuplers#  ''166 r@ceZdZdZy)NonUniformShardingErrorz5Raised when sharding is not uniform across processes.N)r]r^r_rrar@r5rrs=r@rc|js |jry|j|jf|z}|j Dcic] \}}||| }}}t j t}t}|j D]L\} }|j|jf} || jj| |j| Ntt|jj} t|| } t!| t#fd|j%Drt'd|d|dt)|j%D chc] } t| c} }t+d|Dt!|k7rt'd|d||j-| t!|fScc}}wcc} w)a7 Get current process index and number of unique processes for given dimension. This function facilitates mapping of process-level data to individual devices. Each process can use its index to obtain the data corresponding to that index. If process level data is sharded on multiple dimensions this function can be used to build the cross product of indices in each sharded axis. Processes that need to load the same data will have the same index. For shardings whose per-process data is not distributed on a grid, the number of distinct shards will be such that it is possible to build the target shape while maintaining a "cube" shape of local-process data. For example, in case of 4 hosts with sharding distributed like so: 1234 2143 For dim 0 (rows): all processes need to access all rows, so we return (0, 1) For dim 1 (cols): process 1 and 2 returns index 0 out of 2 (need cols 0 and 1), process 3 and 4 returns index 1 out of 2 (need cols 2 and 3). On the other hand, for a sharding like: 1212 3434 Dim 0 (rows): process 1 and 2 returns (0, 2), process 3 and 4 returns (1, 2) Dim 1 (cols): process 1 and 3 returns (0, 2), process 2 and 4 returns (1, 2) Note: This function requires sharding to be process uniform in dimension `dim`: each process has the same number of addressable indices in that dimension and all index sets across processes are either disjoint or the same. For sharding to be process uniform the addressable shards doesn't need to form contiguous subtensor, or even a sparse grid and in case of interleaved high-dimensional tensor it is possible for sharding to be process uniform only in some dimensions but not others. For example: 1111 and 12 and 1212 and 1212 2222 21 2121 1212 are all sharding uniform, in both dimensions. However 1122 2121 1121 1222 is uniform in dimension 0 (both hosts access all rows), but is not uniform in dimension 1 (host 1 accesses columns: 0, 1, and 3), while host 2 accesses (0, 1, 2, 3). Returns: A tuple of (index, num_distinct_shards) for the given dimension. It is guaranteed that `index` will cover 0 to `num_distinct_shards - 1`, across all processes. Raises: NonUniformShardingError: if the sharding is not process uniform in dimension `dim`. )rrDc3:K|]}t|k7ywr.r)r3r>slices_per_processs r5r6z.get_process_index_and_count..FsI!Q% %Isztensor_sharding=z is non-uniform on dim=z3 as some processes have different number of slices.c32K|]}t|ywr.r)r3hs r5r6z.get_process_index_and_count..Ps*AQ*s)rrrErrJrH defaultdictrr8r9raddrraddressable_devicesrmranyrjrrsumrP)tensor_shardingdimndims device_mapkr4 global_sliceprocess_to_slice all_slicesrkey current_pidaddressable_slicesr>unique_processesrs @r5get_process_index_and_countrsH**))  22""$u,.*)3(8(8(:;1!QsV);,;!,,S1u*  "da 77AFF CQ__%))#.NN3T/==>?MM+ !1+!>?-.I/?/F/F/HII ! ? 4v6& & 1A1H1H1JKA9Q<KL  *)**c*o= ! ? 4v6   !3 4c:J6K LLE<2Ls G.Gcdgt|z}t|D](\}} t||t|\}}||z||<*t |S#t$rd||<YFwxYw)aVComputes the global shape given the per process if possible. The returned shape will have the size of the global tensor in that dimension or None, if it is not computable. The latter can happen when sharding is not uniform along that dimension, e.g. different hosts require different shapes, or if different processes have partial data overlap. If at most one dimension is sharded the shape is always computable. Generally, global shape is computable for most practical meshes (including topology aware such as meshes returned by mesh_utils.create_device_mesh) Some examples: Suppose mesh is {'a': 2, 'b': 2, 'c': 2} with 2 devices per host, 4 hosts total. For different specs we get: - P(): global_shape = local_shape - P(('a', 'b', 'c'), None): global_shape = (4 * local_shape[0], local_shape[1]) Note: per device shape is (local_shape[0] / 2, local_shape[1]) - P(('a', 'b'), None) global_shape = (4 * local_shape[0], local_shape[1]) # NB: the same global shape as above, since sharding along 'c' dimension # happens to be within process, and thus doesn't affect the global shape. # The underlying difference will be in the per *device* shape, which # would be (local_shape[0], local_shape[1]) in this case. - P(None, ('a', 'c')) global_shape = (local_shape[0], 2 * local_shape[1]) # Per device shape is (local_shape[0], local_shape[1] / 2) - P(('a', 'c'), 'b'): global_shape = (2 * local_shape[0], 2 * local_shape[1]) # Per device shape is (local_shape[0] / 2, local_shape[1]) - If devices in the Mesh are randomly permuted: For any partition spec which shards more than 1 axis: e.g. P('a', ('b', 'c')): global_shape = (None, None) Args: local_shape: global shape of the tensor. Returns: global_shape with Nones in non-uniform dimensions. N)r)rrrrr<)r local_shaperKr  local_dimr shard_counts r5local_to_global_shaperWs\&*FS-=$=, ,la2 AS-/na!K/l1o  |  #l1os"AA"!A"c |j|}tttjt f|}t |jDchc]}t||c}}|j||}||zScc}w)aReturns the number of indices for given dimension this host has access to. Each host can have multiple number of devices that are spanning possibly discontiguous slices of data. This function computes the total number of unique indices for dimension `dim` that any of its addressable devices hold. In most cases the addressable indices form a sparse grid (and in some cases a subcube), and thus each host will hold the same of number of indices for each dimension. However, it is possible to design a mesh that addressable shards form a complicated pattern. In that case, the returned value is the number of indices that are addressable by at least one device. For example, suppose the sharding looks like this: (number indicates the host index) 1221 1221 0000 Then on host 1 and 2, both dim 0 (rows), and dim=1 (cols) will have size 2, while on host 0, dim 0 will have size 1, and dim 1 will have size 4. Args: tensor_sharding: Sharding of the tensor. dim: dimension along which to compute the number of addressable indices. global_shape: global shape of the tensor. Returns: The number of indices for dimension `dim` that this host holds. ) addressable_devices_indices_maprrrJriIndexrrjrr)rrrK addressables addressablenum_unique_slices shard_sizes r5num_addressable_indicesrsD!@@N,gi..56 E,;G;N;N;P,7ok#&'**<8=* ' '' sBc$tj|j}|jj }t |\}}|dk(rgn|g}|dg|j zz|z}||_tjj|SNrD) r physical_element_avalr|rcloner&rrrrr)avalrelt_avalnew_op_shardingr num_replicassuffixtads r5physical_hlo_shardingrs  ' ' 3( ))+113/5lC*l"2&aS8==((61#/2/,  " "? 33r@cF|jdk(xrt|t Sr)rr/rr s r5is_single_device_shardingrs$    " M:h +M'MMr@ct|r|St|trtj|j }t jg|jz}t jg|jj||jj}t|j|St|trVtj|j }dg|jz}|jt!g|j"|S|j%|j}t'|j(t+||SN)r mesh_mapping)rr)rN)rr/rr rr|r NoShardingr ShardingSpecrr rrrupdater(rNrrrr)rr rtrailing_shardingphys_sharding_spec trailing_spechloss r5make_key_array_phys_shardingrs"x( O(L)))$**5H'2245 E'44G8))22G5FG++88:  0 0&8 ::(M*))$**5HFX]]*M ?? Mx}} M} M? NN  ( ( 3D &&(=dD(I KKr@ct||Sr.)r)rr s r5physical_shardingrs %dH 55r@ctj|}|jt||jz}t |\}}|dk(rgn|g}|j j}|d|j |z} | |_t|jtjj|Sr)r rrrrr&rrrrrrrr) logical_shaper| phys_shardingrphys_hlo_shardingrrrlogical_op_shardingrs r5get_logical_gspmd_shardingrs  ' ' .(#88 -8==(*56GH*l"2&)224::<?X]]N#f,#360 }::~~001DE GGr@crt|tryt|tr|jjryt j ||}|jt|}t|\}}t|t|z }td|| dDstd|d|d|y)Nc3&K|] }|dk( yw)rDNrar0s r5r6z1check_replicated_trailing_dims..s =Q!V =r1zIThe trailing dims of extended dtypes should be replicated. Got sharding: z, partitions: z, num_trailing_dims: ) r/rrr _any_axis_manualr physical_shaperrr&r:r)r rr| phys_shapehlo_srrnum_trailing_dimss r5check_replicated_trailing_dimsrs,' -(X]]-K-K ""=%8*  ' 'J 8%*51-*a*oM(:: =Z):(:(;< = =  Z~j\:/0 2 33 >r@ct|||t|r|St|tryt j |}t j|jjd|j |jj}t|j|St|trt j |}t j||}t|j t|kr5g|j dgt|t|j z z}n |j }|j#|d|j St%|||Sr)rrr/rr rrrrr rrrrrrrNrr)rr|rrlogical_sharding_specr phys_specs r5logical_shardingrsH! uE}- -.))%0H*77,,55o ~F"00==?  5 5&; ==-/))%0H$$]E:J =  Z0JM&&JFc*oM4F4F0GGHJi $$i   Y%?  @@ %mUM JJr@ct|||Src)r)r pspecres r5cached_named_shardingrs tU < N     ' ' )(-(~~  ..6Z 9 ::Xx0H  ##55""$x~~ ##x/  ($mm)) *+hZ@(()9)A)A)CD E G HH(--/x}}::HMMJh  & Hay}}""1% 1 1 8 8*::  $XJ/mm_-aS1mm))!, -.(()9)A)A)CD E  G HH H /r@)rc|tj}tj|}|t d|~t j |}|t|kDrt dt|d||t|kr|d|}|djtjtjfvrd}nd}tj|||}|jd}|jd k(rGt|d r;t|jD chc]} | jc} d kDr t d |ct!j"d r2t$j&j(ft|j*z}nt-j.dt0dt%j2|||Scc} w)a Creates an efficient mesh with the shape and axis names specified. This function attempts to automatically compute a good mapping from a set of logical axes to a physical mesh. For example, on a TPU v3 with 8 devices: >>> mesh = jax.make_mesh((8,), ('x')) # doctest: +SKIP >>> [d.id for d in mesh.devices.flat] # doctest: +SKIP [0, 1, 2, 3, 6, 7, 4, 5] The above ordering takes into account the physical topology of TPU v3. It orders the devices into a ring, which yields efficient all-reduces on a TPU v3. Now, let's see another example with 16 devices of TPU v3: >>> mesh = jax.make_mesh((2, 8), ('x', 'y')) # doctest: +SKIP >>> [d.id for d in mesh.devices.flat] # doctest: +SKIP [0, 1, 2, 3, 6, 7, 4, 5, 8, 9, 10, 11, 14, 15, 12, 13] >>> mesh = jax.make_mesh((4, 4), ('x', 'y')) # doctest: +SKIP >>> [d.id for d in mesh.devices.flat] # doctest: +SKIP [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15] As you can see, logical axes (`axis_shapes`) affect the ordering of the devices. You can use `jax.experimental.mesh_utils.create_device_mesh` if you want to use the extra arguments it provides like `contiguous_submeshes` and `allow_split_physical_axes`. Args: axis_shapes: Shape of the mesh. For example, axis_shape=(4, 2) axis_names: Names of the mesh axes. For example, axis_names=('x', 'y') axis_types: Optional tuple of :class:`jax.sharding.AxisType` entries corresponding to the ``axis_names``. See `Explicit Sharding`_ for more information. devices: Optional keyword only argument, that allows you to specify the devices you want to create a mesh with. Returns: A :class:`jax.sharding.Mesh` object. .. _Explicit Sharding: https://docs.jax.dev/en/latest/notebooks/explicit-sharding.html NzF`axis_shapes` passed to `make_mesh` should be a sequence of ints. Got zNumber of devices z& must be >= the product of mesh_shape rTF)allow_split_physical_axestpu slice_indexrDzy`jax.make_mesh` does not support multi-slice topologies. Please use jax.experimental.mesh_utils.create_hybrid_device_meshzjax-make-mesh-default-explicitaThe default axis_types will change in JAX v0.9.0 to jax.sharding.AxisType.Explicit. To maintain the old behavior, pass `axis_types=(jax.sharding.AxisType.Auto,) * len(axis_names)`. To opt-into the new behavior, pass `axis_types=(jax.sharding.AxisType.Explicit,) * len(axis_names))category stacklevel) axis_types)rrr_canonicalize_axis_sizesrGmathprodr device_kind _TPU_V5_LITE_TPU_V5Ecreate_device_meshrrr|rris_acceleratedr5r ExplicitrwarningswarnDeprecationWarningr ) axis_shapesr=rrnew_axis_shapesrr mesh_devicesfirst_drs r5 make_meshr+bs\ _jjlG77 D/  }  ii()W  S\N+() + ,,3w<jy!G QZ 7 79L9LMM $ %..w 9;,   a ' %GG]$C ,"3"3 4Q1== 459  A BB""#CD%%..03|7I7I3JJjmm M & |ZJ GG# 5sF;c,eZdZdZddgZddZdZdZy) set_meshaSets a concrete mesh in a thread-local context. ``jax.set_mesh`` has dual behavior. You can use it as a global setter or as a context manager. When a mesh is in context via ``jax.set_mesh``, you can use pass raw PartitionSpecs to all APIs that accept sharding as an argument. Using ``jax.set_mesh`` is also required for enabling explicit sharding mode: https://docs.jax.dev/en/latest/notebooks/explicit-sharding.html For example:: mesh = jax.make_mesh((2,), ('x',)) jax.set_mesh(mesh) # use the API as a global setter with jax.set_mesh(mesh): # use the API as a context manager ... Note: ``jax.set_mesh`` can only be used outside of ``jax.jit``. prev_abstract_mesh prev_meshc~t|tjstdt |t j s td|jrtd|dtjj|j|_ tjj||_y)Nz/Expected mesh of type `jax.sharding.Mesh`. Got z1`set_mesh` can only be used outside of `jax.jit`.zmesh z contains manual axes which is not allowed when using `jax.set_mesh`. Please use `jax.shard_map` to enter into `Manual` mode instead.)r/r5r rGrr trace_state_cleanrr abstract_mesh_context_manager swap_localrrr.device_contextr/)r[r s r5rnzset_mesh.__init__s dHMM *  ;DJ< H JJ  ! ! # J KK   $   %BBMM D**55d;DNr@cyr.rarps r5 __enter__zset_mesh.__enter__sr@ctjj|jtjj|j yr.)r r2 set_localr.r4r/)r[exc_type exc_value tracebacks r5__exit__zset_mesh.__exit__s4 ((2243J3JK ##DNN3r@Nr rb)r]r^r_r __slots__rnr6r<rar@r5r-r-s"($[1)<  4r@r-chtjs tdtjS)Nzd`get_mesh` can only be used outside of `jax.jit`. Maybe you want `jax.sharding.get_abstract_mesh()`?)r r1rGr5get_concrete_meshrar@r5get_meshrAs1   !  / 00  # # %%r@c#Kt|tjsJtjj |} dtjj |y#tjj |wxYwwr.)r/r5r r r4r3r8)r prev_vals r5_internal_use_concrete_meshrDs` D(-- (( (  " " - -d 3(.  ##H-F##H-sr\s#- ((%*#%%!)*((((   2AA WW^^  c3h  eSjFCK(!**R T6 7 & ;;;..224?r&&'O 9--O (O b#1 T644"8474r]9%%] ]~) Fr l(I&&l(!l(^ 5:"8Njd#??$?6d#$d#00$04<(4!F8-Q0-Q /-Q4K-Q` >j> T6pM'pM.1pM:=pMBQpM7pMf8 8/489O8v(('((.1((AF((KN((V4N K*6 G3 K6 #= /=8E=='4== > > E>> HL4$'4AE4344pBFXH7;XH>XH4XH@MXHt,4,4^& ..r@