
    VpfY                    0   d Z ddlmZ ddlZddlZddlmZ ddl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 ddlm	Z ddlmZ ddlmZ ddlm Z! ej"        Z"ej#        e$cZ$Z%ej&        e'cZ'Z( e
j)        d          Z*d Z+e*j,        d             Z-dedZ. ej/        e*e.            e
j)        d          Z0dfdZ1e0j,        d             Z2dedZ3 ej/        e0e3            e
j)        d          Z4ddddgd#Z5e4j,        d$             Z6ded%Z7 ej/        e4e7            G d& d'ej8                  Z9dhd(Z: e
j)        d)          Z;d*e;_<        d+ Z=e;j,        d,             Z> e
j)        d-          Z?d.e?_<        	 dide9j@        dd0djd8ZAe?j,        dkd9            ZBdld@ZCeCe
jD        e?<    e
j)        dA          ZEd.eE_<        didmdCZFeEj,        dD             ZGdldEZHeHe
jD        eE<   ejI         G dF dG                      ZJ e
j)        dH          ZKd.eK_<        eKjL        dI             ZMdldJZNeNe
jD        eK<   dK ZO  ejP        eK          eO            e
j)        dL          ZQd.eQ_<        eQj,        dM             ZRdldNZSeSe
jD        eQ<   dO ZT  ejP        eQ          eT           dP ZUdQ ZVdR ZWe9j@        fdkdSZXe9j@        fdkdTZY e
j)        d3          ZZeZj,        dU             Z[eZj\        Z] e
j)        dV          Z^e^j,        dW             Z_dX Z` e
j)        dY          Zad.ea_<        eaj,        dZ             Zbd[ Zc e
j)        d\          Zdd.ed_<        edj,        d]             ZedndaZf e
j)        db          Zgegj,        dc             Zedd ZhdS )oz<Module for Pallas:TPU-specific JAX primitives and functions.    )annotationsN)Any)core)dtypes)pretty_printer)state)	tree_util)util)indexing)
primitives)mlir)	discharge)	DTypeLikerepeatc                <    t                               | ||          S )N)repeatsaxis)repeat_pbind)xr   r   s      a/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/_src/pallas/mosaic/primitives.pyr   r   -   s    	q'	5	55    c               ~    t          | j                  }||xx         |z  cc<   t          j        || j                  S N)listshapejax_coreShapedArraydtype)r   r   r   r   s       r   _repeat_abstract_evalr    0   s:    
qw--%++++++		eQW	-	--r   ctxmlir.LoweringRuleContextc               R    fd} t          j        |d          | |          S )Nc                0    t          j        |           S r   )jnpr   )r   r   r   s    r   _repeatz&_repeat_lowering_rule.<locals>._repeat8   s    :a$'''r   Fmultiple_resultsr   	lower_fun)r!   r   r   r   r&   s     `` r   _repeat_lowering_ruler+   7   sC    ( ( ( ( ( (	8%	8	8	8a	@	@@r   bitcasttyr   c                   t          j        |          }t          | j                  dk     rt	          d          | j        d         | j        j        z  |j        z  rt	          d          t                              | |          S )N   zNot implemented: bitcast 1DzJNot implemented: the 2nd minor dim can not be perfectly packed or unpacked)r-   )	r   canonicalize_dtypelenr   
ValueErrorr   itemsize	bitcast_pr   )r   r-   s     r   r,   r,   @   s     $$"\\A
2
3
33WR[17##bk1 
	   
b	!	!!r   c                   t          | j                  }|d         | j        j        z  |j        z  |d<   t	          j        ||          S )Nr0   )r   r   r   r4   r   r   )r   r-   r   s      r   _bitcast_abstract_evalr7   L   s@    
qw--%Bi!'**bk9%)		eR	(	((r   c               N    fd} t          j        |d          | |          S )Nc                   | j         j        j        k     rk| j        ^ }}}j        | j         j        z  } | j        g |||z  ||R  } t	          j        | dd          } t          j                            |           S | j         j        j        k    rTt          j                            |           }|j        ^ }}}} t	          j        |dd          j        g |||z  |R  S t          j                            |           S )Nr0   )	r   r4   r   reshaper%   swapaxesjaxlaxbitcast_convert_type)r   leadingmnpackingyr-   s         r   _bitcastz(_bitcast_lowering_rule.<locals>._bitcastT   s   w"+%%wnw1qw//g
!)
7W
7a7l
7G
7Q
7
7
7a
,q"b
!
!aW))!R000w"+%%
'
&
&q"
-
-a !w1g,S\!R$$,FgFq7{FAFFFF7''2...r   Fr'   r)   )r!   r   r-   rE   s     ` r   _bitcast_lowering_rulerF   S   s?    / / / / / 
:5	9	9	9#q	A	AAr   roll)stridestride_axisr   intrH   
int | NonerI   c                  t          |t                    r|dk     rt          d          |dk     s|t          | j                  k    rt          d          |d u |d u k    rt          d          |Y|W|dk     rt          d          |dk     s|t          | j                  k    rt          d          ||k    rt          d          t
                              | ||||          S )	Nr   zshift must be non-negative.zaxis is out of range.z5stride and stride_axis must be both specified or not.zstride must be non-negative.zstride_axis is out of rangez,expected axis and stride_axis are different.)r   rH   rI   )
isinstancerJ   r3   r2   r   roll_pr   )r   shiftr   rH   rI   s        r   rG   rG   i   s    s 4		
2
3
33	AXXQW%%
,
-
--n+-..
L
M
MMK3zz5666Q+QW554555{EFFF	T&k 
 
 
 r   c                ,    ~t          j        |           S r   )r   raise_to_shaped)r   rO   _s      r   _roll_abstract_evalrS      s    		!!	$	$$r   c               X    fd} t          j        |d          | ||          S )Nc                    t          j        |           S fdt          t          j        | | j                                     D             }t          j        |          S )Nc                N    g | ]!\  }}t          j        ||z  z             "S  )r%   rG   ).0ixsr   rO   rH   s      r   
<listcomp>z6_roll_lowering_rule.<locals>._roll.<locals>.<listcomp>   sD       Ar 	UQZ'..  r   )r%   rG   	enumeratesplitr   concatenate)r   rO   outputsr   rH   rI   s    ` r   _rollz"_roll_lowering_rule.<locals>._roll   s    ~Xa%%%     syAGK,@+NNOO  G ?7K000r   Fr'   r)   )r!   r   rO   r   rH   rI   r`   s      ``` r   _roll_lowering_rulera      sM    1 1 1 1 1 1 1 
7	6	6	6sAu	E	EEr   c                      e Zd ZdZdZdS )DeviceIdTypemeshlogicalN)__name__
__module____qualname__MESHLOGICALrW   r   r   rc   rc      s        	$'''r   rc   c                |   |t           j        t           j        h}t          | t          j                  st          d| d|            | j        }|r|d                                         }|rt          d| d|           | j	        t          fd|D                       st          d| d|           d S )NzCannot z on a non-semaphore Ref: r:   z on a non-()-shaped semaphore: c              3  B   K   | ]}t          j        |          V  d S r   )r%   
issubdtype)rX   sem_type	sem_dtypes     r   	<genexpr>z"check_sem_avals.<locals>.<genexpr>   sC        
 
nY))     r   zMust z$ semaphores of the following types: )tpu_core	semaphorebarrier_semaphorerM   r   AbstractRefr3   r   get_indexer_shaper   any)sem_avalsem_indexers_avalsnameallowed_semaphore_types	sem_shapero   s        @r   check_sem_avalsr|      s   $'183MN	He/	0	0 J
HtHHhHH
I
IIn) ;"2&88::I Q
OtOOIOO
P
PPn)	    -   
 
  	& 	& 	&#	& 	&  	 r   semaphore_readFc                ~    t          |           \  }}||g}t          j        |          \  }}t          j        |d|iS )N	args_tree)_get_ref_and_indexersr	   tree_flattensemaphore_read_pr   )sem_or_viewrefindexersargs	flat_argsr   s         r   r}   r}      sG    '44-#x
x$"/55)Y				?Y	?	??r   c                    t          j        | |          \  }}t          ||dt          j        t          j        t          j        h           t          j        dt          j
        d                    S )Nread)rz   rW   int32)r	   tree_unflattenr|   rq   dma_semaphorerr   rs   r   r   r%   r   )r   avalsrw   rx   s       r   _semaphore_read_abstract_evalr      sp    
 "+!9)U!K!K(

 ("4h6P	    
	b#)G"4"4	5	55r   semaphore_signalT   )	device_iddevice_id_type
core_indexincint | jax.Arrayr   4int | jax.Array | None | tuple[int | jax.Array, ...]r   r   int | jax.Array | Nonec                   t          |           \  }}t          j        |t          j                  }|||||g}t	          j        |          \  }}	t          j        ||	|d d S )Nr   )r   r   )r   r%   asarrayr   r	   r   semaphore_signal_pr   )
r   r   r   r   r   r   r   r   r   r   s
             r   r   r      s}     (44-#xCsy)))#
xi	4$"/55)Y#     r   c                F   ~t          j        | |          \  }}}}}t          ||d           |j        t	          j        d          k    rt          d          |Et          j        |          }|D ].}	|	j        t	          j        d          k    rt          d          /g S )Nsignalr   zMust signal an int32 value.z$`device_id`s must be an int32 value.)r	   r   r|   r   r%   r3   tree_leaves)
r   r   r   rw   rx   
value_avaldevice_id_avalscore_index_avaldevice_id_flat_avalsavals
             r   _semaphore_signal_abstract_evalr      s     y%00 M(
O_ (.9997++++
2
3
33 $0AA$ A A	sy))	)	)?@@@ 
*	)r   eqnjax_core.JaxprEqncontextjax_core.JaxprPpContextsettingsjax_core.JaxprPpSettingsc                2   ~| j         }| j        d         }t          j        ||          \  }}}}}	t	          j        t	          j        d          t	          j        d          t          j        |||          t	          j        d          t	          j        t          j
        ||                    g          }
|t          j        |          }|s|
S t	          j        t          j
        |d         |                    g}|dd          D ]c}|                    t	          j        d                     |                    t	          j        t          j
        ||                               dt	          j        |
t	          j        |          g          }
|
S )Nr   r    r   r   )invarsparamsr	   r   ppconcattextsppp_ref_indexersr   pp_varr   append)r   r   r   r   treesemsem_indexersvalue
device_idsrR   outflat_device_idsdevice_ids_ppr   s                 r   _semaphore_signal_pp_eqnr      ss    :&	K	 $ tV,,	
	g !!gcll#|44gcllghoeW--.. 	 	# +J77O jWX__Q-?IIJJKM$QRR( I I	273<<(((278?9g#F#FGGHHHH
)S")M223
4
4C	*r   semaphore_waitdecc                    t          |           \  }}t          j        |t          j                  }|||g}t	          j        |          \  }}t          j        |d|i d S )Nr   r   )r   r%   r   r   r	   r   semaphore_wait_pr   )r   r   r   r   r   r   r   s          r   r   r     sf    '44-#xCsy)))#
x	$"/55)Y8i88888r   c                    t          j        | |          \  }}}t          ||d           |j        t	          j        d          k    rt          d          g S )Nwaitr   zMust wait an int32 value.)r	   r   r|   r   r%   r3   )r   r   rw   rx   r   s        r   _semaphore_wait_abstract_evalr   &  s\    -6-EiQV-W-W*(
(.7777++++
0
1
11	)r   c                n   ~| j         }| j        d         }t          j        ||          \  }}}t	          j        t	          j        d          t	          j        d          t          j        |||          t	          j        d          t	          j        t          j
        ||                    g          S )Nr   r   r   )r   r   r	   r   r   r   r   r   r   r   r   )r   r   r   r   r   r   r   r   s           r   _semaphore_wait_pp_eqnr   .  s     :&	K	 $
 tV,,			ggcll#|44gcllghoeW--.. 
 
 r   c                      e Zd ZU ded<   ded<   ded<   ded<   ded<   ded	<   d
ed<   ded<   d
ed<   ej        Zded<   d Zed             Z	d Z
d Zd Zd ZdS )AsyncCopyDescriptorr   src_refztuple[indexing.NDIndexer, ...]src_indexersdst_refdst_indexersr   dst_semdst_sem_indexersr   src_semz%tuple[indexing.NDIndexer, ...] | Nonesrc_sem_indexersr   rc   r   c                J    | j         d u | j        d u z  rt          d          d S )Nz<Either both or neither `src_sem` and `device_id` can be set.)r   r   r3   selfs    r   __post_init__z!AsyncCopyDescriptor.__post_init__P  s>    4!78 & % & & && &r   c                    | j         d uS r   )r   r   s    r   	is_remotezAsyncCopyDescriptor.is_remoteU  s    <t##r   c                    t          j        | j        | j        | j        | j        | j        | j        | j        | j	        | j
        f	          \  }}t          j        ||| j        d d S Nr   r   )r	   r   r   r   r   r   r   r   r   r   r   dma_start_pr   r   )r   r   r   s      r   startzAsyncCopyDescriptor.startY  sq    ,
. 
 
OIt id4;NOOOOOOr   c                d    | j         r|                                  |                                  d S r   )r   	wait_send	wait_recvr   s    r   r   zAsyncCopyDescriptor.waitg  s2    ~ 
nnNNr   c                    t          j        | j        | j        | j        | j        f          \  }}t          j        ||| j        d d S r   )	r	   r   r   r   r   r   
dma_wait_pr   r   r   	wait_argsr   s      r   r   zAsyncCopyDescriptor.wait_recvl  s]    ,	t,dlD<MN OIt O	d.A     r   c                    | j         st          d          t          j        | j        | j        | j        | j        f          \  }}t          j	        ||| j
        d d S )Nz#Cannot `wait_send` on a local copy.r   )r   r3   r	   r   r   r   r   r   r   r   r   r   s      r   r   zAsyncCopyDescriptor.wait_sendt  sv    > ><===,	t,dlD<MN OIt O	d.A     r   N)rf   rg   rh   __annotations__rc   ri   r   r   propertyr   r   r   r   r   rW   r   r   r   r   C  s         ,,,....,,,....2222!!!!9999####!-!2.2222& & &
 $ $ 8$P P P  
      r   r   	dma_startc           	        t          j        | |          \	  }}}}}}}	}
}|j        }|r|d                                         }|rt	          d|           |	7|	j        }|
r|
d                                         }|rt	          d|           t          t          j        |                    }g t          j        d          t          j	        |dz             hfS )Nr:   z,Cannot signal on a non-()-shaped semaphore: r   r   )
r	   r   r   ru   r3   r2   r   r   
ReadEffectWriteEffect)r   r   r   src_ref_avalsrc_indexers_avalsdst_ref_avaldst_indexers_avalsdst_sem_avaldst_sem_indexers_avalssrc_sem_avalsrc_sem_indexers_avalsdevice_id_avaldst_sem_shapesrc_sem_shapen_src_indexerss                  r   _dma_start_abstract_evalr     s    tT**
$- C*2.@@BBM 
F}FF    &M E,R0BBDDm 
H
H
H   y,-?@@AA.	eq!!5#4^a5G#H#HI	IIr   c                   | j         }| j        d         }t          j        ||          \	  }}}}}	}
}}}~|s|rt	          j        | ||          S t          j        t          j        d          t          j        d          t          j
        |||          t          j        d          t          j
        |||          t          j        d          t          j
        ||	|
          g          S )Nr   r   r   z -> )r   r   r	   r   r   _pp_eqnr   r   r   r   r   )r   r   r   r   r   r   r   r   r   r   r   r   r   r   s                 r   _dma_start_pp_eqnr     s     :&	F	$ tV,,
 4	 4C(333	gkgcll'<88gfoo'<88gcll'+;<< 
 
 r   c          	        t          j        ||          \	  }}}}}	}
}}}t          j        ||           ^}}}}}~~	~
|d uo|d u}|r |t          j        k    rt	          d          n|J |J |J t          t          j        |                    }t          t          j        |                    }|rt          j        ||          }n|}|rt          j
        j        j        }t          d |D                       }t          d |D                       }t          |          dk    rt	          d          |d         }t          j                            |          }t          j                            ||          |k    }t%          j        |d          }t          j                            ||          }t          j                            ||dd	          }|rt          j        |||          \  }}n|}d
}|d
|z  z  }||fz  }|d
|z  z  }|d
z  }|r|dz  }t          |          t          |           k    s+J t          |          |f dt          |                        |g fS )Nz"Mesh device_id_type not supported.c              3  $   K   | ]}|j         V  d S r   )ry   )rX   frames     r   rp   z+dma_start_discharge_rule.<locals>.<genexpr>  s$      88euz888888r   c              3     K   | ]}||V  	d S r   rW   )rX   ry   s     r   rp   z+dma_start_discharge_rule.<locals>.<genexpr>  s'      PPt?O?O?O?O?OPPr   r   zFSharding with more than one named axis not implemented in dma_start_p.r   )r   F)r   keepdimsr   )NNz != )r	   r   rc   ri   NotImplementedErrorr2   r   state_dischargeindex_arrayr   thread_local_statetrace_stateaxis_envtupler=   r>   
axis_index
all_gatherr%   argmaxdynamic_index_in_dimindex_swap_array)in_avals	out_avalsr   r   r   r   r   r   r   r   r   r   r   r   rR   r   r   r   num_src_index_valsnum_dst_index_valsupdatesr  
axis_namesnonempty_axis_names
shard_axismy_axiswho_copy_to_meindexglobal_updatesnew_dst	new_avalss                                  r   dma_start_discharge_ruler    s    tT**
 tX..*T!;it&;) *** DEEE + ???###901CDDEE901CDDEE )'<@@GGG 7 *6?H88x88888JPPPPPPP
!## !> ? ? ?$Q'Jg  ,,G W''	:>>'IN J~A...EW''<<Ng**A + 7 7G   1w JAww G )w+++)z)w+++)w) I
i..
h--   ^^Y7KKCMMKK  	Br   dma_waitc                    ~~ ~g S r   rW   )r   r   r   s      r   _dma_wait_abstract_evalr    s    
D.	)r   c                N   ~| j         }| j        d         }t          j        ||          \  }}}}t	          j        t	          j        d          t	          j        d          t          j        |||          t	          j        d          t          j        |||          g          S )Nr   r  r   )	r   r   r	   r   r   r   r   r   r   )	r   r   r   r   r   r   r   r   r   s	            r   _dma_wait_pp_eqnr    s     :&	F	$%.%=dF%K%K"#|S(	gjgcll#x00gcll#|44 
 
 r   c               2    ~~~~dt          |           z  g fS )Nr   )r2   )r  r  r   r   r   s        r   dma_wait_discharge_ruler!  -  s!    t^	3x==	 "	$$r   c                Z    t          | t          j                  r| j        | j        fS | dfS NrW   )rM   r   RefViewr   r   )r   s    r   r   r   4  s.    U]## !7CL  	b.r   c                    t          |           \  } }t          |          \  }}t          |          \  }}t          | |||||dddt          j        
  
        S )-Issues a DMA copying from src_ref to dst_ref.N)r   r   rc   ri   )r   r   r   r   r   r   s         r   make_async_copyr'  9  sc    /88'</88'<+C00#|	WlG\ ,dD).
0 
0 0r   c                P    t          | ||          }|                                 |S )r&  )r'  r   )r   r   r   copy_descriptors       r   
async_copyr*  B  s+    #GWc::/	r   c                    t          |           \  } }t          |          \  }}t          |          \  }}t          |          \  }}	t          | |||||	||||
  
        S )a  Creates a description of a remote copy operation.

  Copies data from src_ref on the current device to dst_ref on the device
  specified by device_id. Both semaphores should be waited on using the
  descriptor on both source and target devices.

  Note that device_id can also refer to the current device.

  Args:
    src_ref: The source Reference.
    dst_ref: The destination Reference.
    send_sem: The semaphore on the source device.
    recv_sem: The semaphore on the destination device.
    device_id: The device id of the destination device.
    device_id_type: The type of the device id.
  Returns:
    An AsyncCopyDescriptor.
  )r   )r   r   )
r   r   send_semrecv_semr   r   r   send_sem_indexersr   recv_sem_indexerss
             r   make_async_remote_copyr0  H  s    ( 088'< 5h ? ?(/88'< 5h ? ?(	|WlH>O!9^
M 
M 
M Mr   c                V    t          | |||||          }|                                 |S r   )r0  r   )r   r   r,  r-  r   r   r)  s          r   async_remote_copyr2  d  s8    *7GXx+4nF F/	r   c                 P    t          j        dt          j        d                    S )NrW   r   r   r   r%   r   rW   r   r   _device_id_abstract_evalr5  m  s    		b#)G"4"4	5	55r   get_barrier_semaphorec                     t          j        t          j        dt	          j                              t          j        j                  S r#  )pl_coreAbstractMemoryRefr   r   rq   BarrierSemaphoreTyTPUMemorySpace	SEMAPHORErW   r   r   $_get_barrier_semaphore_abstract_evalr=  u  s9    		"2x:<<=='
 
 r   c                 4    t                                           S )a  Returns a barrier semaphore.

  This function returns a barrier semaphore based on the collective_id of the
  current pallas kernel.

  It's very important that the semaphore is wait-ed back down to 0, or else the
  semaphores will become corrupted.

  It's also very important that the collective_id is different for each pallas
  kernel with communication. E.g. if you have two pallas kernels, one that syncs
  across the X axis of the device mesh and the second that syncs across the Y
  axis, they must have different collective_ids.
  However it is legal for two kernels that perform the same synchronization
  pattern (e.g. only communicating with neighbours on the same mesh axis)
  to share a collective_id. However, if in doubt, prefer not sharing
  collective_ids, as doing so incorrectly can lead to silent data corruption or
  crashes.
  Note that re-using the same collective_id doesn't guarantee that the same
  semaphore is provided by XLA.
  )get_barrier_semaphore_pr   rW   r   r   r6  r6  |  s    * 
!	%	%	'	''r   delayc                    ~ g S r   rW   nanoss    r   _delay_abstract_evalrD    s
    	)r   c                :    t                               |            dS )z;Delays vector execution for the given number of nanosconds.N)delay_pr   rB  s    r   r@  r@    s    	,,ur   	prng_seedc                     g S r   rW   )rR   s    r   rR   rR     s    	)r   seedsreturnNonec                 "    t          j        |   dS )zSets the seed for PRNG.

  Args:
    seeds: One or more integer seeds for setting the PRNG seed. If
      more than one seed is passed in, the seed material will be
      mixed before setting the internal PRNG state.
  N)prng_seed_pr   )rI  s    r   rG  rG    s     Er   prng_random_bitsc                P    t          j        | t          j        d                    S )Nr   r4  r   s    r   rR   rR     s    		eSYw%7%7	8	88r   c                8    t                               |           S )NrP  )prng_random_bits_pr   rP  s    r   rN  rN    s    		 	 u	 	-	--r   )r!   r"   )r-   r   )r   rJ   rH   rK   rI   rK   r   )r   )r   r   r   r   r   rc   r   r   )r   rc   )r   r   r   r   r   r   )r   r   )rI  r   rJ  rK  )i__doc__
__future__r   dataclassesenumtypingr   r=   jax._srcr   r   r   r   r   r   r	   r
   jax._src.stater   r   r   jax._src.interpretersr   jax._src.pallasr8  jax._src.pallas.mosaicrq   r   r  jax._src.typingr   	jax.numpynumpyr%   Slicesafe_mapmap
unsafe_mapsafe_zipzip
unsafe_zip	Primitiver   r   def_abstract_evalr    r+   register_loweringr5   r,   r7   rF   rN   rG   rS   ra   Enumrc   r|   r   r(   r}   r   r   ri   r   r   r   pp_eqn_rulesr   r   r   r   	dataclassr   r   def_effectful_abstract_evalr   r   r  register_discharge_ruler   r  r  r!  r   r'  r*  r0  r2  device_id_pr5  r   r   r?  r=  r6  rF  rD  r@  rM  rR   rG  rR  rN  rW   r   r   <module>rp     s   C B " " " " " "            



 % % % % % %       ) ) ) ) ) )                   # # # # # # + + + + + + & & & & & & + + + + + + 3 3 3 3 3 3 7 7 7 7 7 7 % % % % % %      -Z-Z8h''6 6 6 
. . .A A A A  x!6 7 7 7Hy))		" 	" 	" 	" ) ) )B B B B"  y"8 9 9 9		F	#	# "     4 % % %
F F F F  v2 3 3 3    49   
   * &8%&677 $)  !@ @ @ #6 6 $#6  (X'(:;; &*  #
  GK#/#4)-     & %   &%(   < -E ( )%8%&677 $(  !9 9 9 9 9 #  $#   $ +A & ' 8 8 8 8 8 8 8 8v !h --# (J J )(J<   < &7 k "N N N` 5 ' ' 4 45M N N N  X
++
"
       %5 j !% % %
 4 ' '
 3 34K L L L  
0 0 0   ;G:KM M M M M: 6B5F     !h --6 6 6 	,(,-DEE *  +*( ( (. (
W
%
%  	  
   !h --#       (X'   %9 9 &%9. . . . .r   