
    ciF                         d dl mZmZmZmZmZ d dl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  ed	      Zefd
edeg ef   fdZ G d d      Z G d dee         Zy)    )SequenceListTypeVarTupleCallableN)TritonSemantic   )_core)
AutoLayoutDistributedLayoutSliceLayout)GluonOpBuilder)flatten_values_to_irunflatten_ir_valuesTensorTycondmsg_fnc                 "    | s | |             y N )r   r   categorys      f/var/www/html/engine/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/_semantic.py_checkr      s    vx       c                   $    e Zd ZdefdZd Zd Zy)GluonCallerContext	num_warpsc                     || _         y r   r   )selfr   s     r   __init__zGluonCallerContext.__init__   s	    "r   c                      d| j                    S )N_NWr   )r    s    r   manglezGluonCallerContext.mangle   s    T^^$%%r   c                 Z    |j                  d|j                  | j                               y )Nzttg.num-warps)set_attrget_int32_attrr   )r    fnbuilders      r   initialize_calleez$GluonCallerContext.initialize_callee   s    
OW%;%;DNN%KLr   N)__name__
__module____qualname__intr!   r$   r*   r   r   r   r   r      s    ## #&Mr   r   c            
           e Zd ZU ej                  ZeZeed<   defdZd Z	d Z
dee   dee   fdZded	ed
efdZdeded
ef fdZded
eeef   f fdZdedee   d
ef fdZdedee   d
efdZdeded
ef fdZ fdZdedee   def fdZd Zd Zd5dZd Zd Zd  Zd! Zd" Z d# Z!d$ Z"d% Z#d& Z$d' Z%d( Z&e'd)        Z(d*e)e   d	ed+ed
eed,f   fd-Z*d*e)e   d	ed
eed,f   fd.Z+ded/ed0ed
efd1Z,d2e)e   d3e)e   fd4Z- xZ.S )6GluonSemanticr)   c                     || _         y r   )r)   )r    r)   s     r   r!   zGluonSemantic.__init__#   s	    r   c                     |g k(  r|}n0t        j                  ||| j                  j                  |            }| j	                  ||      S r   )ttgldistributed_typer)   get_gluon_layout_from_tensortensor)r    handle	scalar_tyshapetys        r   _wrap_handle_infer_layoutz'GluonSemantic._wrap_handle_infer_layout&   sC    B;B&&y%9b9bci9jkB{{62&&r   c                 x    | j                  |j                  |j                  j                  |j                        S r   )r;   r7   typescalarr9   )r    r6   s     r   _wrap_tensor_infer_layoutz'GluonSemantic._wrap_tensor_infer_layout-   s,    --fmmV[[=O=OQWQ]Q]^^r   	lhs_shape	rhs_shapec                 X   t        |      t        |      k7  rt        d| d|       g }t        |      D ]q  \  }}||   }|dk(  r|j                  |       "|dk(  s||k(  r|j                  |       >t        dt	        |      z   dz   t	        |      z   dz   t	        |      z          |S )N!Cannot broadcast, rank mismatch: , r	   z?Cannot make_shape_compatible: incompatible dimensions at index : z and )len
ValueError	enumerateappendstr)r    r@   rA   	ret_shapeileftrights          r   _broadcast_shapeszGluonSemantic._broadcast_shapes0   s    y>S^+@2i[YZZ	 + 	aGAtaLEqy  '1*%4-  &  "-/21v"68<"=?B4y"IKR"SUXY^U_"` a a	a r   inputaxisreturnc                    j                   D cg c]  }t        j                  |       }}|j                  d       dk  rt	        j                         z  t        t        j                  t        j                        fd       j                  j                  t        t        t        t        f      fd       t        t        t              xs j                  k(  fd       | j                  j                  j                        }| j!                  |j                  j"                  |      S c c}w )Nr	   r   c                  "    d j                   S Nz=expected expand_dims input to be a distributed_type but got: r=   rP   s   r   <lambda>z+GluonSemantic.expand_dims.<locals>.<lambda>H       VW\WaWaVde r   c                      d  S )Nz;expected expand_dims input to have a SliceLayout, but got: r   )layouts   r   rX   z+GluonSemantic.expand_dims.<locals>.<lambda>K   s    TU[T\] r   c                  (    d  dj                    S )Nz7expected expand_dims input layout to be sliced in axis z	 but got )dim)rQ   r[   s   r   rX   z+GluonSemantic.expand_dims.<locals>.<lambda>N   s    MdVS\]c]g]g\hi r   )r9   r3   _unwrap_if_constexprinsertrF   r   
isinstancer=   r4   r[   r   r   r]   r)   create_expand_dimsr7   r;   r>   )r    rP   rQ   x	dst_shaper7   r[   s    ``   @r   expand_dimszGluonSemantic.expand_dims@   s    ;@;;GaT..q1G	Gq!!8C$$Dz%**d&;&;<e	g""z&;
";<]	_vz*@fjjD.@i	k 00tD--fejj6G6GSS! Hs   Eabc                     | j                  ||      \  }}t        |j                  g k7  d       t        |   ||      }| j                  |      S )NzCannot join scalars in gluon)broadcast_impl_valuer   r9   superjoinr?   )r    re   rf   value	__class__s       r   rj   zGluonSemantic.joinS   sM    ((A.1qww"}<=Q"--e44r   c                 l    t         |   |      \  }}| j                  |      | j                  |      fS r   )ri   splitr?   )r    re   lhsrhsrl   s       r   rn   zGluonSemantic.splitY   s7    7=#S--c2D4R4RSV4WWWr   dimsc                 F    t         |   ||      }| j                  |      S r   )ri   permuter?   )r    rP   rq   rk   rl   s       r   rs   zGluonSemantic.permute]   s$    t,--e44r   r9   c                    t        t        j                  t        j                        fd       j                  j                         t        t              t              k(  fd       k(  rS t              D ]0  \  }}|   |k7  s|dk7  st        d|    d| d| d d 
       t        j                  j                  j                  j                  j                        }| j                  j                  j                  |j                  | j                              }| j                  ||      S )	Nc                  "    d j                   S rU   rV   rW   s   r   rX   z4GluonSemantic.broadcast_impl_shape.<locals>.<lambda>c   rY   r   c                      d d  S )NrC   rD   r   )r9   	src_shapes   r   rX   z4GluonSemantic.broadcast_impl_shape.<locals>.<lambda>e   s    7XYbXccefkel5m r   r	   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension rE   rD   )r   r`   r=   r3   r4   get_block_shapesrF   rH   rG   r>   r[   r)   create_broadcastr7   to_irr6   )r    rP   r9   rL   itemret_tyr7   rw   s    ``    @r   broadcast_impl_shapez"GluonSemantic.broadcast_impl_shapea   s-   z%**d&;&;<e	gJJ//1	s9~U+-mnIL + 	@GAtQx4DAI #VW\]^W_V` aCCG& I%%&Cr)Bug"? @ @	@
 &&uzz'8'8%ARARS..u||V\\$,,=WX{{66**r   ro   rp   c                 P  	 |j                   |j                   	j                         r	j                         st        
|   ||      S t	        t        t        j                        fd       t	        t        	t        j                        	fd       j                         }	j                         }| j                  ||      }t        j                  t              }t        	j                  t              }|r|s| j                  |	j                        }n_|r|s| j                  |j                        }n>j                  	j                  k7  r%t        dj                   d	j                         | j                  ||      }| j                  ||      }||fS )Nc                      d S )Nz@expected broadcast left input to be a distributed_type but got: r   )lhs_tys   r   rX   z4GluonSemantic.broadcast_impl_value.<locals>.<lambda>y   s    YZ`Ycd r   c                      d S )NzAexpected broadcast right input to be a distributed_type but got: r   )rhs_tys   r   rX   z4GluonSemantic.broadcast_impl_value.<locals>.<lambda>{   s    Z[aZde r   zLayout mismatch in broadcast: z vs )r=   is_blockri   rh   r   r`   r3   r4   rx   rO   r[   r   set_auto_layoutrG   r}   )r    ro   rp   r@   rA   rK   is_lhs_autois_rhs_autor   r   rl   s           @@r   rh   z"GluonSemantic.broadcast_impl_valueq   sX    (97/S99z&$"7"78d	fz&$"7"78e	g ++-	++-	**9i@	 
; 
;{&&sFMM:C&&sFMM:C]]fmm+=fmm_DQWQ^Q^P_`aa''Y7''Y7Cxr   c                     ||z
  g}|
t               }t        j                  t        j                  ||      }t        |   |||      S )N)r|   )r   r3   r4   int32ri   arange)r    startendr[   r9   r|   rl   s         r   r   zGluonSemantic.arange   sG    u>\F&&tzz5&Aw~eS~88r   rc   can_reorderc                 b    t        | d       t        | 	  |||      }| j                  |      S )Nz%can_reorder is not supported in gluon)r   ri   reshaper?   )r    rP   rc   r   rk   rl   s        r   r   zGluonSemantic.reshape   s3    ; GHy+>--e44r   c                     t        j                  |j                  ||      }| j                  j	                  |j                  | j                        |j                        }t        j                  ||      S r   )r3   r4   dtyper)   create_splatrz   r7   r6   )r    rk   r9   r[   r|   r7   s         r   splatzGluonSemantic.splat   sR    &&u{{E6B**6<<+Eu||T{{66**r   c                 d    | j                  ||      }|
t               }| j                  |||      S r   )make_scalarr   r   )r    r9   rk   r   r[   r>   s         r   fullzGluonSemantic.full   s3    !!%/>\Fzz&%00r   c                    |j                   t        t        t        j                        fd       t        j                  j
                  j                  |      }|j                  | j                        }|rB| j                  j                  ||j                        st        dj                   d| d      | j                  j                  ||j                        }t        j                  ||      S )Nc                      d S )Nz@expected convert_layout input to be a distributed_type but got: r   )r:   s   r   rX   z.GluonSemantic.convert_layout.<locals>.<lambda>   s    YZ\Y_` r   zlayout conversion from z to z is not trivial)r=   r   r`   r3   r4   
element_tyr9   rz   r)   is_convert_layout_trivialr7   	TypeErrorr[   create_convert_layoutr6   )r    rk   r[   assert_trivialr|   	ret_ty_irr7   r:   s          @r   convert_layoutzGluonSemantic.convert_layout   s    ZZz"d334`	b&&r}}bhhGLL.	$,,"H"HTYT`T`"a5bii[VHO\]]33Iu||L{{66**r   c                 P   t        j                  ||||      }|@| j                  j                  |j	                  | j                        |j
                        }n4| j                  j                  |j	                  | j                              }t        j                  |||||      S r   )r3   shared_memory_descriptor_typer)   create_local_allocrz   r7   shared_memory_descriptor)r    r   r9   r[   rk   r:   r7   s          r   allocate_sharedzGluonSemantic.allocate_shared   s    //
E65Q\\44RXXdll5KU\\ZF\\44RXXdll5KLF,,VZPUVVr   c                    t        j                  |j                  |j                  |      }| j                  j                  |j                  | j                        |j                        }t        j                  ||      S r   )	r3   r4   r   r9   r)   create_local_loadrz   r7   r6   )r    mem_descr[   r|   r7   s        r   shared_loadzGluonSemantic.shared_load   sV    &&x~~x~~vN//T\\0JHOO\{{66**r   c                 P   |j                   |j                   k(  s"J d|j                    d|j                    d       |j                  |j                  k(  s"J d|j                   d|j                   d       | j                  j                  |j                  |j                         y )Nzsource shape z and destination shape z must matchzsource dtype z and destination dtype )r9   r   r)   create_local_storer7   )r    r   rk   s      r   shared_storezGluonSemantic.shared_store   s    {{hnn,}ekk]Jabjbpbpaqq|.}},{{hnn,}ekk]Jabjbpbpaqq|.}},''Fr   c                 N    | j                   j                  |j                         y r   )r)   create_local_deallocr7   )r    r   s     r   shared_dealloczGluonSemantic.shared_dealloc   s    ))(//:r   c                    |j                   }t        |t              s
J d|        t        |j                  t              sJ d|j                   j                          | j
                  j                  |j                  | j
                        |j                        }t        j                  |j                  |j                  |      }| j                  ||      S )Nz9set_auto_layout must set to a distributed layout but got z4set_auto_layout input must have auto layout but got )r=   r`   r   r[   r   r)   create_set_auto_layout_to_irr7   r3   r4   r   r9   r6   )r    rk   r[   src_tyr7   res_tys         r   r   zGluonSemantic.set_auto_layout   s    &+- 	s0ijpiq.r	s -&--$& 	r)]^c^h^h^o^o]p'q	r &44V]]4<<5PRWR^R^_&&v'8'8&,,O{{66**r   c                    dg|j                   z  }|||<   t        |j                        }|||<   |j                  }t	        j
                  |j                  |||j                  j                        }| j                  }	|	j                  |j                  |	      |j                  |      }
t	        j                  |
fi |j                  S )Nr   )ranklistr9   r[   r3   r   r   r=   alloc_shaper)   create_memdesc_subslicerz   r7   r   __dict__)r    r   r   lengthr]   offsetsr9   r[   r:   r)   r7   s              r   memdesc_slicezGluonSemantic.memdesc_slice   s    #%X^^$c
//vx}}OhOhi,,00'1BHOOU\],,VCr{{CCr   c                    |j                   dd  }| j                  |      j                  }|j                  }t	        j
                  |j                  |||j                  j                        }| j                  }|j                  |j                  |      |j                  |      }t	        j                  |fi |j                  S )Nr	   )r9   	to_tensorr7   r[   r3   r   r   r=   r   r)   create_memdesc_indexrz   r   r   )r    r   indexr9   r[   r:   r)   r7   s           r   memdesc_indexzGluonSemantic.memdesc_index   s    qr"u%,,//vx}}OhOhi,,--bhhw.?RWX,,VCr{{CCr   c                 \   t        |      t        |j                        k(  s!J d|j                   dt        |       d       |D cg c]  }|j                  |    }}|j                  j                  }|d t        |      |j                  z
   }||D cg c]   }|t        |      |j                  z
  d  |   " c}z  }| j
                  j                  |j                  |      }| j
                  j                  |      }t        j                  ||j                  |||      S c c}w c c}w )Nzsource rank (z) and order length (z) must matchr   r9   r   r[   )rF   r9   r   r=   r   r)   create_memdesc_transr7   get_gluon_layout_from_memdescr3   r   r   )	r    r   orderrL   r9   r   new_alloc_shaper7   r[   s	            r   memdesc_transzGluonSemantic.memdesc_trans   s%   5zSNN  	i,X]]O;OPSTYPZ|[gh	i  -22q"22mm//%&Gs;'7(--'GHW\]RSKK(88==(H(IJ1M]]228??EJ;;FC,,VV[9HQWY 	Y 3 ^s   D$%D)c                    t        t        j                        t        j                  j                        k(  fd       | j                  j                  j                        }| j                  j                  |      }j                  j                  }t        |      j                  z
  }|d | t              z   }t        j                  |j                  ||      S )Nc                  (    d j                    d S )Nz)memdesc_reshape total elements mismatch: z -> )r9   )r   r9   s   r   rX   z/GluonSemantic.memdesc_reshape.<locals>.<lambda>   s    @ 'tE74 r   r   )r   mathprodr9   r)   create_memdesc_reshaper7   r   r=   r   rF   r   r   r3   r   r   )r    r   r9   r7   r[   r   
prefix_lenr   s    ``     r   memdesc_reshapezGluonSemantic.memdesc_reshape   s    IIe		(.. 995	
 44X__eL;;FCmm//%5
%kz2T%[@,,~~'
 	
r   c                     t        j                  ||||      }| j                  j                  |j	                  | j                        |j
                        }t        j                  |fi |j                  S r   )r3   r   r)   create_memdesc_reinterpretrz   r7   r   r   )r    r   r   r9   r[   r:   r7   s          r   memdesc_reinterpretz!GluonSemantic.memdesc_reinterpret  s\    //ufeL88$,,9OQYQ`Q`a,,VCr{{CCr   c                 ^    |rt        j                  |||      }n|}| j                  ||      S r   )r3   r4   r6   )r    rb   r8   rK   r[   r   s         r   wrap_tensorzGluonSemantic.wrap_tensor  s/    **9iHFF{{1f%%r   c                    | D ]3  t        t        j                  t        j                        fd       5 | D cg c]  }|j                  j
                   c}d   t        t        fddd  D              fd       y c c}w )Nc                  "    d j                   S )Nz#expected distributed_type but got: rV   )rb   s   r   rX   z2GluonSemantic._check_same_layout.<locals>.<lambda>  s    HklmlrlrkuFv r   r   c              3   (   K   | ]	  }|k(    y wr   r   ).0ll0s     r   	<genexpr>z3GluonSemantic._check_same_layout.<locals>.<genexpr>  s     0q170s   r	   c                      d  S )Nz3Expected inputs to have matching layouts, but got: r   )layoutss   r   rX   z2GluonSemantic._check_same_layout.<locals>.<lambda>  s    LWIV r   )r   r`   r=   r3   r4   r[   all)xsrb   r   r   s    `@@r   _check_same_layoutz GluonSemantic._check_same_layout  su     	xA:affd&;&;<>vw	x*,-Q166==-QZs0GABK00V	X .s    B
inputsreverse.c                     d   j                   j                  t              }| |cxk  r|k  sn J d| d| d       |dk  r||z  }D ]"  }|j                   j                  k(  rJ d         j                  j	                  D cg c]  }|j
                   c}||       |       j                         sJ t         fdt        t                    D              S c c}w )Nr   z
scan axis z must be < inputs rank ()z(all scan inputs must have the same shapec              3      K   | ]=  }j                  j                  |      |   j                  j                         ? y wr   r;   
get_resultr=   r>   )r   rL   r   scan_opr    r9   s     r   r   z1GluonSemantic.associative_scan.<locals>.<genexpr>,  sB      ) **7+=+=a+@&)..BWBWY^_)   AA)	r=   r9   rF   r)   create_scanr7   verifytuplerange)	r    r   rQ   region_builder_fnr   r   tr   r9   s	   ``     @@r   associative_scanzGluonSemantic.associative_scan  s    q	$$5zu#t#Wz$7OPTvUV%WW#!8DLD 	UA66<<5(T*TT(	U ,,**f+EAHH+EtWU'"~~ )3v;') ) 	)	 ,Fs   C.c                 V   	
 t        d ud        d   j                  j                  
t        
      t        dcxk  xr k  nc fd        j	                         t        
      D cg c]  \  }}|k7  s| c}}	t        
fdD              sJ d        j                  j                  D cg c]  }|j                   c}       |       j                         sJ t        	 fdt        t                    D              S c c}}w c c}w )Nc                       y)Nz*All-reduce is not yet implemented in gluonr   r   r   r   rX   z)GluonSemantic.reduction.<locals>.<lambda>1      r   r   c                      d d  S )Nz/expected reduction axis to be in the range [0, z
) but got r   )rQ   r   s   r   rX   z)GluonSemantic.reduction.<locals>.<lambda>5  s    +Z[_Z``jkojp)q r   c              3   P   K   | ]  }|j                   j                  k(    y wr   )r=   r9   )r   r   r9   s     r   r   z*GluonSemantic.reduction.<locals>.<genexpr>8  s     9Q166<<5(9s   #&z-all reduction inputs must have the same shapec              3      K   | ]=  }j                  j                  |      |   j                  j                         ? y wr   r   )r   rL   r   	reduce_oprK   r    s     r   r   z*GluonSemantic.reduction.<locals>.<genexpr>>  sB      ) **9+?+?+BF1INNDYDY[de)r   )r   r=   r9   rF   r   rH   r   r)   create_reducer7   r   r   r   )r    r   rQ   r   rL   sr   r   r   rK   r9   s   ```    @@@@r   	reductionzGluonSemantic.reduction0  s    t4!UVq	$$5zqD4!qr'#,U#3A41aqDyQA	9&99j;jj9LL..&/IQ/I4P	)$!!! )3v;') ) 	) B 0Js   8D D D&num_binsmaskc                     t        t        |j                        dk(  d        t        |j                  j	                         d        t        |d ud        |P| j                  ||      \  }}t        |j                  j                  j                         d        |j                  }|j                  | j                        }| j                  j                  |j                  |||      }| j                  |t        j                  |g|      S )Nr	   c                       y)Nz histogram only supports 1D inputr   r   r   r   rX   z)GluonSemantic.histogram.<locals>.<lambda>C  r   r   c                       y)Nz%histogram only supports integer inputr   r   r   r   rX   z)GluonSemantic.histogram.<locals>.<lambda>D  r   r   c                       y)Nz'histogram requires a destination layoutr   r   r   r   rX   z)GluonSemantic.histogram.<locals>.<lambda>E  r   r   c                       y)Nz"Mask must have boolean scalar typer   r   r   r   rX   z)GluonSemantic.histogram.<locals>.<lambda>H  r   r   )r   rF   r9   r   is_intrh   r=   r>   is_boolr7   r   r)   create_histogramr   r3   r   )r    rP   r   r   r[   layout_attrr7   s          r   	histogramzGluonSemantic.histogramB  s    s5;;1$&PQu{{!!#%TUvT!#TU33D%@KD%499##++-/[\;;DmmDLL1..u||Xt[Y

XJGGr   worker_num_warpsworker_num_regsc           	         t        |      }|t        |      k(  sJ d| dt        |       d       |t        |      k(  sJ d| dt        |       d       | j                  }	|	j                         }
|	j                         }|	j	                  |       |j                  ||i       }g }|t        |      }|	j                  |       |D cg c]  }|j                          }}|	j                  |
       t        |      }|	j                  |||      }|j                         j                  |       |j                  |       |	j                  |j                         g        |	j!                  |      }|D cg c]  }|j                          }}t#        |      D ]  }t%        ||         }|	j                  |j'                  |      |      }t#        t        |            D cg c]  }|j)                  |       }}t+        ||D cg c]  }|j,                   c}      }|j                  ||   |i |       |	j/                           |	j1                  |j3                                t#        t        |            D cg c]  }|j5                  |       }}|y t7        t+        ||D cg c]  }|j,                   c}            S c c}w c c}w c c}w c c}w c c}w c c}w )Nzwarp specialize got z partitions but z warp countsz register counts)kwargsr   )r	  caller_context)rF   r)   get_insertion_point	new_blockset_insertion_point_to_startcall_JitFunctionr   create_warp_yieldget_typerestore_insertion_pointcreate_warp_specializeget_default_region	push_backset_requested_registerscreate_block_with_parentget_partition_op_holder!create_warp_specialize_partitionsr   r   
get_regionget_argumentr   r=   create_warp_returnset_insertion_point_afterget_operationr   r   )r    default_argsdefault_partitionworker_argsworker_partitionsr  r  	generatornum_partitionsr)   	insert_ptdefault_blockdefault_resultsmlir_resultsrresult_types	mlir_argsws_oppartitions_oparg	arg_typesrL   r
  blockj
block_argss                             r   warp_specializezGluonSemantic.warp_specializeN  s   ./"
 
 	f!.!11A#FVBWAXXde	f 
 "
 
 	i!.!11A#oBVAWWgh	i 
 ,,//1	  ))+,,];#445F]_4`&/@L!!,/.:;

;; 	''	2(5	..|YHXY  ",,];%%o6 	(()F)F)H"MAA.Q/89S\\^9	9~& 	)A/:J1:MNN44]5M5Ma5PR[\E9>s9~9NOA%,,Q/OJO,Zk9Zs#((9Z[J&&'8';ZPRcq&r&&(	) 	))%*=*=*?@5:3|;L5MN((+NN"(7X17XYZZ3 < : P9Z
 O 8Ys$   K8K%K	KK .K%)F)/r+   r,   r-   r3   r6   langr   __annotations__r!   r;   r?   r   r.   rO   r   rd   rj   r   rn   rs   r}   rh   r   boolr   r   r   r   r   r   r   r   r   r   r   r   r   r   r   staticmethodr   r   r   r   r  r2  __classcell__)rl   s   @r   r0   r0      s!   [[FD '_49 c  T T T T&5h 58 5 5Xx XE(H*<$= X5X 5U3Z 5H 5+( +5: +( +  x H :95X 5$s) 5$ 5
+
1	+W+
G
;+	DDY
*D
& X X)x'9 ) )"&)+03+?)*) 2 )# )UZ[ceh[hUi )$
Hx 
H3 
Hh 
HS[ 
H.[*23-.[JRSV-.[r   r0   )typingr   r   r   r   r   r   triton.language.semanticr    r
   r3   _layoutsr   r   r   triton._C.libtriton.gluon_irr   triton.compiler.code_generatorr   r   r   rG   r5  rJ   r   r   r0   r   r   r   <module>r>     sm    ; ;  3  @ @ 7 T: <F ! !xC0 !
	M 	M_[N8, _[r   