
    Vpf}                       d dl mZ d dl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
Z
d dlZd dlmZmZ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m Z m!Z! d dl"m#Z# d dl$m%Z%m&Z&m'Z'm(Z( d dl)Z*e+e,df         Z-ej.        Z.e+e/df         Z0e+e.df         Z1ej2        Z3 ej4        d           G d d                      Z5 ej6        dd          d             Z7ddZ8 ej6        dd          dd"            Z9 ej6        dd          dd%            Z: e'ej;                   G d& d'ej2                              Z; ej6        dd          d(             Z< e'ej=                   G d) d*ej2                              Z= ej6        dd          dd,            Z> e'ej?                   G d- d.ej2                              Z?	 ddd6Z@ ej6        dd          dd7            ZA G d8 d5ej2                  ZB G d9 d:          ZC e'ejD                   G d; d<ej2                              ZD G d= d>          ZEd? ZF G d@ dA          ZG eG            ZHdB ZIdC ZJeZK	 eeKe,f         ZLeeLeEeGf         ZMddFZNddJZOdK ZP G dL dMejQ                  ZR G dN dO          ZS G dP dQeS          ZT eU            fdRZVedSk     reVZW	 ddTZXdU ZY G dV dWe          ZZ ej4        d           G dX dY                      Z[ ej4        d           G dZ d[                      Z\ ej4        d           G d\ d]                      Z]d^ Z^d_ Z_d` Z`da ZaddfZbddiZc G dj dked          ZeddqZfdduZgddvZhddwZiddyZjdz Zkdd{Zld| Zmdd}Zndd~Zo ej6                    	 	 ddd            ZpddZqdS )    )annotationsN)OrderedDict)MappingSequence)Any
NamedTupleUnioncast)coremeshsharding)sharding_specs)	tree_util)util)
xla_bridge)
xla_client)xla_extension_version)are_op_shardings_equalget_num_ways_dim_shardedis_op_sharding_replicated)PartitionSpec)safe_mapsafe_zipuse_cpp_classuse_cpp_method.T)frozenc                      e Zd ZU ded<   dS )TransferToMemoryKindstrmemory_kindN)__name__
__module____qualname____annotations__     W/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/_src/sharding_impls.pyr    r    2   s         r(   r       F)max_sizetrace_context_in_keyc                   	 |D ]F}|B|D ]?}| j         |          ||v r,t          d| d|                                 d| d          d @Gd S # t          $ r,}t          d|j        d          d|j         d          d d }~ww xY w)NzAxis: z of z is also found in manual_axes: .zResource axis: r   z is undefined.)shape
ValueErrorget_partition_specKeyErrorargs	user_spec)r   parsed_pspec_manual_axespres         r)   _check_mesh_resource_axisr:   7   s   - L L	
 	L 	LA
*Q--,A A A ? ? A A A A1=A A AB BGKL 	L L 
 - - -
 "qvay " "l6L " " " # #(,--s   A	A 
B'A>>Breturnintc                    t          d | D                       sJ t          t          d | D                                 S )Nc              3  R   K   | ]"}t          |t                    |j        d u V  #d S N)
isinstanceslicestep.0vs     r)   	<genexpr>zhashed_index.<locals>.<genexpr>J   s7      ??*Q*>*>?QVt^??????r(   c              3  `   K   | ])}t          |t                    r|j        |j        fn|V  *d S r?   )r@   rA   startstoprC   s     r)   rF   zhashed_index.<locals>.<genexpr>K   s?      PPAu)=)=DQWaf%%1PPPPPPr(   )allhashtuplexs    r)   hashed_indexrO   G   sL     
??Q???	?	????	ePPaPPPPP	Q	QQr(   i   global_shapeShapeMapping[Device, int]c                $   	 | j         }n"# t          $ r t          d|  d          d w xY wt          j                    }i } ||                                          D ]1\  }}t          |          }||         }||xx         dz  cc<   |||<   2|S )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_mapAttributeErrorr0   collectionsCounteritemsrO   )	r   rP   device_indices_map_fnindex_to_replicaoutdeviceindexh_index
replica_ids	            r)   device_replica_id_mapra   N   s    -$8	 - - -
	"x 	" 	" 	"# # )--- &1%8%:%:
#,,\::@@BB  mfe5!!G!'*JW"CKK	*s   
 )num_dimensionsxc.HloShardingc                   | j         j        }t          | j                  }d t	          | j         j                  D             }i }| j        rB| j         j        }| j        D ].}t          j        j	        j
        ||                    |          <   /g }t	          |                                          D ]#\  }	\  }
}|
|vr|                    |	|f           $t          |          t          |          k    r |st          j                                        S g }dg|z  }t#          |                                d           D ]6\  }}||xx         ||         z  cc<   |                    ||                    7g }|r:t%          j        t(                    }t%          j        d           }d |D                                 t-          |                                                    sJ |D ]Z\  }	}|                    |	t          j        j	        j                  }||                             |	           ||xx         |z  cc<   [t#          |                                d           D ]J\  }}|                    |           |                    ||                    |                    |           Kt          j                            |t9          | j         j                                                  ||          S )	Nc                    i | ]\  }}||	S r'   r'   rD   inames      r)   
<dictcomp>z6named_sharding_to_xla_hlo_sharding.<locals>.<dictcomp>g   s    JJJwq$4JJJr(   rT   c                    | d         S NrT   r'   rM   s    r)   <lambda>z4named_sharding_to_xla_hlo_sharding.<locals>.<lambda>y   s
    qt r(   )keyc                     dS rk   r'   r'   r(   r)   rl   z4named_sharding_to_xla_hlo_sharding.<locals>.<lambda>   s    1 r(   c                    h | ]
}|d          S )r   r'   rD   rN   s     r)   	<setcomp>z5named_sharding_to_xla_hlo_sharding.<locals>.<setcomp>   s    ///QAaD///r(   c                    | d         j         S Nr   )valuerM   s    r)   rl   z4named_sharding_to_xla_hlo_sharding.<locals>.<lambda>   s    qtz r(   )dimsreshape_dimstranspose_permsubgroup_types)r   r/   get_array_mapping_parsed_pspec	enumerate
axis_namesr6   xc
OpShardingTypeMANUALr^   rY   appendlenHloSharding	replicatesortedrW   defaultdictlist
issupersetsetkeysget
REPLICATEDextend	iota_tilerL   values)selfrb   
mesh_shapearray_mappingmesh_axis_posspecial_axesr|   manual_axisreplicated_mesh_axesrg   	axis_nameaxis_valmesh_permutationnew_mesh_shaperh   poslast_tile_dimsaxes_by_typesize_by_typesizetyaxess                         r)   "named_sharding_to_xla_hlo_shardingr   b   s$    y*#D$677-JJ)DI4H*I*IJJJ-,	 N%J( N N46M4F4Ml:##K0011"+J,<,<,>,>"?"? 1 1a	)X%%!!1h-000	#j//11,1>##%%%3'.---//^^DDD 1 1idC3:d++M$/0000. $*400L*9955L//.///::3|?P?P?R?R;S;STTTTT'  4Ar}1<==b2a   2$<--//5I5IJJJ $ $DBL,---d####$ 
	!	!dio.D.D.F.F(G(G%n 
" 
F 
F Fr(   c                     e Zd ZU dZded<   ded<   ded<   ded	<   d
ed<    e            dd e            dd-d            Zd Zd Z	e
d.d            Zd Zd Zd/dZed e            dd            Ze
d0d            Ze
d1d            Ze
d2d!            Ze
d0d"            Zej        d2d#            Zd3d&Zd4d*Zd5d,ZdS )6NamedShardinga  A :class:`NamedSharding` expresses sharding using named axes.

  A :class:`NamedSharding` is a pair of a :class:`Mesh` of devices and
  :class:`PartitionSpec` which describes how to shard an array across that
  mesh.

  A :class:`Mesh` is a multidimensional NumPy array of JAX devices,
  where each axis of the mesh has a name, e.g. ``'x'`` or ``'y'``.

  A :class:`PartitionSpec` is a tuple, whose elements can be a ``None``,
  a mesh axis, or a tuple of mesh axes. Each element describes how an input
  dimension is partitioned across zero or more mesh dimensions. For example,
  ``PartitionSpec('x', 'y')`` says that the first dimension of data
  is sharded across ``x`` axis of the mesh, and the second dimension is sharded
  across ``y`` axis of the mesh.

  The Distributed arrays and automatic parallelization
  (https://jax.readthedocs.io/en/latest/notebooks/Distributed_arrays_and_automatic_parallelization.html#namedsharding-gives-a-way-to-express-shardings-with-names)
  tutorial has more details and diagrams that explain how
  :class:`Mesh` and :class:`PartitionSpec` are used.

  Args:
    mesh: A :class:`jax.sharding.Mesh` object.
    spec: A :class:`jax.sharding.PartitionSpec` object.

  Examples:

    >>> from jax.sharding import Mesh
    >>> from jax.sharding import PartitionSpec as P
    >>> mesh = Mesh(np.array(jax.devices()).reshape(2, 4), ('x', 'y'))
    >>> spec = P('x', 'y')
    >>> named_sharding = jax.sharding.NamedSharding(mesh, spec)
  mesh_lib.Meshr   r   spec
str | None_memory_kindParsedPartitionSpecrz   frozenset[MeshAxisName]r6   Nr"   rz   r6   r"   c               ~    || _         || _        || _        || _        t	          | j         | j        |          | _        d S r?   )r   r   r   r6   
preprocessrz   )r   r   r   r"   rz   r6   s         r)   __init__zNamedSharding.__init__   s@    
 DIDI#D$D#DIty-HHDr(   c                    d                     d | j        j                                        D                       }| j        dn	d| j         }d| d| j         | dS )N, c              3  ,   K   | ]\  }}d | d| V  dS )'z': Nr'   )rD   krE   s      r)   rF   z)NamedSharding.__repr__.<locals>.<genexpr>   s2      LLDAq-a--A--LLLLLLr(    , memory_kind=zNamedSharding(mesh=Mesh(z), spec=))joinr   r/   rY   r"   r   )r   	mesh_reprmems      r)   __repr__zNamedSharding.__repr__   so    		LLDIO4I4I4K4KLLLLLI ("".Qt?O.Q.QCJiJJJCJJJJr(   c                X    t          |           | j        | j        f| j        | j        dfS )Nr"   r6   )typer   r   r"   r6   r   s    r)   
__reduce__zNamedSharding.__reduce__   s6    JJDI. ,!.0 01 1r(   r;   c                    | j         S r?   r   r   s    r)   r"   zNamedSharding.memory_kind       r(   c                    t          | d          s,t          | j        | j        | j        | j        f          | _        | j        S N_hash)hasattrrK   r   r"   rz   r6   r   r   s    r)   __hash__zNamedSharding.__hash__   sG    4!! P9d&(:D<M
NP Pdj:r(   c                    t          |t                    sdS | |u rdS | j        |j        k    s | j        |j        k    s| j        |j        k    rdS | j        |j        u p| j        |j        k    S NFT)r@   r   rz   r"   r6   r   r   others     r)   __eq__zNamedSharding.__eq__   s|    e]++ Uu}}Te111u000 222U9
"=di5:&==r(   
aval_shaperQ   Nonec           
        | j         J t          |          t          | j                   k     rQt          |          dk    rdnd}t          d|  dt          | j                    dt          |           d|           d S )Nr   z, For scalars the PartitionSpec should be P()r   	Sharding + is only valid for values of rank at least %, but was applied to a value of rank r.   )rz   r   r0   )r   r   	extra_msgs      r)   check_compatible_avalz#NamedSharding.check_compatible_aval   s    )))
:T/0000*oo** BA02 +d + +#$$+ +__+ +(+ +, , , 10r(   r   c               F     | ||                                 |||          S )Nr   )r1   )clsr   r5   r"   r6   s        r)   _from_parsed_pspecz NamedSharding._from_parsed_pspec   s5     3t\4466'|)+ + + +r(   set[Device]c                    | j         j        S r?   )r   _flat_devices_setr   s    r)   
device_setzNamedSharding.device_set  s    9&&r(   XLADeviceAssignmentc                    | j         j        S r?   )r   _flat_devices_tupler   s    r)   _device_assignmentz NamedSharding._device_assignment
  s    9((r(   boolc                    | j         j         S r?   )r   is_multi_processr   s    r)   is_fully_addressablez"NamedSharding.is_fully_addressable  s     y)))r(   c                    | j         j        S r?   )r   _local_devices_setr   s    r)   addressable_devicesz!NamedSharding.addressable_devices  s     9''r(   c                    | j         j        dk    rdS t          t          t	          | j                            }| j         j        }d}|D ]}|||         z  }|dk    S )NrT   T)r   r   r
   r   ry   rz   r/   )r   r   r   num_partitionsrh   s        r)   is_fully_replicatedz!NamedSharding.is_fully_replicated  sk    y~T,.?@R.S.STTMJN ) )
4((nnQr(   kindr!   c                :    t          | j        | j        |          S Nr"   )r   r   r   r   r   s     r)   with_memory_kindzNamedSharding.with_memory_kind%  s    DI4@@@@r(   rb   r<   rc   c                "    t          | |          S r?   )r   r   rb   s     r)   _to_xla_hlo_shardingz"NamedSharding._to_xla_hlo_sharding(  s    -dNCCCr(   sharding.SdyArrayShardingc                    d t          |          D             }t          | j                  D ]%\  }}|d||         _        |s|||         _        &t          j        d|          S )Nc                :    g | ]}t          j        g d           S T)r   	is_closedr   SdyDimShardingrD   _s     r)   
<listcomp>z2NamedSharding._to_sdy_sharding.<locals>.<listcomp>,  s8     5 5 5 ,"EEE 5 5 5r(   Fr   )ranger{   rz   r   r   r   SdyArraySharding)r   rb   dim_shardingsrg   dim_specs        r)   _to_sdy_shardingzNamedSharding._to_sdy_sharding+  s    5 5#N335 5 5M !344 ) )8		%*a"" ) (a$V];;;r(   )r   r   r   r   r"   r   r;   r   r   rQ   r;   r   r;   r   r;   r   r;   r   )r   r!   r;   r   rb   r<   r;   rc   rb   r<   r;   r   )r#   r$   r%   __doc__r&   r   	frozensetr   r   r   propertyr"   r   r   r   classmethodr   r   r   r   r   	functoolscached_propertyr   r   r   r  r'   r(   r)   r   r      s1           D $$$$''''> !%D9;;I I I I I IK K K
1 1 1
    8  	> 	> 	>, , , , .2+ + + + ;+ ' ' ' 8' ) ) ) 8) * * * 8*
 ( ( ( 8(
    A A A AD D D D< < < < < <r(   r   c                 >    t           j                                        S r?   )r}   r   r   r'   r(   r)   get_replicated_hlo_shardingr  9  s    		!	!	#	##r(   c                     e Zd ZU dZded<   ded<    e            ddd%d
            Zd Zd Zd Z	d Z
ed&d            Zed'd            Zd(dZd)dZed*d            Zd+dZd,d!Zed-d#            Zed-d$            ZdS ).SingleDeviceShardingzA :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_devicer   r   Nr   r]   r"   c               "    || _         || _        d S r?   )r  r   )r   r]   r"   s      r)   r   zSingleDeviceSharding.__init__N  s    DL#Dr(   c                @    t          |           | j        fd| j        ifS Nr"   )r   r  r   r   s    r)   r   zSingleDeviceSharding.__reduce__S  s!    ::8I(JJJr(   c                B    | j         dn	d| j          }d| j        | dS )Nr   r   zSingleDeviceSharding(device=r   )r   r  r   r   s     r)   r   zSingleDeviceSharding.__repr__V  s7    !)""/S@Q/S/SC@$,@#@@@@r(   c                p    t          | d          s t          | j        | j        f          | _        | j        S r   )r   rK   r  r"   r   r   s    r)   r   zSingleDeviceSharding.__hash__Z  s5    4!! :t'7899dj:r(   c                |    t          |t                    sdS | |u rdS | j        |j        k    o| j        |j        k    S r   )r@   r  r  r"   r   s     r)   r   zSingleDeviceSharding.__eq___  sL    e122 Uu}}TLEM) 2 113r(   r;   r   c                    | j         hS r?   r  r   s    r)   r   zSingleDeviceSharding.device_setg  s    L>r(   c                    | j         S r?   r   r   s    r)   r"   z SingleDeviceSharding.memory_kindk  r   r(   r   r!   c                .    t          | j        |          S r   )r  r  r   s     r)   r   z%SingleDeviceSharding.with_memory_kindo  s    $????r(   rP   rQ   Mapping[Device, Index]c                P    | j         t          d           ft          |          z  iS r?   )r  rA   r   r   rP   s     r)   rU   z(SingleDeviceSharding.devices_indices_mapr  s$    L5;;.3|+<+<<==r(   r   c                    | j         fS r?   r  r   s    r)   r   z'SingleDeviceSharding._device_assignmentu  s    L?r(   rb   r<   rc   c                    t                      S r?   )r  r   s     r)   r   z)SingleDeviceSharding._to_xla_hlo_shardingy  s    &(((r(   r   c                Z    t          j        dd t          |          D                       S )Nr   c                :    g | ]}t          j        g d           S r   r   r   s     r)   r   z9SingleDeviceSharding._to_sdy_sharding.<locals>.<listcomp>  s8     	) 	) 	) 
	 bD	9	9	9 	) 	) 	)r(   )r   r   r   r   s     r)   r  z%SingleDeviceSharding._to_sdy_sharding|  s>    $	) 	)''	) 	) 	)* * *r(   r   c                    dS NTr'   r   s    r)   r   z(SingleDeviceSharding.is_fully_replicated      4r(   c                    dS r(  r'   r   s    r)   r   z)SingleDeviceSharding.is_fully_addressable  r)  r(   )r]   r  r"   r   r  r  )r   r!   r;   r  rP   rQ   r;   r   r  r  r  r  )r#   r$   r%   r	  r&   r   r   r   r   r   r   r  r   r"   r   rU   r   r   r  r   r   r'   r(   r)   r  r  >  s        	 	 ///>BF $ $ $ $ $ $K K KA A A  
3 3 3    8    8@ @ @ @> > > >    8) ) ) )* * * *    8    8  r(   r  r   c                    |                      |           t          j        || j                  }t	          t          | j        j        |                    S r?   )shard_shaper   spec_to_indicessharding_specdictr   devicesflat)r   rP   indicess      r)   !pmap_sharding_devices_indices_mapr4    sJ     <   *<9KLL'	ht|('22	3	33r(   c                  f   e Zd ZU dZded<   ded<   ded<    e            d1d	            Zd
 Zd Zd Z	d Z
d Zd2dZe	 	 d3d4d            Zej        d5d            Zd6d!Zej        d7d#            Zed8d%            Zd9d(Zd:d+Zd;d-Zej        d<d.            Zej        d<d/            Zd=d0ZdS )>PmapShardingz.Describes a sharding used by :func:`jax.pmap`.
np.ndarrayr1  sharding_specs.ShardingSpecr/  xc.DeviceList_internal_device_listSequence[Device] | np.ndarrayc                F    t          j        |          | _        || _        d S r?   )npasarrayr1  r/  )r   r1  r/  s      r)   r   zPmapSharding.__init__  s#     :g&&DL&Dr(   c                L    t          |           | j        | j        fd| j        ifS r  )r   r1  r/  r"   r   s    r)   r   zPmapSharding.__reduce__  s,    JJt'9:D,-/ /r(   c                    t          |t                    sdS | |u rdS | j        |j        k    o)| j        j        |j        j        k    o| j        |j        k    S r   )r@   r6  r/  r1  r/   r:  r   s     r)   r   zPmapSharding.__eq__  sg    e\** Uu}}T%"55 FL%-"55F&%*EEGr(   c                p    t          | d          s t          | j        | j        f          | _        | j        S r   )r   rK   r:  r/  r   r   s    r)   r   zPmapSharding.__hash__  s7    4!! J3T5GHIIdj:r(   c           	         d | j         j        D             }d| j         d|d| j         j        d         j                                         d| j         j         d	S )Nc                    g | ]	}|j         
S r'   )idrD   ds     r)   r   z(PmapSharding.__str__.<locals>.<listcomp>  s    2221!$222r(   PmapSharding(sharding_spec=z, device_ids=z, device_platform=r   z, device_shape=r   )r1  r2  r/  platformupperr/   )r   
device_idss     r)   __str__zPmapSharding.__str__  s    22 1222J2$*< 2 22 2#|03<BBDD2 2 !L.2 2 2 3r(   c                (    d| j          d| j         dS )NrG  z
, devices=r   )r/  r1  r   s    r)   r   zPmapSharding.__repr__  s/    '$*< ' '|' ' ' (r(   r   r   ndimr<   r;   r   c                    | |k    S r?   r'   )r   r   rM  s      r)   is_equivalent_tozPmapSharding.is_equivalent_to  s    5=r(   r   Nr/   rQ   sharded_dimSequence[xc.Device] | Nonec                   t          j        t          |          |          }d}|j        D ]{}t	          |t           j                  r|J |j        }(t	          |t           j                  r9|J t          |j	                  dk    r|j	        d         }lt          d          ||t          d          |.t          j        t          j                    d|                   }nt          j        |          } | ||          S )a  Creates 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`.
    NrT   r   z3Multiple chunks in Chunked dimension not supported.zY`None` to sharded_dim is not supported. Please file a jax issue if you need this feature.)r   create_pmap_sharding_specrL   r   r@   	Unstackedr   Chunkedr   chunksNotImplementedErrorr=  arrayr   local_devices)r   r/   rP  r1  r/  num_ways_shardedspmap_devicess           r)   defaultzPmapSharding.default  s@    #<ek# #M # 
E 
E	A~/	0	0 	E'''6a/00 E'''qx==AXa[

#CE E EE ,- - - !#

"
$
$%6&6%6
7"9 "9ll Xg&&l3|]+++r(   r   c                4    t          | j        j                  S r?   )r   r1  r2  r   s    r)   r   zPmapSharding.device_set  s    t| !!!r(   rP   r   c                "    t          | |          S r?   )r4  r"  s     r)   rU   z PmapSharding.devices_indices_map  s    ,T<@@@r(   r   c                4    t          | j        j                  S r?   )rL   r1  r2  r   s    r)   r   zPmapSharding._device_assignment  s    "###r(   r   c                ,    	 | j         j        S #  Y d S xY wr?   )r:  default_memory_kindr   s    r)   r"   zPmapSharding.memory_kind  s#    ';;TTs    r   r!   c                     t          d          )Nzpmap does not support memories.rW  r   s     r)   r   zPmapSharding.with_memory_kind   s    
?
@
@@r(   rb   rc   c                     t          d          )Nzpmap doesn't use OpSharding.rd  r   s     r)   r   z!PmapSharding._to_xla_hlo_sharding  s    
<
=
==r(   r   c                     t          d          )Nz"pmap doesn't use SdyArraySharding.rd  r   s     r)   r  zPmapSharding._to_sdy_sharding  s    
B
C
CCr(   c                v    | j         j        D ]+}t          |t          j        t          j        f          r dS ,dS r   )r/  r   r@   r   rT  rU  )r   r[  s     r)   r   z PmapSharding.is_fully_replicated	  sE    (  	A0.2HI	J	J uu4r(   c                    | j         j        S r?   r:  r   r   s    r)   r   z!PmapSharding.is_fully_addressable      %::r(   c                   d }d }t          | j        j                  D ]\  }}t          |t          j                  r |}|j        }t          j        ||          } nit          |t          j	                  rN|}t          |j                  dk    sJ |j                    |j        d         }t          j        ||d          } n||S ||         |k    r3t          d| d||          d| dt          | j                             |S )NrT   r   zkThe 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=)r{   r/  r   r@   r   rT  r   r   tuple_deleterU  r   rV  tuple_updater0   r   )r   rP   rP  sharded_dim_sizerg   r[  sharded_shapes          r)   r-  zPmapSharding.shard_shape  s^   K$,566  1	A~/	0	0 
6),DDa/00 18}}!!!18!!!8A;),QGG K $444BCNB B$[1B B=IB B $'t'>#?#?B BC C C
 r(   )r1  r;  r/  r8  )r   r6  r   r6  rM  r<   r;   r   )r   N)r/   rQ   rP  r<   r1  rQ  r;   r6  r  r+  r  r  )r   r!   r  r  r  )rP   rQ   r;   rQ   )r#   r$   r%   r	  r&   r   r   r   r   r   rK  r   rO  r  r]  r  r  r   rU   r   r  r"   r   r   r  r   r   r-  r'   r(   r)   r6  r6    s        66,,,,&&&&>' ' ' '/ / /G G G  
3 3 3( ( (   
 4548(, (, (, (, ;(,T " " " "A A A A $ $ $ $    8A A A A> > > >D D D D     ; ; ; ;     r(   r6  op_shardingxc.OpSharding | xc.HloShardingdevice_assignmentSequence[xc.Device]r"   r   PositionalShardingc                   t          | t          j                  rt          j                            |           } |                                 r#t          ||                                          S t          | 	                                          dk    rt          d          |d         j                                        t          j        fd|                                 D                       }t                              t#          |          ||          }|                    |                                           }|                                 r|                    dd          }|S )	Nr   rT   5Unhandled HloSharding type. Please open a bug report!r   c                0    g | ]}t          |          S r'   DeviceIdSetrf   s     r)   r   z0_op_sharding_to_pos_sharding.<locals>.<listcomp>@  s#    KKK{4KKKr(   F)keepdims)r@   r}   r~   r   
from_protois_replicatedrt  r   r   rx   rW  rH  rI  r=  rX  tile_assignment_devices_remakerL   reshapetile_assignment_dimensionsreplicate_on_last_tile_dim)rp  rr  r"   idsr7   rh   s        @r)   _op_sharding_to_pos_shardingr  .  se    R]++ 9.++K88K   @{4 4 44=IKK@ 		#	#	%	%&&**
?   
1		&	,	,	.	.$
KKKK[%H%H%J%JKKK	 	#   '8!9!93-8 ! : :!ii668899!++-- (	B''A	
(r(   c                \   | j         d| j        z  k    rt                      S t          j                    }| j         | j        |z
  d          }d | j        j        D             \  }t          j        j        j        |_	        |dk    r)t          j        j        j
        g|_        g ||R |_        n||_        d | j        j        D             |_        t          j        |j                  }t!          |j                  }||k    sJ ||f            t          j                            |          S )NrT   c                ,    h | ]}t          |          S r'   r   )rD   r   s     r)   rq   z;_positional_sharding_to_xla_hlo_sharding.<locals>.<setcomp>R  s    @@@:s:@@@r(   rT   c                    g | ]	}|D ]}|
S r'   r'   )rD   r  rg   s      r)   r   z<_positional_sharding_to_xla_hlo_sharding.<locals>.<listcomp>Y  s%    !K!K!Ks!K!K!!!K!K!K!Kr(   )r/   rM  r  r}   r~   _idsr2  r   OTHERr   r   r   r  r~  mathprodr   r   r|  )r   rb   pbufr/   set_sizeproduct_of_dimsnum_devicess          r)   (_positional_sharding_to_xla_hlo_shardingr  J  s    
Z4$)###&(((	$
*TY/00
1%@@@@@)(m &$)\\=-89D&8&8x&8&8D##&+D#!K!Kty~!K!K!K$Id=>>/D011+	K	'	'	'/;)G	'	'	'		"	"4	(	((r(   c                      e Zd ZU ded<   ded<   ded<   ddd. fdZed             Zed             Zd/dZd0dZ	d0dZ
 ee
          Zd1d0dZd2dZeddd3d            Zd4dZd5dZej        d6d!            Zed7d"            Zd8d$Zej        d5d%            Zed9d'            Zd:d*Zd;d,Zej        d5d-            Z xZS )<rt  tuple[xc.Device, ...]_devicesr   r   r7  r  Nr   r1   Sequence[xc.Device] | np.ndarrayr"   c                  t                                                       t          |t          j                  st          j        |d          }|j        st          | j        j	         d|           t          |j                  | _        || _        | j        d         j                                        t          j        fdt!          |j                  D             d                              |j                  | _        t)          j        | j                  | _        t)          j        | j        | j                  | _        d S )Nobjectdtypez,.__init__ requires at least one device, got r   c                0    g | ]}t          |          S r'   rx  rf   s     r)   r   z/PositionalSharding.__init__.<locals>.<listcomp>p  s#    LLL1+dA..LLLr(   )superr   r@   r=  ndarrayrX  r   r0   	__class__r#   rL   r2  r  r   rH  rI  r   r  r/   r  r}   
DeviceListr:  "check_and_canonicalize_memory_kind)r   r1  r"   rh   r  s      @r)   r   zPositionalSharding.__init__e  s8   	GGgrz** 2111g< 5$.1 4 4*14 4 5 5 5',''DM#D=$**,,DLLLLgl8K8KLLL') ) ))0)?)? 	I!#t}!=!=D=457 7Dr(   c                    | j         j        S r?   )r  r/   r   s    r)   r/   zPositionalSharding.shapev  s    9?r(   c                    | j         j        S r?   )r  rM  r   s    r)   rM  zPositionalSharding.ndimz  s    9>r(   r;   r!   c                     j         j        } j                                        } j        d         j                                        }t          j        |          D ]!\  }}t          |g fd|D             R  ||<   "t          j
        ||dz   dd          } j        dn	d j         }| d| | d	 j         dS )
Nr   c              3  <   K   | ]}j         |         j        V  d S r?   )r  rD  )rD   rg   r   s     r)   rF   z.PositionalSharding.__repr__.<locals>.<genexpr>  s-      -M-MadmA.>.A-M-M-M-M-M-Mr(   (r   d   )prefixsuffixmax_line_widthr   r   z, shape=)r  r#   r  copyr  rH  rI  r=  ndenumeratery  array2stringr   r/   )r   cls_namer  platform_nameidxrN   bodyr   s   `       r)   r   zPositionalSharding.__repr__~  s    ~&H
)..

CM!$-3355M.%% O OQ]N-M-M-M-M1-M-M-MNNNc#hh?3x#~c*-/ / /D!)""/S@Q/S/SC999s99DJ9999r(   c                R    |                      | j         | j        j        |           S r?   )r  r  r  r  )r   r/   s     r)   r  zPositionalSharding.reshape  s%    <<'8ty'8%'@AAAr(   c                R    |                      | j         | j        j        |           S r?   )r  r  r  	transposer   r   s     r)   r  zPositionalSharding.transpose  s%    <<':ty':D'ABBBr(   Tc                p    | j                             ||          }|                     | j        |          S )N)axisr{  )r  sumr  r  )r   r  r{  new_idss       r)   r   zPositionalSharding.replicate  s/    immm99G<<w///r(   r   rQ   r   c           
         t          |          t          | j                  k    r>| j        s9t          d|  dt          | j                   dt          |                     d S d S )Nr   z" is only valid for values of rank r   )r   r/   r   r0   )r   r   s     r)   r   z(PositionalSharding.check_compatible_aval  s|    
:#dj//))$2J)d  __ __       *)))r(   r  c                   |                      |           }||_        ||_        t          j        |j                  |_        t          j        ||j                  |_        |S r?   )__new__r  r  r}   r  r:  r  r   )r   r1  r  r"   r   s        r)   r  zPositionalSharding._remake  sY     ;;sDDMDI!#t}!=!=D=T/1 1DKr(   r<   c                p    t          | d          s t          | j        | j        f          | _        | j        S r   )r   rK   r:  r"   r   r   s    r)   r   zPositionalSharding.__hash__  s7    4!! H3T5EFGGdj:r(   r   c                    t          |t                    sdS | |u rdS t          j        | j        |j                  }| j        |j        k    }| j        |j        u r|r|rdS |o|o| j        |j        k    S r   )r@   rt  r=  array_equalr  r"   r  r:  )r   r   all_ids_equalmem_kind_equals       r)   r   zPositionalSharding.__eq__  s    e/00 Uu}}TN49UZ88M%)::N}&&>&m&T F} F&%*EEGr(   set[xc.Device]c                *    t          | j                  S r?   r   r  r   s    r)   r   zPositionalSharding.device_set      t}r(   c                    | j         S r?   r   r   s    r)   r"   zPositionalSharding.memory_kind  r   r(   r   c                .    t          | j        |          S r   )rt  r  r   s     r)   r   z#PositionalSharding.with_memory_kind  s    dm>>>>r(   c                (    | j         d| j        z  k    S )Nr  )r/   rM  r   s    r)   r   z&PositionalSharding.is_fully_replicated  s    :	)))r(   r   c                    | j         S r?   r  r   s    r)   r   z%PositionalSharding._device_assignment  
    =r(   rb   rc   c                "    t          | |          S r?   )r  r   s     r)   r   z'PositionalSharding._to_xla_hlo_sharding  s    3D.IIIr(   r   c                     t          d          )Nz=PositionalSharding can't be converted to an SdyArraySharding.rd  r   s     r)   r  z#PositionalSharding._to_sdy_sharding  s    
GI I Ir(   c                    | j         j        S r?   ri  r   s    r)   r   z'PositionalSharding.is_fully_addressable  rj  r(   )r1  r  r"   r   r;   r!   )r;   rt  r(  r  )r1  r  r  r7  r"   r   r;   rt  r;   r<   r  )r;   r  r  )r   r!   r;   rt  r  r  r  )r#   r$   r%   r&   r   r  r/   rM  r   r  r  Tr   r   r  r  r   r   r  r  r   r"   r   r   r   r   r  r   __classcell__r  s   @r)   rt  rt  `  ss        !!!! -17 7 7 7 7 7 7 7"   8   8	: 	: 	: 	:B B B BC C C Chy!0 0 0 0 0         $(	 	 	 	 	 ;	   

G 
G 
G 
G        8? ? ? ? * * * *
    8J J J JI I I I ; ; ; ; ; ; ; ;r(   c                  X    e Zd ZU ded<   ded<   d Zd ZddZdd
ZddZddZ	ddZ
dS )ry  r!   _namezfrozenset[int]r  c                <    || _         t          |          | _        d S r?   )r  r
  r  )r   rh   r  s      r)   r   zDeviceIdSet.__init__  s    DJ#DIIIr(   c                D    t          t          | j                            S r?   )iterr   r  r   s    r)   __iter__zDeviceIdSet.__iter__  s    ty!!"""r(   r;   c                n    t          |t                    sJ t          | j        g| j        |j        z  R  S r?   r@   ry  r  r  r   s     r)   __add__zDeviceIdSet.__add__  s9    e[)))))tz=TY%;====r(   r<   c                *    t          | j                  S r?   )r   r  r   s    r)   __len__zDeviceIdSet.__len__  s    ty>>r(   c                    d                     t          t          t          | j                                      }d| j         d| dS )Nr   { })r   r   r!   r   r  r  )r   r  s     r)   r   zDeviceIdSet.__repr__  sB    
))HS&"3"344
5
5C$
$$S$$$$r(   c                8    t          | j        | j        f          S r?   )rK   r  r  r   s    r)   r   zDeviceIdSet.__hash__  s    TY'(((r(   r   c                l    t          |t                    o| j        |j        k    o| j        |j        k    S r?   r  r   s     r)   r   zDeviceIdSet.__eq__  s5    uk** $tzU[/H $I#%r(   N)r;   ry  r  r  r  )r#   r$   r%   r&   r   r  r  r  r   r   r   r'   r(   r)   ry  ry    s         ***  # # #> > > >   % % % %) ) ) )% % % % % %r(   ry  c                  ~   e Zd ZU ded<   ded<   ded<   ded<   d	ed
<    e            dddd/d            Zd Zej        d             Z	d Z
d Zd Zd0dZej        d1d            Zed2d            Zd3d"Zed4d$            Zd5d'Zd6d)Zej        d7d+            Zej        d7d,            Zedd-d8d.            ZdS )9GSPMDShardingztuple[Device, ...]r  rc   _hlo_shardingr   r   xc.DeviceList | None_device_listr9  r:  N)r"   r  r1  Sequence[Device]rp  rq  r"   c                   t          |          | _        t          |t          j                  r%t          j                            |          | _        n|| _        || _        d S r?   )	rL   r  r@   r}   r~   r   r|  r  r   )r   r1  rp  r"   r  s        r)   r   zGSPMDSharding.__init__  sV    
 'NNDM+r}-- '>44[AAd&d#Dr(   c                p    t          |           | j        | j                                        fd| j        ifS r  )r   r  r  to_protor   r   s    r)   r   zGSPMDSharding.__reduce__  s8    JJ(:(C(C(E(EFD-.0 0r(   c                n    | j         rt          t                                S t          | j                  S r?   )r   rK   r  r  r   s    r)   _hlo_sharding_hashz GSPMDSharding._hlo_sharding_hash  s3     1-//000"###r(   c                    t          |t                    sdS | |u rdS t          | j        |j                  o| j        |j        k    o| j        |j        k    S r   )r@   r  r   r  r"   r:  r   s     r)   r   zGSPMDSharding.__eq__  sh    e]++ Uu}}T"4#5u7JKK J E$55J*e.IIKr(   c                |    t          | d          s&t          | j        | j        | j        f          | _        | j        S r   )r   rK   r:  r  r"   r   r   s    r)   r   zGSPMDSharding.__hash__!  sB    4!! +3T5L(* + +dj:r(   c                B    | j         dn	d| j          }d| j        | dS )Nr   r   zGSPMDSharding(r   )r   r  r  s     r)   r   zGSPMDSharding.__repr__'  s8    !)""/S@Q/S/SC8D.8#8888r(   r   rQ   r;   r   c           
         t          | j                  \  }}t          |          t          |          k     r2t          d|  dt          |           dt          |                     d S )Nr   r   r   )r   r  r   r0   )r   r   num_ways_dim_shardedr   s       r)   r   z#GSPMDSharding.check_compatible_aval+  s    6t7IJJ!
:12222d  %&& __       32r(   r   c                *    t          | j                  S r?   r  r   s    r)   r   zGSPMDSharding.device_set3  r  r(   c                    | j         S r?   r   r   s    r)   r"   zGSPMDSharding.memory_kind7  r   r(   r   r!   c                :    t          | j        | j        |          S r   )r  r  r  r   s     r)   r   zGSPMDSharding.with_memory_kind;  s    (:MMMMr(   r   c                    | j         S r?   r  r   s    r)   r   z GSPMDSharding._device_assignment>  r  r(   rb   r<   c                    | j         S r?   )r  r   s     r)   r   z"GSPMDSharding._to_xla_hlo_shardingB  s    r(   r   c                     t          d          )Nz5GSPMDSharding can't be converted to SdyArraySharding.rd  r   s     r)   r  zGSPMDSharding._to_sdy_shardingE  s    
?A A Ar(   r   c                *    t          | j                  S r?   )r   r  r   s    r)   r   z!GSPMDSharding.is_fully_replicatedI  s    $T%7888r(   c                    | j         j        S r?   ri  r   s    r)   r   z"GSPMDSharding.is_fully_addressableM  rj  r(   r   c               P     | t          |          t                      |          S r   )rL   r  )r   rr  r"   s      r)   get_replicatedzGSPMDSharding.get_replicatedQ  s2    3u&'')D)F)F&( ( ( (r(   )r1  r  rp  rq  r"   r   r  r  r  r  r  )r   r!   r;   r  r  r  r  r  )r"   r   )r#   r$   r%   r&   r   r   r   r  r  r  r   r   r   r   r   r  r"   r   r   r   r  r   r   r  r  r'   r(   r)   r  r    s        $$$$&&&&> -148	$ 	$ 	$ 	$ 	$ 	$0 0 0 $ $ $
K K K  9 9 9               8N N N N    8   A A A A 9 9 9 9 ; ; ; ; JN ( ( ( ( ( ;( ( (r(   r  c                      e Zd ZddZdS )AUTOr   r   c                    || _         d S r?   r   )r   r   s     r)   r   zAUTO.__init__Y  s    DIIIr(   N)r   r   )r#   r$   r%   r   r'   r(   r)   r  r  W  s(             r(   r  c                ,    t          | t                    S r?   )r@   r  rM   s    r)   is_autor  ]  s    	At		r(   c                      e Zd Zd ZdS )UnspecifiedValuec                    dS )Nr  r'   r   s    r)   r   zUnspecifiedValue.__repr__b  s    r(   N)r#   r$   r%   r   r'   r(   r)   r  r  a  s#            r(   r  c                ,    t          | t                    S r?   )r@   r  rM   s    r)   is_unspecifiedr  f  s    	A'	(	((r(   c                >    t          |           pt          |           S r?   )r  r  rM   s    r)   is_unspecified_or_autor	  i  s    		(~a(((r(   r   ArrayMappingc                   | st                      S d}t          j        t                    }|                                 D ](\  }}||                             |           ||k    r|})g }t          |dz             D ]_}||         }|r>|                    t          |          dk    r|d         nt          |                     J|                    d            `t          | S )Nrz  rT   r   )	r   rW   r   r   rY   r   r   r   rL   )r   	max_indexreverse_mapr  r^   
partitionsrg   s          r)   array_mapping_to_axis_resourcesr    s    	 ??)'--+"((**  kdEd###yi*Q  aq>D 3t99>>QuT{{CCCC	
	##r(   axis_resources-ParsedPartitionSpec | AUTO | UnspecifiedValueArrayMappingOrAutoOrUnspecifiedc                    t          | t          t          f          r| S t          d t	          |           D                       S )Nc              3  2   K   | ]\  }}||D ]}||fV  	d S r?   r'   )rD   rg   r   r  s       r)   rF   z$get_array_mapping.<locals>.<genexpr>  sI       ; ; D(T((T AY(((((; ;r(   )r@   r  r  r   r{   )r  s    r)   ry   ry     sX    
 '7 899 	 ; ;$-n$=$=; ; ; 
; 
; ;r(   c                `    t          t          t          t          |                               S r?   )r  r
   r
  ry   )r7   s    r)   rl   rl     s'    <(++,,. . r(   c                      e Zd ZdZdZdZdZdS )SpecSynczEncodes how much out of sync the real value of partitions is compared to the user specified one.

  We use this to make sure we don't show garbage modified values while claiming
  that the users have specified them like that.
  r   rT      N)r#   r$   r%   r	  OUT_OF_SYNCDIM_PERMUTEIN_SYNCr'   r(   r)   r  r    s)         
 ++'''r(   r  c                      e Zd ZdZej        fdZed             ZddZ	d Z
d Zedd
            Zd Zd Zd Zd Zd Zd ZdS )r   )unsafe_user_specr  syncc                J    || _         t          |          | _        || _        d S r?   )r  rL   r  r  )r   r4   r  r  s       r)   r   zParsedPartitionSpec.__init__  s&    %D J''DODIIIr(   c                @    |                      t          j                  S r?   )unsynced_user_specr  r  r   s    r)   r4   zParsedPartitionSpec.user_spec  s    ""8#3444r(   r;   r   c                    | j         t          j        k     rt          |           S t	          | j        t                    r| j        S t          |           S r?   )r  r  r  get_single_pspecr@   r  r   r   s    r)   r1   z&ParsedPartitionSpec.get_partition_spec  sL    y8###d###	D)=	9	9 &$$%%%r(   c                \    | j         |k     rt          d| j          d| d          | j        S )NzPlease open a bug report! (z >= r   )r  AssertionErrorr  )r   min_syncs     r)   r!  z&ParsedPartitionSpec.unsynced_user_spec  s<    y8SSSSSSTTT  r(   c                    | j         }|t          |          z
  }|dk    r|d|z  z  }t          j        |||          }|dk    s|t          j        nt          j        }t          | j        ||          S )Nr   )r'   r'   )r  )	r  r   r   tuple_insertr  r  r  r   r  )r   dimvalparts	too_shortnew_partitionsnew_syncs          r)   insert_axis_partitionsz*ParsedPartitionSpec.insert_axis_partitions  sy    OEc%jj I1}}uy  e&uc377N(+r		S[x##xG[Ht4n8TTTTr(   Fc                   | | |d          S t          |t                    st          | d|           g }|D ]r}|d}nVt          |t          t          f          rt	          |          }n*|t          j        k    r|st          d|           d }n|f}|                    |           st          d |D              } | ||          S )Nr'   z= are expected to be PartitionSpec instances or None, but got z$Unconstrained dims are not allowed: c                h    g | ]/}t          |t          t          f          rt          |          n|0S r'   )r@   r   rL   )rD   r9   s     r)   r   z7ParsedPartitionSpec.from_user_input.<locals>.<listcomp>  s5    	I	I	IajT5M22
9%(((	I	I	Ir(   )r@   r   	TypeErrorr   rL   UNCONSTRAINEDr0   r   )r   entryarg_nameallow_unconstrained_dims
axis_specs	axis_spec	new_entrys          r)   from_user_inputz#ParsedPartitionSpec.from_user_input  s#   }S^^e]++ K J JBGJ J K K KJ # #					i$// !)$$		333' 	KI%IIJJ
J		L		""""	I	I5	I	I	IKI3y*%%%r(   c                8    t          | j        | j        f          S r?   )rK   r  r  r   s    r)   r   zParsedPartitionSpec.__hash__  s    $),---r(   c                B    | j         |j         k    o| j        |j        k    S r?   )r  r  r   s     r)   r   zParsedPartitionSpec.__eq__  s%    Ou// $I#%r(   c                *    t          | j                  S r?   )r   r  r   s    r)   r  zParsedPartitionSpec.__len__  s    tr(   c                    | j         |         S r?   )r  )r   rg   s     r)   __getitem__zParsedPartitionSpec.__getitem__  s    ?1r(   c                *    t          | j                  S r?   )r  r  r   s    r)   r  zParsedPartitionSpec.__iter__  s       r(   c                8    d| j          d| j         d| j         dS )NzParsedPartitionSpec(partitions=, unsafe_user_spec=, sync=r   r  r  r  r   s    r)   r   zParsedPartitionSpec.__repr__  s>    !do ! ! $ 5! !I! ! ! "r(   N)r;   r   F)r#   r$   r%   	__slots__r  r  r   r  r4   r1   r!  r/  r  r:  r   r   r  r?  r  r   r'   r(   r)   r   r     s        8)191A     5 5 85& & & &! ! !
U U U & & & ;&.. . .% % %       ! ! !" " " " "r(   r   c                  *     e Zd ZdZd fdZd Z xZS ) CanonicalizedParsedPartitionSpecaU  ParsedPartitionSpecs that are canonicalized.

  ParsedPartitionSpecs may contain trailing empty tuples, that make them
  semantically different in general, and yet in some situations we prefer
  to regard them as equivalent. For example, partitions of () and ((),)
  cannot be always considered equivalent, since the first one is a valid
  spec for a scalar value, while the second is not! However, when either of
  those are applied to a 2D array, they both mean that the array is fully
  replicated.

  So CanonicalizedParsedPartitionSpecs removes the trailing empty tuples from
  partitions.
  r5   r   c                    t          |j                  }|r.|d         dk    r"|                                 |r|d         dk    "t                                          |j        ||j                   d S )Nrz  r'   )r   r  popr  r   r  r  )r   r5   r  r  s      r)   r   z)CanonicalizedParsedPartitionSpec.__init__  s    l-..J
 B2--nn  B2-- 
GG\2J!&( ( ( ( (r(   c                8    d| j          d| j         d| j         dS )Nz,CanonicalizedParsedPartitionSpec(partitions=rB  rC  r   rD  r   s    r)   r   z)CanonicalizedParsedPartitionSpec.__repr__  s>    !4? ! ! $ 5! !I! ! ! "r(   )r5   r   )r#   r$   r%   r	  r   r   r  r  s   @r)   rH  rH    sV         ( ( ( ( ( (" " " " " " "r(   rH  c                p    |"t          |t                      n|dd          }t          | ||           |S )NzNamedSharding specTr6  )prepare_axis_resourcesr   r:   )r   r   r5   r6   s       r)   r   r     sL     )<Tt= = =L D,===	r(   i  c                   t          j        | d           \  }}| d}g }|D ]}t          |          s||                    |           )t	          |t
          j                  rAt	          |t                    rt          d| d| d          |                    |           |                    t          
                    |||                     t          ||           t          j        ||          S )Nc                
    | d u S r?   r'   rM   s    r)   rl   z(prepare_axis_resources.<locals>.<lambda>4  s
    T	 r(   )is_leafz leaf specificationszOne of z got sharding z which is not allowed.rM  )r   tree_flattenr	  r   r@   r   Shardingr6  r0   r   r:  _check_unique_resourcestree_unflatten)r  r5  r6  entriestreedefwhatnew_entriesr4  s           r)   rN  rN  /  sT    +113 3 3'7	*	*	*$+ 
K 
Kee$$ 	K	E8,	-	- K	E<	(	( % $4 $ $u $ $ $ % % 	%,<<
0H = J J K K K K +x000		!';	7	77r(   c                   | D ]}|st          |          st          |t          j                  r/d |D             }t	          j        t          j                            |                    }|so|	                    d          d         d         dk    rOd |
                                D             }|r/t          d| d|j         dt          j        |                     d S )Nc                    g | ]}||S r?   r'   rE  s     r)   r   z+_check_unique_resources.<locals>.<listcomp>N  s    GGGar(   rT   r   c                $    g | ]\  }}|d k    |S r  r'   )rD   r8   cs      r)   r   z+_check_unique_resources.<locals>.<listcomp>S  s!    FFFTQAqr(   z	A single zP specification can map every mesh axis to at most one positional dimension, but z has duplicate entries for )r	  r@   r   rS  rW   rX   	itertoolschainfrom_iterablemost_commonrY   r0   r4   mesh_lib	show_axes)r  r5  arg_axis_resourcesconstrained_dimsresource_countsmultiple_usess         r)   rT  rT  H  s6   * [ ['x122 %x'899GG#5GGG!)%%&6779 9O$H""1%%a(+a//FF_%:%:%<%<FFFm	 [ ZX Z ZEWEaZ Z6>6H6W6WZ Z [ [ 	[[ [r(   c                  2    e Zd ZU dZded<   ded<   ded<   dS )	AxisEnvz5Represents a pmap mesh (only along the replica axes).r<   nrepsztuple[Any, ...]namesztuple[int, ...]sizesNr#   r$   r%   r	  r&   r'   r(   r)   ri  ri  [  s9         ==***r(   ri  c                  p    e Zd ZU dZded<    e            Zded<   ed             Zed             Z	dd
Z
dS )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.
  r   r   r   manual_axesc                    | j         S r?   )unsafe_axis_envr   s    r)   axis_envzSPMDAxisContext.axis_envm  s     r(   c                    t          | j        j        | j        j        t	          | j        j                                                            S )Nrj  rk  rl  )ri  r   r   r|   rL   r/   r   r   s    r)   rr  zSPMDAxisContext.unsafe_axis_envs  sC    ini"DIO**,,--/ / / /r(   r   r;   c                <    t          | j        | j        |z            S r?   )ro  r   rp  r  s     r)   extend_manualzSPMDAxisContext.extend_manualz  s    49d&6&=>>>r(   N)r   r   r;   ro  )r#   r$   r%   r	  r&   r
  rp  r  rs  rr  rw  r'   r(   r)   ro  ro  b  s           )2+4444    8 
 / / 8/? ? ? ? ? ?r(   ro  c                      e Zd ZU dZded<   dS )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.
  ri  rs  Nrm  r'   r(   r)   ry  ry  ~  s*          
 r(   ry  c                  H    e Zd ZU dZded<   dZded<   d Zed             ZdS )	ShardingContextzA hardware axis context for parallel computations that use the sharding
  interface.

  This context also uses the GSPMD partitioner.
  r<   r  Nztuple[xc.Device, ...] | Nonerr  c                    | j         ;t          | j         t                    sJ | j        t	          | j                   k    sJ d S d S r?   )rr  r@   rL   r  r   r   s    r)   __post_init__zShardingContext.__post_init__  sQ    ).66666T%;!<!<<<<< *)<<r(   c                &    t          ddd          S )NrT   r'   ru  )ri  r   s    r)   rs  zShardingContext.axis_env  s    "B////r(   )	r#   r$   r%   r	  r&   rr  r}  r  rs  r'   r(   r)   r{  r{    sj          
 488888= = = 0 0 80 0 0r(   r{  c                x    t          j        | ddd                   ddd         t          j        |           z  S )z5Returns an array of strides for major-to-minor sizes.Nrz  )r=  cumprodr>  )rl  s    r)   strides_for_sizesr    s6    	E$$B$K	 	 2	&"*U*;*;	;;r(   c                N   d |                                  D             } t          j        |                                 t          j                  }t          |          }t          |t          |                    }d t          |||           D             fd|D             S )ai  Recovers 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.
  c                &    i | ]\  }}|d k    ||S r  r'   )rD   rh   r   s      r)   ri   z#unflatten_array.<locals>.<dictcomp>  s#    OOO
dTQYYtYYYr(   r  c                     i | ]\  }}}||f|S r'   r'   )rD   r   striderh   s       r)   ri   z#unflatten_array.<locals>.<dictcomp>  s%    aaa*<$$aaar(   c                     g | ]
}|         S r'   r'   )rD   rF  dim_to_names     r)   r   z#unflatten_array.<locals>.<listcomp>  s    	'	'	'Q+a.	'	'	'r(   )	rY   r=  fromiterr   int64r  explode_superdimsunflatten_superdimszip)named_sizes
assignmentrl  stridesru   r  s        @r)   unflatten_arrayr    s    & POk.?.?.A.AOOO+
+k((**"(
;
;
;%e$$'	5"5j"A"A	B	B$aaE7T_@`@`aaa+	'	'	'	'$	'	'	''r(   c                x   d }t          j        | t           j                  } ||d         dk               g }|j        dk    rt|d         }t	          t          |                    D ]}||         ||z  k    r n|dz  }|}|                    ||f           |dk    sJ |dd|         }|j        dk    t|S )a  Unflatten 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                (    | rd S t          d          )NzKFailed to convert OpSharding into a ShardingSpec. Please open a bug report!rd  )conds    r)   checkz"unflatten_superdims.<locals>.check  s"    OVV
 : ; ; ;r(   r  r   rT   N)r=  r>  r  r   r   r   r   )r  r  flat_assignmentru   r  rg   r   s          r)   r  r    s    ; ; ; Jz:::/%a   	$q  QF3''((  		q6z	)	)55	) 1faDKKv!8888%fff-O 	q   
+r(   c                   d t          | t          |                     D             }t          t          |                    }g }|D ]\  }}||         }g }||k    rB|dk    sJ ||z  dk    sJ |                    ||f           ||z  }||z  }||         }||k    B||k    sJ |                    ||f           |t          |          z  }|S )a  Explode 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)
  c                    i | ]\  }}||	S r'   r'   )rD   r   r  s      r)   ri   z%explode_superdims.<locals>.<dictcomp>  s    \\\|tVfd\\\r(   rT   r   )r  r  r   reversedr   )rl  ru   strides_to_sizes
final_dimsr   r  target_sizenew_dimss           r)   r  r    s    ]\s5BSTYBZBZ7[7[\\\	htnn		$* % %ldF"6*KH


1____K1$$$$oo{F+,,,
{df$V,k 

 ;OOT6N###(8$$$JJ	r(   hlo_shardingr   r   Sequence[ParsedPartitionSpec]c                D   t          | t          j                  rt          j                            |           } |                                 r>g }|                                 D ]%}|                    t          ||                     &|S |                                 r*t          t          t                      d                    gS |                                 r4|j        }t          |j        |                                           }t!          |          }|                                 }g }|D ]n}	g }
|	dk    rBt%          |          }||         }|	|z  dk    sJ |	|z  }	|
                    |           |	dk    B|                    t)          |
                     ot+          |                                           dk    rt/          d          |                                 r
|d d         }t          t          d|                    gS t3          d          )Nr'   rT   r   rv  rz  z<internally generated spec>z4Unhandled OpSharding type. Please open a bug report!)r@   r}   r~   r   r|  tuple_elementsr   parse_flatten_op_shardingr}  rH  r   r   is_tiledr/   r  r~  r  r  nextr   rL   r   rx   rW  r  r%  )r  r   r\   r[  r   mesh_axis_order	mesh_axisr/   r  dim_sizedim_partitionsr  	axis_sizes                r)   r  r    sH   bm,, ;>,,\::L  "" "Q%'C((** 5 5	jj*1d334444J!!## Q,MOOR002 2 3 3 QJ%
L88:: O _%%I3355EJ / /nqLLIt$	)#q((((Yd### qLL n--....
<&&(())A--
A   ..00 #crc?j,9:FFH H I I O
P
PPr(   r[  rA   c                0    | j         J | j        | j        fS r?   )rB   rH   rI   )r[  s    r)   _slice_as_tupler  (  s    	

'16	r(   c                      e Zd ZdZdS )NonUniformShardingErrorz5Raised when sharding is not uniform across processes.N)r#   r$   r%   r	  r'   r(   r)   r  r  -  s        ====r(   r  tensor_shardingsharding.Shardingr)  ndimstuple[int, int]c                   | j         s| j        rdS t          | j                  }|                     |f|z            }fd|                                D             }t          j        t                    }t                      }|                                D ]H\  }}	|	j	        |	j
        f}
||j                                     |
           |                    |
           It          t          | j                            j        }t!          ||                   }t          |          t#          fd|                                D                       rt'          d| dd          t)          d |                                D                       }t+          d |D                       t          |          k    rt'          d| d          |                    |          t          |          fS )	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`.
  )r   rT   c                (    i | ]\  }}||         S r'   r'   )rD   r   rE   r)  s      r)   ri   z/get_process_index_and_count.<locals>.<dictcomp>  s#    ;;;1!QsV;;;r(   c              3  >   K   | ]}t          |          k    V  d S r?   r  )rD   rN   slices_per_processs     r)   rF   z.get_process_index_and_count.<locals>.<genexpr>  s/      II!Q%	%IIIIIIr(   ztensor_sharding=z is non-uniform on dim=z3 as some processes have different number of slices.c                ,    h | ]}t          |          S r'   )r
  rp   s     r)   rq   z.get_process_index_and_count.<locals>.<setcomp>  s    KKKA9Q<<KKKr(   c              3  4   K   | ]}t          |          V  d S r?   r  )rD   hs     r)   rF   z.get_process_index_and_count.<locals>.<genexpr>  s(      **AQ******r(   )r   r   r   r   rU   rY   rW   r   r   rH   rI   process_indexaddr  r  r   r
  anyr   r  r   r  r^   )r  r)  r  r  
device_mapglobal_sliceprocess_to_slice
all_slicesrF  rE   rm   current_pidaddressable_slicesunique_processesr  s    `            @r)   get_process_index_and_countr  1  s9   F * )6O.//+
 22K>E3IJJ* <;;;
(8(8(:(:;;;, !,S11uu*   ""  da7AF
CQ_%))#...NN3 T/=>>??M+ !1+!>?? -..IIII/?/F/F/H/HIIIII 
!	&? 	& 	& 	& 	& 	&   KK1A1H1H1J1JKKKLL
 	**)*****c*oo==
!6?6666   
 
 !3
4
4c:J6K6K	LLr(   r   local_shapetuple[int | None, ...]c                    dgt          |          z  }t          |          D ]E\  }}	 t          | |t          |                    \  }}||z  ||<   1# t          $ r d||<   Y Bw xY wt	          |          S )aV  Computes 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  )r   r{   r  r  rL   )r   r  rP   rg   	local_dimr   shard_counts          r)   local_to_global_shaper    s    \ &*FS-=-=$=,,,  la2
AS--/ / /na!K/l1oo"   l1oh 
|		s   *AA&%A&c                    |                      |          }t          t          t          j        t
          f         |          }t          fd|                                D                       }|                     |                   }||z  S )a  Returns 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.
  c                :    h | ]}t          |                   S r'   )r  )rD   addressabler)  s     r)   rq   z*num_addressable_indices.<locals>.<setcomp>  s3       ,7ok#&''  r(   )	addressable_devices_indices_mapr
   r   r   r  Indexr   r   r-  )r  r)  rP   addressablesnum_unique_slices
shard_sizes    `    r)   num_addressable_indicesr    s    D !@@NN,ghou45|DD,    ;G;N;N;P;P     **<88=*	'	''r(   c                >   | j         j                            | j                   }|                                                                }t          |          \  }}|dk    rg n|g}|dg|j        z  z   |z   }||_        t          j	        
                    |          S rk   )r  _rulesphysical_element_avalr  cloner   rM  r  r}   r   r|  )avalr  elt_avalnew_op_shardingr  num_replicasr  tads           r)   physical_hlo_shardingr  
  s    Z44TZ@@( ))++1133/5lCC*l""22&aS8=((61#/2/,		"	"?	3	33r(   r   c                ^    t          | j                  dk    ot          | t                     S rk   )r   r   r@   r6  r   s    r)   is_single_device_shardingr    s,     
X 	!	!Q	&	Qz(L/Q/Q+QQr(   c                   t          |          r|S t          |t                    r| j        j                            | j                  }t          j                    g|j        z  }t          j	        g |j
        j        |R |j
        j                  }t          |j        |          S t          |t                    rU| j        j                            | j                  }d g|j        z  }t          |j        t!          g |j        |R            S  |j        | j                  }t'          |j        t+          | |                    S N)r   mesh_mapping)r1  r/  )r  r@   r6  r  r  r  r   
NoShardingrM  ShardingSpecr/  r   r  r1  r   r   r   r   r   r  r   r  )r  r   r  trailing_shardingphys_sharding_spectrailing_spechloss          r)   make_key_array_phys_shardingr    s_   x(( HO(L)) Hz 66tzBBH'2445E'4G8)2G5FGG+8: : :  0&8: : : :(M** 	Hz 66tzBBHFX]*M5x}5}5557 7 7 )8(33D#%:4%F%FH H Hr(   c                "    t          | |          S r?   )r  )r  r   s     r)   physical_shardingr  /  s    	%dH	5	55r(   c                   | j         j                            | j                   }|                    | j        |j        z             }t          |          \  }}|dk    rg n|g}|                                                                }|d |j                  |z   }||_        t          |j
        t          j                            |                    S rk   )r  r  r  r   rM  r   r  r  r  r  r   r}   r   r|  )	r  phys_shardingr  phys_hlo_shardingr  r  r  logical_op_shardingr  s	            r)   get_logical_gspmd_shardingr  4  s    Z44TZ@@(#88
i(-! !56GHH*l""22&)2244::<<?X]N?#f,#360	}7~001DEE
G 
G Gr(   c                <   t          | t                    rd S t          j        |          } | j        |j                  }t          |          \  }}|j        |j        z
  }t          d || d          D                       st          d|  d| d|           d S )Nc              3  "   K   | ]
}|d k    V  dS )rT   Nr'   )rD   rg   s     r)   rF   z1check_replicated_trailing_dims.<locals>.<genexpr>H  s&      ==Q!V======r(   zIThe trailing dims of extended dtypes should be replicated. Got sharding: z, partitions: z, num_trailing_dims: )	r@   r6  r   physical_avalr   rM  r   rJ   r%  )r   r  	phys_avalhlo_sr  r   num_trailing_dimss          r)   check_replicated_trailing_dimsr  A  s    ,'' 
F &&)
'(
'	
7
7%*511-*anty0	==Z):(:(;(;<===	=	= 3
	2	2 	2.8	2 	2/	2 	23 3 33 3r(   c                   t          ||            t          |          r|S t          |t                    rr| j        j                            | j                  }t          j        |j	        j
        d |j                  |j	        j                  }t          |j        |          S t          |t                    r%t          | |          }t!          ||j                  S t          | |          S r  )r  r  r@   r6  r  r  r  r   r  r/  r   rM  r  r1  r   r  !_gspmd_to_named_sharding_via_meshr   )r  r  r  logical_sharding_spec
logical_gss        r)   logical_shardingr  N  s     555}-- ;-.. ;z 66tzBBH*7,5o~oF"0=? ? ?  5&;= = = =-// ;+D-@@J,M&( ( ( &dM:::r(   pspecPartitionSpec | Nonec                L    |t                      d }}t          | |||          S )N)rz   r"   )r   r   )r   r  r5   r"   s       r)   create_mesh_pspec_shardingr  c  s6     ]'//4<E	tU,#.
0 
0 
0 0r(   out_sc                    t          | j        |          d         }t          ||                                || j                  S rs   )r  r  r  r1   r"   )r  r   r5   s      r)   r  r  m  sL    *4! !!"$,	#
L++--|
 
 r(   r  )rP   rQ   r;   rR   r  r+  r?   )rp  rq  rr  rs  r"   r   r;   rt  )r   r
  )r  r  r;   r  rE  )r  rq  r   r   r;   r  )r[  rA   )r  r  r)  r<   r  r<   r;   r  )r   r  r  rQ   r;   r  )r  r  r)  r<   rP   rQ   r;   r<   )r  rc   r;   rc   )r   r  r;   r   )r   r  r;   r  )r   r  )r;   r  )NN)r   r   r  r  r"   r   r;   r   )r  r  r   r   r;   r   )r
__future__r   rW   r   collections.abcr   r   dataclassesenumr  r^  r  typingr   r   r	   r
   jax._srcr   r   rb  r   r   r   r   r   jax._src.libr   r}   r   jax._src.op_shardingsr   r   r   jax._src.partition_specr   jax._src.utilr   r   r   r   numpyr=  rL   r<   rQ   r  rA   r  r   rS  XLACompatibleSharding	dataclassr    cacher:   rO   ra   r   r   r  r  r4  r6  r  r  rt  ry  r  r  r  r  UNSPECIFIEDr  r	  MeshAxisNamer
  r  r  ry   r#  IntEnumr  r   rH  r
  r   preprocess_with_manualrN  rT  ri  ro  ry  r{  r  r  r  r  r  r  r0   r  r  r  r  r  r  r  r  r  r  r  r  r  r'   r(   r)   <module>r     s	   # " " " " "     # # # # # # - - - - - - - -               / / / / / / / / / / / /       % % % % % %       # # # # # #                   ) ) ) ) ) ) . . . . . .Q Q Q Q Q Q Q Q Q Q 1 1 1 1 1 1 K K K K K K K K K K K K     	c3h	eSjFCK(  ) d###       $# Su555- - 65-R R R R T666   76& T666:F :F :F 76:Fz r  U< U< U< U< U<H% U< U< ! U<p Su555$ $ 65$ r&''I I I I I8, I I ('IX T6664 4 4 764 rW W W W W8$ W W  Wz #    8 T666) ) ) 76)*z; z; z; z; z;* z; z; z;z% % % % % % % %: r  Y( Y( Y( Y( Y(H% Y( Y( ! Y(x                   ) ) )) ) )  <,-"'d<L(L"M $ $ $ $$	; 	; 	; 	;. .     t|   P" P" P" P" P" P" P" P"d" " " " "': " " ": 7@ikk      3% 5:8 8 8 82[ [ [&    j    d###? ? ? ? ? ? ? $#?6 d###       $# d###0 0 0 0 0 0 0 $#02< < <( ( (4! ! !F  8&Q &Q &Q &QR   
> > > > >j > > >pM pM pM pMf8 8 8 8v(( (( (( ((V4 4 4 4R R R R
H H H.6 6 6 6
G G G3 3 3 3; ; ; ;* CG"0 0 0 0 0     r(   