
    ^ja                       d dl mZ d dlZd dlZd dlZd dlZd dlmZ d dlm	Z	m
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mZ d d
lmZ ddlmZmZmZ ddlm Z m!Z!m"Z" ddl#m$Z$m%Z%m&Z&m'Z'm(Z(m)Z) ddl*m+Z+m,Z,m-Z- e
rddl.m/Z/m0Z0 ddl1m2Z2m3Z3 ddl#m4Z4  ejj                  e6      Z7ejp                  dejr                  dejt                  dejv                  dejx                  dejz                  dej|                  dej~                  dej                  di	ZAd'dZB G d de      ZC G d  d!e(      ZDeDj                  d"       eDj                           G d# d$e,      ZG G d% d&e-      ZHy)(    )annotationsN)Path)AnyTYPE_CHECKING)
PRECEDENCE)_embed_headers)
OrderedSet)Min)
CppPrinterExprPrinter)ValueRanges   )ceildivget_bounds_index_exprget_kernel_metadata)ops
OpsWrapperV   )CSEVariableDeferredLineDTYPE_TO_COMPUTATION_DTYPEIndentedBufferOpOverridesPythonPrinter)IterationRangesEntry
SIMDKernelSIMDScheduling)ReductionType	StoreMode)	SchedulerSchedulerNode)OpVarTboolcharshortintlongucharfloathalfbfloatc                    t        | t              r:| t        j                  k(  ry| t        j                   k(  ry| | k7  ryt	        |       S t        | t
              r| rdS dS t	        |       S )N	HUGE_VALFz
-HUGE_VALFNANtruefalse)
isinstancer*   torchinfstrr$   )vals    f/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_inductor/codegen/mps.pyvalue_to_metalr8   7   s_    #u%))UYYJCZ3x	C	v)')s8O    c                      e Zd ZdZddZddZddZedd       ZddZ	ddZ
ddZdd	Zdd
ZddZddZddZddZddZeZddZddZddZy)MetalExprPrinterz/Converts sympy expression to Metal code snippetc                    |j                   \  }}| j                  |      }| j                  |      }|j                  r	d| d| dS d| d| dS )Nc10::metal::floor_divide(, )metal::floor() / (argsdoprint
is_integer)selfexprxdivs       r7   _print_FloorDivz MetalExprPrinter._print_FloorDivH   s[    3LLOll3??.qcC5::qcse1--r9   c                   |j                   \  }}}|dk(  xr ||dz
  z  dk7  }| j                  |      }|dk7  r0| j                  |      }|j                  r
d| d| d}n	d| d| d}| j                  |      }|r	d| d	| dS d| d
| dS )Ni   r   r   (rA   r?   r@   zc10::metal::safe_mod(r>   z) % (rB   )rF   rG   rH   rI   moduse_safe_mods         r7   _print_ModularIndexingz'MetalExprPrinter._print_ModularIndexingP   s    ii3e|>aQ(>LLO!8,,s#Cs%uA&#A3eC52ll3*1#RuA661#U3%q!!r9   c                    t        t        | j                  |j                              }|d   }|dd  D ]  }| j	                  |||      } |S )Nr   r   )listmap_printrC   _print_binary_min_max)rF   rG   fnrC   resultargs         r7   _print_min_maxzMetalExprPrinter._print_min_maxa   sT    CTYY/0a8 	AC//R@F	Ar9   c                J    d|  d| d|  d}d|  d| d| d}d| d| d| dS )Nstatic_cast<decltype(+)>(r?   zmetal::rL   r>    )abrU   
typecast_a
typecast_bs        r7   rT   z&MetalExprPrinter._print_binary_min_maxi   sQ    ,QCq3qc;
,QCq3qc;
Aj\J<q99r9   c                &    | j                  |d      S )NminrX   rF   rG   s     r7   
_print_MinzMetalExprPrinter._print_Mino       ""4//r9   c                &    | j                  |d      S )Nmaxrd   re   s     r7   
_print_MaxzMetalExprPrinter._print_Maxr   rg   r9   c                z    t        |j                        dk(  sJ d| j                  |j                  d          dS )Nr   metal::abs(r   r?   lenrC   rS   re   s     r7   
_print_AbszMetalExprPrinter._print_Absu   s9    499~"""T[[167q99r9   c                z    t        |j                        dk(  sJ d| j                  |j                  d          dS )Nr   zstatic_cast<long>(metal::rint(r   ))rm   re   s     r7   _print_RoundToIntz"MetalExprPrinter._print_RoundToIntz   s9    499~"""/DIIaL0I/J"MMr9   c                    t        |j                        dk(  sJ |j                  \  }}|j                  r|dk  sJ t        d| d      | j	                  |t
        d         }d| d| d|  d	S )
Nr   r   zOFor integer inputs, only non-negative ndigits are currently supported, but got .Mulz!static_cast<float>(metal::rint(1e * z) * 1er?   )rn   rC   rE   
ValueErrorparenthesizer   )rF   rG   numberndigits
number_strs        r7   _print_RoundDecimalz$MetalExprPrinter._print_RoundDecimal   s    499~"""))Q;;abiajjkl  &&vz%/@A
27)3zl&RYQYPZZ[\\r9   c                n    |j                   \  }}d| j                  |       d| j                  |       dS )Nstatic_cast<float>(z) / static_cast<float>(r?   )rC   rS   )rF   rG   lhsrhss       r7   _print_IntTrueDivz"MetalExprPrinter._print_IntTrueDiv   s=    99S %T[[%5$66MdkkZ]N^M__`aar9   c                    t        |j                        dk(  sJ t        | j                  |j                        \  }}d| d| dS )Nr   z'metal::precise::pow(static_cast<float>(z), static_cast<float>(rq   )rn   rC   rR   rD   )rF   rG   rH   ys       r7   _print_PowByNaturalz$MetalExprPrinter._print_PowByNatural   sI    499~"""4<<+18;QRSQTTVWWr9   c                ~    t        |j                        dk(  sJ | j                  |j                  d         }d| dS )Nr   r   r~   r?   rn   rC   rD   rF   rG   rH   s      r7   _print_ToFloatzMetalExprPrinter._print_ToFloat   s=    499~"""LL1&$QCq))r9   c                X    |j                   rt        t        |            S t        |      S N)rE   r5   r'   re   s     r7   _print_FloatzMetalExprPrinter._print_Float   s#    ?? s4y>!t9r9   c                ~    t        |j                        dk(  sJ | j                  |j                  d         }d| dS )Nr   r   z1static_cast<int>(metal::floor(static_cast<float>(z)))r   r   s      r7   _print_FloorToIntz"MetalExprPrinter._print_FloorToInt   s=    499~"""LL1&B1#SIIr9   c                ~    t        |j                        dk(  sJ | j                  |j                  d         }d| dS )Nr   r   zstatic_cast<int>(metal::trunc(rq   r   r   s      r7   _print_TruncToIntz"MetalExprPrinter._print_TruncToInt   s=    499~"""LL1&/s"55r9   c                ~    t        |j                        dk(  sJ | j                  |j                  d         }d| dS )Nr   r   zmetal::precise::log2(r?   r   r   s      r7   _print_OpaqueUnaryFn_log2z*MetalExprPrinter._print_OpaqueUnaryFn_log2   s=    499~"""LL1&&qc++r9   c                J      fd|j                   D        \  }}}| d| d| S )Nc              3  V   K   | ]   }j                  |t        d    dz
         " yw)Atomg      ?N)rx   r   ).0rW   rF   s     r7   	<genexpr>z0MetalExprPrinter._print_Where.<locals>.<genexpr>   s-      
ADDc:f#5#;<
s   &) ? z : )rC   )rF   rG   cpqs   `    r7   _print_WherezMetalExprPrinter._print_Where   s5    
HL		
1a Cs#aS!!r9   N)rG   
sympy.Exprreturnr5   )rG   r   rU   r5   r   r5   )r^   r5   r_   r5   rU   r5   r   r5   )__name__
__module____qualname____doc__rJ   rO   rX   staticmethodrT   rf   rj   ro   rr   r|   r   r   r   r   r   _print_floorr   r   r   r]   r9   r7   r;   r;   E   su    9."" : :
00:
N

]bX
*
J
 %L6
,
"r9   r;   c                  t   e Zd ZdZe	 	 d1	 	 	 	 	 	 	 	 	 d2d       Ze	 	 	 	 	 	 	 	 d3d       Zed4d       Zed5d       Zed5d       Z	ed6d       Z
ed7d	       Zed8d
       Zed9d       Zed9d       Zed9d       Zed9d       Zed:d       Zed:d       Zed:d       Zed:d       Zed:d       Zed:d       Zed:d       Zed:d       Zed:d       Zed:d       Zed:d       Zed:d       Zed:d       Zed;d       Zed:d       Zed:d       Z ed:d       Z!ed:d        Z"ed:d!       Z#ed9d"       Z$ed:d#       Z%ed:d$       Z&ed9d%       Z'ed:d&       Z(ed9d'       Z)ed:d(       Z*ed<d)       Z+ed<d*       Z,e	 	 	 	 	 	 	 	 	 	 d=d+       Z-ed:d,       Z.ed9d-       Z/d>d.Z0d?d/Z1e2d@d0       Z3y)AMetalOverrideszXImplements Metal-specific overrides for ops. Base class emits Python-friendly overrides.Nc                ~    |t         j                  k(  rt        j                  d       d|  dS dt        |    d|  dS )Nz>float64 cast requested, probably from tensorify_python_scalarsr~   r?   static_cast<>()r3   doublelogwarningDTYPE_TO_METAL)rH   dtype	src_dtypeuse_compute_typess       r7   to_dtypezMetalOverrides.to_dtype   sK     ELL KKP )1--nU34Bqc;;r9   c                6    dt         |    dt         |    d|  dS )Nzas_type<z>(static_cast<r   rq   r   )rH   r   r   s      r7   to_dtype_bitcastzMetalOverrides.to_dtype_bitcast   s/     ./0~i?X>YY[\][^^`aar9   c                    t        |       S r   r8   )r6   r   s     r7   constantzMetalOverrides.constant   s    c""r9   c                @   t         j                  j                  t         j                  j                  |             }t         j                  j                  j                  t         j                  j                  |t        |             }t        j                  ||      S )N)bounds)
r   kernelindex_to_strprepare_indexingcsegeneratecomputer   r   r   )rG   r   idx_strvars       r7   
index_exprzMetalOverrides.index_expr   sl    ((''(A(A$(GHhhll##HHg.CD.I $ 
 ||C''r9   c                .    t         j                  | |      S r   )r   r   )rG   r   s     r7   
value_exprzMetalOverrides.value_expr   s     ((u55r9   c           	        t        |      }t               }t        j                  j	                  |      5  |j                         5  t        j                         t        j                  j                  _	        t        j                  j                  xj                  dz  c_
         |       }d d d        d d d        |  d|j                          d| }t        j                  j                  j                  |      }|sct        j                  j                  j                  j                        }t        j                  j                  j                  ||       t        j                  j                   j#                  t$        |j                      d| dd|  dg       t        j                  j                   j                         5  t        j                  j                   j'                  |       t        j                  j                   j)                  | d| d	| d
       d d d        t        j                  j                   j)                  d| d| d	| d
       |S # 1 sw Y   xY w# 1 sw Y   xY w# 1 sw Y   XxY w)N_scoped:r    ;if () {z = static_cast<decltype(r\   );z} else )r8   r   r   r   swap_buffersindent	itertoolscountr   iter_buffer_idsname_prefixgetvaluetry_getnewvarr   putr   
writelinesr   splice	writeline)maskbodyother	other_strscoped_bodyrc	cache_keyr   s           r7   maskedzMetalOverrides.masked   s   
 #5)	$&XX"";/ 	1C1C1E 	 ,5??+<AHHLL(HHLL$$	1$B	 	 fAk2245QykB	hhll""9-((,,%%BHH%5CHHLLY,HH''"288,-Qse15dV47HI !!((*   ''4  **e3C5B4rB
 HH&&3%7uC	{"M 
5	 	 	 	$ s1   I2A&I%,I2AI?%I/	*I22I<?Jc           	     0    |  d| d| dt        |       dS )Nr   z : static_cast<decltype(r\   r?   r   )r^   r_   r   s      r7   wherezMetalOverrides.where  s*     Cs21#S9J8K1MMr9   c                    d|  d| dS )Nzc10::metal::remainder(r>   r?   r]   r^   r_   s     r7   	remainderzMetalOverrides.remainder  s    's"QCq11r9   c                D    d|  d| d|  d}d|  d| d| d}d| d| dS )NrZ   r[   r\   r?   zc10::metal::max(r>   r]   r^   r_   r`   ra   s       r7   maximumzMetalOverrides.maximum  M     -QCq3qc;
,QCq3qc;
!*R
|1==r9   c                D    d|  d| d|  d}d|  d| d| d}d| d| dS )NrZ   r[   r\   r?   zc10::metal::min(r>   r]   r   s       r7   minimumzMetalOverrides.minimum  r   r9   c                    |  d| S )Nz || r]   r   s     r7   
logical_orzMetalOverrides.logical_or       D}r9   c                    |  d| S )Nz && r]   r   s     r7   logical_andzMetalOverrides.logical_and$  r   r9   c                    d|  dS )Nzmetal::isnan(r?   r]   rH   s    r7   isnanzMetalOverrides.isnan)       qc##r9   c                    d|  dS )Nzmetal::isinf(r?   r]   r   s    r7   isinfzMetalOverrides.isinf.  r   r9   c                    d|  dS )Nzmetal::precise::log(r?   r]   r   s    r7   r   zMetalOverrides.log3       &aS**r9   c                    d|  dS )Nzmetal::precise::exp(r?   r]   r   s    r7   expzMetalOverrides.exp8  r   r9   c                    d|  dS )Nrl   r?   r]   r   s    r7   abszMetalOverrides.abs=  s     QCq!!r9   c                    d|  dS )Nzmetal::signbit(r?   r]   r   s    r7   signbitzMetalOverrides.signbitB  s     !1%%r9   c                    d|  dS )Nzmetal::precise::sin(r?   r]   r   s    r7   sinzMetalOverrides.sinG  r   r9   c                    d|  dS )Nzc10::metal::sinc(r?   r]   r   s    r7   sinczMetalOverrides.sincL  s    "1#Q''r9   c                    d|  dS )Nzmetal::precise::cos(r?   r]   r   s    r7   coszMetalOverrides.cosP  r   r9   c                    d|  dS )Nzmetal::precise::tan(r?   r]   r   s    r7   tanzMetalOverrides.tanU  r   r9   c                    d|  dS )Nzmetal::precise::asin(r?   r]   r   s    r7   asinzMetalOverrides.asinZ       'qc++r9   c                    d|  dS )Nzmetal::precise::acos(r?   r]   r   s    r7   acoszMetalOverrides.acos_  r
  r9   c                    d|  dS )Nzmetal::precise::atan(r?   r]   r   s    r7   atanzMetalOverrides.atand  r
  r9   c                    d|  d| dS )Nz::metal::precise::atan2(r>   r?   r]   )rH   r   s     r7   atan2zMetalOverrides.atan2i  s     *!Bqc33r9   c                    d|  dS )Nzmetal::precise::sqrt(r?   r]   r   s    r7   sqrtzMetalOverrides.sqrtn  r
  r9   c                    d|  d|  dS )NrZ   z)>(-r?   r]   r   s    r7   negzMetalOverrides.negs  s     'qcaS22r9   c                    d|  dS )Nzmetal::precise::rsqrt(r?   r]   r   s    r7   rsqrtzMetalOverrides.rsqrty       (s!,,r9   c                    d|  dS )Nzmetal::precise::tanh(r?   r]   r   s    r7   tanhzMetalOverrides.tanh~  r
  r9   c                    d|  dS )Nzmetal::precise::atanh(r?   r]   r   s    r7   atanhzMetalOverrides.atanh  r  r9   c                    d|  d| dS )Nr=   r>   r?   r]   r   s     r7   floordivzMetalOverrides.floordiv  s     +1#Rs!44r9   c                    d|  dS )Nr@   r?   r]   r   s    r7   floorzMetalOverrides.floor  r   r9   c                    d|  dS )Nzmetal::sign(r?   r]   r   s    r7   signzMetalOverrides.sign       aS""r9   c                D    d|  d| d|  d}d|  d| d| d}d| d| dS )NrZ   r[   r\   r?   zmetal::fmod(r>   r]   r   s       r7   fmodzMetalOverrides.fmod  sM     -QCq3qc;
,QCq3qc;
j\J<q99r9   c                    d|  dS )Nmetal::trunc(r?   r]   r   s    r7   trunczMetalOverrides.trunc  r   r9   c                    |  d| }| j                   | j                   j                  s"|j                   |j                   j                  rd| dS |S )Nz / r&  r?   )r   is_floating_point)r^   r_   quots      r7   truncdivzMetalOverrides.truncdiv  sS     Cs|GGAGG$=$=GGAGG$=$="4&**r9   c                    d|  dS )Nzmetal::ceil(r?   r]   r   s    r7   ceilzMetalOverrides.ceil  r"  r9   c                f    t         j                  j                  j                  d       d|  d| dS )Nrandomzc10::metal::rand(r>   r?   r   r   headersaddseedoffsets     r7   randzMetalOverrides.rand  s/    	X&"4&6(!44r9   c                f    t         j                  j                  j                  d       d|  d| dS )Nr/  zc10::metal::randn(r>   r?   r0  r3  s     r7   randnzMetalOverrides.randn  s/    	X&#D6F8155r9   c           	     r    t         j                  j                  j                  d       d|  d| d| d| d	S )Nr/  zc10::metal::randint64(r>   r?   r0  )r4  r5  lowhighs       r7   	randint64zMetalOverrides.randint64  s=     	
X&'vRxr#baHHr9   c                    d|  dS )Nzmetal::rint(r?   r]   r   s    r7   roundzMetalOverrides.round  r"  r9   c                D    d|  d| d|  d}d|  d| d| d}d| d| dS )NrZ   r[   r\   r?   zmetal::pow(r>   r]   )r^   r_   cast_acast_bs       r7   powzMetalOverrides.pow  sK    (1QCs1#Q7(1QCs1#Q7VHBvha00r9   c                f    t         j                  j                  j                  d       d| d| dS )Nspecial_mathc10::metal::rL   r?   r0  )rF   r^   names      r7   _special_unaryzMetalOverrides._special_unary  s/    	^,dV1QCq))r9   c                l    t         j                  j                  j                  d       d| d| d| dS )NrD  rE  rL   r>   r?   r0  )rF   r^   r_   rF  s       r7   _special_binaryzMetalOverrides._special_binary  s5    	^,dV1QCr!A..r9   c           
        dD ].  }t        | |t        j                  | j                  |             0 t        j                  | j                  d      | _        dD ]1  }t        | |t        j                  | j                  |dz                3 dD ].  }t        | |t        j                  | j
                  |             0 dD ]1  }t        | |t        j                  | j
                  |dz                3 y )N)erferfinvi0i0ei1i1edigammaspherical_bessel_j0)rF  	log_gamma)
	bessel_j0	bessel_j1	bessel_y0	bessel_y1modified_bessel_i0modified_bessel_i1modified_bessel_k0modified_bessel_k1scaled_modified_bessel_k0scaled_modified_bessel_k1_forward)	polygammaigammaigammaczeta)
chebyshev_polynomial_tchebyshev_polynomial_uchebyshev_polynomial_vchebyshev_polynomial_whermite_polynomial_hhermite_polynomial_heshifted_chebyshev_polynomial_tshifted_chebyshev_polynomial_ushifted_chebyshev_polynomial_vshifted_chebyshev_polynomial_w)setattr	functoolspartialmethodrG  lgammarI  )clsrF  s     r7   _initialize_special_opsz&MetalOverrides._initialize_special_ops  s    	
 
	WD Cy66s7I7IPTUV
	W ,,S-?-?kR

 	D ''(:(:
ARS	&
 	XD Cy66s7J7JQUVW	X
 	D ''(;(;$BST	r9   NT)
rH   r   r   torch.dtyper   ztorch.dtype | Noner   r$   r   r5   )rH   r   r   rt  r   rt  r   r5   )r6   zbool | float | intr   rt  r   r5   )rG   r   r   rt  r   r5   )r   r   r   r   r   r   r   r5   )r^   r#   r_   r#   r   r#   r   r5   )r^   r#   r_   r#   r   r5   )r^   r   r_   r   r   r5   )rH   r   r   r5   )rH   r   r   r   r   r5   )r4  r   r5  r   r   r5   )
r4  r   r5  r   r:  r   r;  r   r   r5   )r^   r   rF  r5   r   r5   )r^   r   r_   r   rF  r5   r   r5   r   None)4r   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-  r6  r8  r<  r>  rB  rG  rI  classmethodrr  r]   r9   r7   r   r      s   b )-"&	<<< &<  	<
 
< < bb*b7Bb	b b
 # # ( ( 6 6
    D N N 2 2 > >
 > >
     $ $ $ $ + + + + " " & & + + ( ( + + + + , , , , , , 4 4 , , 3 3
 - - , , - - 5 5 $ $ # # : :
 $ $   # # 5 5 6 6 II#.I5@IHSI	I I # # 1 1
*/ = =r9   r   mpsc                      e Zd ZU dZeZdZdZdZdZ	 e
       j                  Z e       j                  Z e       j                  ZeZ edg      Zded<   g Zd	ed
<   	 	 	 	 	 	 d fdZddZddZ	 d	 	 	 	 	 	 	 	 	 ddZd dZddd ej8                         f	 	 	 	 	 	 	 	 	 	 	 d!dZ	 	 	 	 	 	 	 	 	 	 d"dZ	 	 	 	 	 	 	 	 	 	 d"dZd#dZ d$dZ!d%d&dZ"	 d'	 	 	 	 	 	 	 d(dZ#d)dZ$	 	 	 	 	 	 	 	 	 	 d*dZ% xZ&S )+MetalKernelz;Implement Metal codegen based on the SIMDKernel abstractionr   auto i       utilszOrderedSet[str]r1  zlist[IterationRangesEntry]multistage_reduction_entryc                X    t        |   |fi | t        j                         | _        y r   )super__init__r   r   acc_var_ids)rF   tilingkwargs	__class__s      r7   r  zMetalKernel.__init__*  s&    
 	*6*$??,r9   c                    t         |   S r   r   )rF   r   s     r7   dtype_to_strzMetalKernel.dtype_to_str2  s    e$$r9   c                   | j                   j                  |      }| j                  |      }t        j                  j                  |      }| d| j                  |       d}|t        j                  t        j                  fv rd| d}t        j                  }| j                  j                  | j                  ||      S )z"Codegen a load from an InputBuffer[]r~   r?   r   )rC   inputr   r   graph	get_dtyper   r3   float16bfloat16float32r   r   loads)rF   rF  indexr   r   lines         r7   loadzMetalKernel.load5  s    iiood#%%e,!!$'a))%013U]]ENN33 )a0DMMExx  T ??r9   Nc                h   | j                   j                  |      }| j                  |      }| j                  t        j
                  j                  |            }d| d| d}|| d| j                  |       d| d}n[|dk(  rH| j                  j                  d       d	| d
}	d|	 d| d}
|	 d|
 d| j                  |       d| d}nt        d|       | j                  r&| j                  j                  t        ||             y | j                  j                  t        ||             y )Nr   r   r?   r  ] = r   
atomic_addatomiczc10::metal::AtomicType<>zreinterpret_cast<device z
::type *>(z::atomic_add(r>   r   zUnimplemented store mode )rC   outputr   r  r   r  r  r   r1  r2  RuntimeErrorinside_reductionr   r   r   stores)rF   rF  r  valuemoder   	dtype_strcast_valr  atomic_typecast_vars              r7   storezMetalKernel.storeC  s<    iit$%%e,%%agg&7&7&=>	!)BugQ7<U!D--e45T(1ED\!LLX&3I;a@K1+jQOH!]-zD<M<Me<T;UUWX`WaacdD!:4&ABB  LL""<d#;<KK!!,tT":;r9   c                   | j                   j                  |      }| j                  |      }| j                  t        j
                  j                  |            }t        d | j                  D              }| d| j                  |       d| d| d}d|j                   d| }| j                  j                  t        ||             y )Nc              3  :   K   | ]  }|j                   s|  y wr   is_reductionr   ts     r7   r   z.MetalKernel.store_reduction.<locals>.<genexpr>]  s     K1ANNQK   r  z] = static_cast<r   r   r   z == 0) )rC   r  r   r  r   r  r  nextrange_treesr   rF  r  r   r   )rF   rF  r  r  r   r  reduction_dimr  s           r7   store_reductionzMetalKernel.store_reductionX  s    iit$%%e,%%agg&7&7&=>	K(8(8KKa))%011A)BugUWXm(()7l467r9   Tc                   t        |t        j                        r| j                  |      }dt	        | j
                         }t        j                  j                  |||      }|rdnd}|| d| z  }|r|d| j                  |       dz  }||rJ d       |d| z  }| j                  j                  || j                  z          |S )	Ntmp_acc_zthreadgroup  r   r  r  z+Thread group var can not have default value = )r2   r3   r   r  r  r  r   r   create_cse_varsexprindexing_coder   suffix)	rF   r   
elem_countdefault_valueis_threadgroupr   var_namer   var_defs	            r7   _new_idxvarzMetalKernel._new_idxvarc  s     eU[[)%%e,Ed4#3#3456hh%%h>$2.eWAhZ((4::j12!44G$%T'TT%]O,,G$$Wt{{%:;
r9   c                    |||f}|| j                   j                  v r| j                   j                  |   S | j                  ||||      }|| j                   j                  |<   |S )z)Caching wrapper around _reduction_nocache)r   reduction_cache_reduction_nocache)rF   r   r   reduction_typer  r   rV   s          r7   	reductionzMetalKernel.reductiony  sf     6	00088++I66((	>5Q.4  +r9   c                $   | j                   sJ | j                  rJ dAd}d}d}| j                  D ]  }|j                  s|r|dz  }||j                   d| z  }t        |j                  t        j                        r||j                  z  }_|t        j                  |j                   ddd      z  } t        || j                        }| j                  |      }	t        |t        j                        r|n| j                  }
t        |t        j                        rt        || j                        n| j                  }|d	k(  r| j!                  |      }| j"                  j%                  | d
       | j"                  j%                  d       | j&                  j)                  d| d| d       | j*                  j%                  d       |S | j,                  j/                  d       |dv rt0        |   }| j!                  ||      }| j2                  s|}nD|dk(  rdnd\  }}| j!                  ||d      }| j&                  j)                  | d| d| d       | j4                  j7                  | j*                  d| d| d| d| d|	 dt0        |         S |dv r| j!                  ||      }t8        |   }d| d | d}| j2                  s|}nM|d!k(  rd"nd#}d$| d%| d&}| j!                  ||d      }| j&                  j)                  | d'| d| d| d(       | j4                  j7                  | j*                  d| d| d| d| d|	 dt0        |         S |d)v r5| j!                  ||      }| j!                  ||      }t8        |   }d| d | d}| j2                  s|}dt8        |    d | d}n|d*k(  rd"nd#}d$| d%| d&}| j!                  ||d      }| j!                  |d+d      }t;        d, | j<                  j?                         D              }| j&                  j)                  d-| d%| d.| d| d/| d0| d1| d0|j                   d2       | j4                  j7                  | j*                  d| d| d| d| d| d| d|	 d|      S |d3k(  r+| j2                  s~| j!                  ||
      }| j&                  j)                  | d4| d5| d       | j4                  j7                  | j&                  d| d| d|	 dt@        jB                        } ||      S | j!                  d6|
      }| d4| d7}| j"                  j)                  | d8       | j&                  j%                  | d9| d:| d;       | j4                  j7                  | j*                  d<| d|	 dt@        jB                        } ||      S |d=k(  r4t        |tD              sJ d>       | j!                  d6|
      }| d4| d7}d?|d+    d|d    d|d@    d}| j"                  j)                  | d8       | j2                  rC| j"                  j)                  | d8       | j&                  j%                  | d9| d| d(       n!| j&                  j%                  | d0| d       | j4                  j7                  | j2                  r| j*                  n| j&                  d| d| d|	 dt@        jB                        } ||      S tG        |      )BzeCodegen a reduction operation.
        Only sum and prod operations are somewhat reasonable optimizedc           
         t        j                  dD cg c](  }t        |  d| | j                  | j                        * c}      S c c}w )Nxyzrt   )r   _unwrapr   r   r   )res3r  s     r7   _unwrap_helperz6MetalKernel._reduction_nocache.<locals>._unwrap_helper  sA    %%NSTvQqc]DKKDT Ts   -Ar  r    + rv   numelTintegerpositiveanyz	 = false;z7threadgroup_barrier(metal::mem_flags::mem_threadgroup);z
                if (z) {
                    z' = true;
                }
            reduction_utils)prodsumr  )r   r[   )r   *F)r  r  r   z= r   zc10::metal::threadgroup_rL   r>   r?   r   )ri   rc   r   r   ri   MaxOpMinOpz::c10::metal::<z>::identity()z = ::c10::metal::r   )argminargmaxr  r   c              3  :   K   | ]  }|j                   s|  y wr   r  r  s     r7   r   z1MetalKernel._reduction_nocache.<locals>.<genexpr>  s      Ar  z#
                if (::c10::metal::z>::replace(z)) {
                    r  z;
                    z$;
                }
                welford_reducer  r  float3r  z = 0.0;z! = ::c10::metal::welford_combine(z	, float3(z, 0.0, 1.0));z(c10::metal::threadgroup_welford_combine(welford_combinez&Input to welford combine must be tuplezfloat3(r   )r  r   r   ztuple[CSEVariable, ...])$r  
_load_maskr  r  rF  r2   r  sympyIntegerSymbolprefixr
   max_threadgroup_sizer  r   simd_group_sizer  r  r   r   r   r  r1  r2  r   r~  r   r   r   r  range_tree_nodesvaluesr3   r  tupleNotImplementedError)rF   r   r   r  r  r  reduction_idxacc_buf_sizerdacc_buf_size_stracc_buf_alloc_sizeshmem_buf_sizeacc	acc_dtypeacc_bufr6   default_valreduction_opsrc_metal_type
cast_value	op_struct	limit_valdata_acc_bufidx_acc_bufidx_validx_varwf_resacc_thread_var	inp_values                                r7   r  zMetalKernel._reduction_nocache  sJ    $$$$??""	 "" 	B??&yL>::M"((EMM2(yyk'! 	 <)B)BC::l3 ,6 ** 	 ,6 L$"6"67%% 	 U"""5)C((C5	):;((I LLG E  KK!!I J*+_,29=I&&y.AG22 !/% 7HX *\ &&[ '  ##se1\N"UG1$EF88$$*>*:!G9Bse2m_\^_o^ppqr07 %  
 ^+&&y.AG+I6N''7r%BJ22 '5'>GG	,YKq8HV	&&Yu '  ##e,^,<AcU"ZLPRS 88$$*>*:!G9Bse2m_\^_o^ppqr07 %  
 11++I~FL**5.AK+I6N''7r%BJ22 ()>(?r-PQR'5'AGw	,YKq8HV	&&Yu '  **5RW*X #44;;=  ## )##,+Q~.>k*UWX[W\ ]EZL )IS /%  88$$*>*:!L>K=XZ%r'"]O26F5GqJ	 %   --22**96HI##wiqtE7!$LM**LL.~.>ayK[J\\]^-- + 
 &f--&&x1CDG 'y-:N%%(8&@ALL""!""CNCSS\]b\ccpq XX&&:7)2FVEWWXYmm ' F
 "&))..eU+U-UU+&&x1CDG 'y-:N!%(2eAhZr%(1EI%%(8&@A..""))^,<G*DE&&%&&GGWWYZcYddfg &&.)9YKq'IJXX&&#>>DLL*>*:!G9BGWFXXYZmm ' F
 "&))!.11r9   c                   | j                  j                        }| j                  |      }j                  rQt	        j
                  j                  t        j                        r\j
                  j                  | j                  k  r9| j                  j                  | j                   dj                   d| d       y t	        j
                  j                  t        j                        rj
                  j                  n.t        j                  j
                  j                   ddd      }t!        fd| j"                  D              }j
                  j                   d}|s\| j"                  j%                         |t'        | j                  d	z
        z   t'        | j                        z  }| j                  |      }j
                  j                  }	| j(                  j                  d
j
                  j                   dj
                  j                   d| dj
                  j                   d	       | j(                  j+                         5  t	        |t        j                        rS| j(                  j                  | j                   d| d| j                   dj
                  j                   d|	 d
       nH| j(                  j                  | j                   d| d| d|	 dj
                  j                   d
       t	        |t        j                        s|| j                  z  |k7  r"| j(                  j                  d| d| d       |j-                  j
                  j                  |      }
| j(                  j                  | j                   dj                   d|
 d       d d d        y | j(                  j+                         5  |j-                  j
                  j                  |      }
| j(                  j                  | j                   dj                   d|
 d       d d d        y # 1 sw Y   y xY w# 1 sw Y   y xY w)Nr   r  r   r  Tr  c              3  N   K   | ]  }|j                   j                   u   y wr   root)r   eentrys     r7   r   z=MetalKernel.codegen_iteration_ranges_entry.<locals>.<genexpr>T  s#      %
%&AFFejj %
s   "%_linear_idxr   z	for(auto z
_cnt = 0; z_cnt < z; ++z_cnt) {rv   z_cnt + r  z_cnt;r    >= z) break;)rename_indexingrG   r  r  r2   r   r  r  r  r  r  r   index_dtyperF  r  r  r  r~  appendr*   r   r   replace)rF   r  r   	index_stracc_sizeroot_already_processedlinear_idx_name	loop_sizeloop_size_str	root_namesub_index_strs    `         r7   codegen_iteration_ranges_entryz*MetalKernel.codegen_iteration_ranges_entry@  s   ))%**5
JJz*	!!uzz''7

  D$=$==((##$Aejj\YKqA  %****EMM: JJ!2!2 3594RVW 	 "% %
*.*I*I%
 "
 #ZZ../{;%++2259 "E$*C*Ca*G$HHU))N I !JJy1M

IIIEJJ--.j9J9J8K7S`Raaefkfpfpfwfwex  yA  B !!# h5II''++,Ao->c445S9J9J8K7S\R]]^`
 II''++,Ao->c-PST]S^^abgblblbsbsattyz x6 4#<#<<HII''$.?tH:X(VW !* 1 1%**//? S		##''(%**SqI) 2 !!#  ) 1 1%**//? S		##''(%**SqI 3 2 s   "E"Q'AQQQ$c                   | j                   rm| j                  j                         5  | j                  j                  | j                         | j                  j                  | j
                         ddd       | j                  j                  dt        | j                         z         | j                  j                  t        d | j                  j                  j                         D                     | j                   D cg c]  }|j                   }}| j                   j                          | j                  j                         D ]'  t!        fd|D              sj#                          ) nJ| j                  j                  | j                         | j                  j                  | j
                         | j                  j                  | j$                         | j                  j                          | j
                  j                          | j$                  j                          y# 1 sw Y   xY wc c}w )a  
        Concat output code from index_code, loads, compute, stores,
        suffix into self.body.

        For pointwise kernels, this is called just once at the end.

        For reduction kernels, this generates a loop over the reduction
        axis.
        N}c              3  T   K   | ]   }t        |t              r|n|fD ]  }|  " y wr   )r2   r  )r   itemvs      r7   r   z+MetalKernel.codegen_body.<locals>.<genexpr>  s:      &0u&=dD7  s   &(c              3  :   K   | ]  }j                   |u   y wr   r  )r   rnodes     r7   r   z+MetalKernel.codegen_body.<locals>.<genexpr>  s     5!tyyA~5s   )r~  r   r   r   r  r   r   rn   r   
invalidater	   r  r  r   clearr  r  cache_clearr  )rF   r  rootsr  s      @r7   codegen_bodyzMetalKernel.codegen_body  s    **!!# /		  ,		  ./ IIc$*I*I&J JK
 HH  $ 8 8 ? ? A  &*%D%DEQVVEEE++113--446 '5u55$$&' IITZZ(IIT\\*		%

7/ /  Fs   AH:I:Ic                t   | j                          t               }|}t        j                  j                  r|j                  d       n|j                  d       | j                         }|j                         5  t        j                  j                  s'| j                  D ]  }|j                  d| d        nr| j                  D cg c]  }d| d
 }}t        |t        t              j                  j                  j                  dz  gt                     }|j                  |       | j                  rwt        j                   d | j"                  D              }t%        |t&        j(                        rt+        || j,                        n| j,                  }	|j                  d|	 d       |j                  d	| d
       |j                         5  | j.                  j0                  j3                         D ]Z  \  }
}|
| j4                  v r| j7                  t        j                  j9                  |
            }|j                  d| d| d       \ | j.                  j:                  j3                         D ]  \  }
}t        j                  j9                  |
      }|t<        j>                  k(  rBt        j                  jA                  |
      }||jC                         g k7  rtE        d      d}n| j7                  |      }|j                  d| d| d        | j.                  jF                  jI                         D ]  }|j                  d| d        |D ]F  }t%        |jJ                  t&        j(                        r(|j                  d|jL                   d       H d| j                  v r|j                  d       tO        |      dk  sJ d       tO        |      dkD  rdtO        |       nd}tO        |      dk(  r|d   jP                  nd}| j                  rdnd}|j                  | d| d|        | j                  r|j                  | d       ddd       |j                  d        |j                         5  tO        |      dkD  rAtS        |      D ]3  \  }}|j                  d!|jP                   d"tU        d#|z          d$       5 |jW                  | jX                         |jW                  | jZ                         ddd       |j                  d%       ddd       t        j                  j                  r!|j                  d&       |j]                         S |j                  d'       |j]                         S c c}w # 1 sw Y   3xY w# 1 sw Y   xY w# 1 sw Y   xY w)(z3Called at the end to generate a final kernel stringz(R"MTL(compile_mps_shader('''z#include <c10/metal/z.h>includec              3  N   K   | ]  }|j                   s|j                    y wr   )r  r  r  s     r7   r   z-MetalKernel.codegen_kernel.<locals>.<genexpr>  s      1 !AGG1s   %%z$[[max_total_threads_per_threadgroup(z)]]kernel void rL   zdevice z* ,Nzfloat64 is not supported by MPSr*   z	constant zconstant long& znumel,errorz,device c10::metal::ErrorMessages* error_buf,   z%Up to 3 index variables are supportedr   uintr   
thread_posr  r   z [[thread_position_in_grid]]z- group_pos [[thread_position_in_threadgroup]]r   r{  z = thread_pos.x   r   r  z)MTL");'''))/r  r   r   r  cpp_wrapperr   active_range_treesr   r1  r   r   __file__parentr	   r  mathr  r  r2   r  r  rc   r  rC   output_buffersitemsremoved_buffersr  r  input_buffersr3   float64try_get_bufferget_sizer  sizevarsr  r  r  rn   rF  	enumeratechrr   r  r   r   )rF   rF  codefn_nameidx_varsheaderr1  header_contentstotal_reduction_sizethreadgroup_sizeouterinnerr  r   	outer_bufr  thread_pos_dtypethread_pos_var_namethread_pos_suffixidxr   s                        r7   codegen_kernelzMetalKernel.codegen_kernel  s/   77NN9%NN34**,[[] U	 77&&"ll GFNN%9&#EFG FJ\\;A*6(#6  #1(^**11889DEL#
 /$$'+yy 1%)%5%51 ($ ""6F ,d.G.GH22 !
 :;K:LCP NN\'!45 -$(II$<$<$B$B$D DLE5 4 44  $ 1 1!''2C2CE2J KINNWYKr%#BC	D
 %)II$;$;$A$A$C 
FLE5GG--e4E-$%GG$:$:5$A	$,	0B0B0D0J"./P"QQ$+	$($5$5e$<	NNYykE7!#DE
F "YY//668 ?ENN_UG1#=>?  ( QG!'--?8H'OP	Q dll*NN#QR8}q(Q*QQ(.1(ma.?d3x=/*V ! ),H(:HQK$$ $ ,0+@+@Cb!'(*=)>>Z[lZmn ((NN+,,YZW-\ NN5! 'x=1$$-h$7 S#CHH:^Cc	N;K1M D../DII&' NN3kU	 n 77NN9% }} NN6"}}o2- -^' '[U	  U	 sR   :AV.
VC?V.JV$)V.BV"V.V.V	V."V+	'V..V7c           
     	   t         j                  j                  }| j                  j                  D ]  }|j                  |        | j                  j                         \  }}}}t        ||      D 	
ci c]  \  }	}
t        |	      |
 }}	}
g | j                  j                  j                         | j                  j                  j                         }|D cg c]  }|| j                  vs| }}|| j                  j                  D cg c]  }t        |       c}z  }|D cg c]  }||   	 }}| j                  D ]  }t        |j                  t         j"                  t$        f      r.t        |j                  t         j&                        r|j                  }n4t         j                  j                  j)                  ||      j*                  }|j,                  r| j.                  s|j1                  t        |             |j1                  t$                t         j                  j2                  r| j4                  n| j6                  }dd}t9        | j;                               dkD  r| j;                         D cg c]@  } ||j,                  r t=        |j                  | j>                        n|j                        B }}|j1                   ||d             |j1                  t@               n%t         j                  j2                  rtC        d      | j.                  r}| j;                         D cg c]6  }|j,                  r& |t=        |j                  | j>                              nd8 }}|j1                   ||d             |j1                  t@               n1t         j                  j2                  r|dgz  }|j1                  d       d| jD                  v rMt         j                  j2                  s3|j1                  d	t9        |D cg c]  }|d
|vs| c}              |jG                  ||tI        jJ                  d      d|       yc c}
}	w c c}w c c}w c c}w c c}w c c}w c c}w )z0
        Codegens a call to this kernel
        threadsc                    t         j                  j                  r(| D cg c]  }d| d
 } }ddj                  |        dS | ddj                  |        dS c c}w )Nzstatic_cast<uint64_t>(r?   {r>   r  z=[r  )r   r  r+  join)rJ  kwargr  s      r7   format_threadsz/MetalKernel.call_kernel.<locals>.format_threadsA  sh    ww""BIJQ3A3a8JJDIIg./r22499W#5"6a88 Ks   Ar   zWe should always have threads?1
group_sizeNr%  zerror_buf_idx==rx  F)devicetriton	arg_types)rJ  z	list[str]rN  r5   r   r5   )&r   r  wrapper_coderC   r7  ensure_size_computedpython_argdefszipr5   r0  keysr3  r2  r  r2   r  r  r  r'   r  generate_numel_exprrB  r  r  r  r+  cexprpexprrn   r,  r
   r  rQ   r  r1  generate_kernel_callr3   rS  )rF   rF  r  deallocate_wswrapperr  _	call_argsrU  call_argarg_typearg_name_to_typerC   rW   treerG   expr_printerrO  rJ  s                      r7   call_kernelzMetalKernel.call_kernel  s    ''&&## 	,A((+	, &*YY%=%=%?"9a>A)Y>W
(:(CM8#
 
 S))..0R4993J3J3O3O3QR#Gs$2F2F'FGG!3!34AQ446:;s%c*;	; $$ 	&D$**u}}c&:;DJJ5zzww++??dKQQ$$(=(=CI&  %	& &'WW%8%8tzzdjj	9 t&&()A- 002  ~~ !:!:;G  KKw	:;T"ww"""#CDD  
 002	  >> S$*C*CDEG  KKw=>T"ww""   & dll"177+>+>KK T%`cS_QT\_Q_c%`!a bc 	$$<<& 	% 	
S

 H4;8 & &as=   5QQ$0Q$Q)*Q.AQ3=;Q8Q=Q=#Q=c                (   t         j                  j                  r-| j                  j	                  | j
                  d| dd       y | j                  j                  d       | j
                  j                  d| ddt        |       dd	d
g       y )Nzif (!) returnF
assignmentr%  r   z"    TORCH_REPORT_ERROR(error_buf, r       return;r  )
r   r  r+  r   r   r   r1  r2  r   repr)rF   condmsgs      r7   device_assert_asynczMetalKernel.device_assert_asyncy  s    77HHdlleD6,BuULLW%LL##D6&8c2F!	r9   c                   |s|sy | j                  |      }| j                  |      }|r|rd| d| d| d}n|r| d}n| d| }t        j                  j                  r-| j                  j                  | j                  d| dd	       y | j                  j                  d
       | j                  j                  d| dd| d| dddg       y )NrL   z < 0 || r  r?   z < 0r   rj  Frk  r%  r   z,    TORCH_REPORT_ERROR(error_buf, "Index ", z, " out of range [0, ", z, ")");rm  r  )
r   r   r  r+  r   r   r   r1  r2  r   )rF   rG   sizelowerupperexpr_strsize_str	conditions           r7   check_boundszMetalKernel.check_bounds  s     $$T*$$T* UH:XhZtH:QGI#*D)I#*D
3I 77HHYKx8U  
 LLW%LL##9+T*B8*Ldemdnnuv!	r9   )r  zdict[str, sympy.Expr]r  r   r   rv  )r   rt  r   r5   )rF  r5   r  r   r   r   r   )
rF  r5   r  r   r  r   r  r    r   rv  )rF  r5   r  r   r  r   r   rv  )r   zstr | torch.dtyper  z
int | Noner  z
Any | Noner  r$   r   zValueRanges[Any]r   r   )
r   rt  r   rt  r  r   r  %CSEVariable | tuple[CSEVariable, ...]r   rz  )r  r   r   rv  ru  )generated_kernel)rF  r5   r   r5   rs  )rF  r5   r  r   r_  r$   r   rv  )ro  r   rp  r5   r   rv  )
rG   r   rs  r   rt  r$   ru  r$   r   rv  )'r   r   r   r   r   	overridesr  newvar_prefixr  r  r   rD   r]  r   r\  r;   r  kexprr	   r1  __annotations__r~  r  r  r  r  r  r   unknownr  r  r  r  r  rH  rh  rq  ry  __classcell__r  s   @r7   rz  rz    s   EIFMOO##EL  E&&EE)7)4G_4=? :?-%- - 
	-%@ SW<< *<3><FO<	<*	8 "&$(##6;#6#6#8   "	
  ! 
,  &	
 5 
/ u2u2 u2 &	u2
 5u2 
/u2nHT&PhV BF[
[
"[
:>[
	[
z""&0"9="FJ"	"r9   rz  c                  L     e Zd ZU eZdZded<   d fdZ	 	 	 	 	 	 	 	 ddZ xZ	S )MetalSchedulingr   r'   _kernel_fn_counterc                $    t         |   |       y r   )r  r  )rF   	schedulerr  s     r7   r  zMetalScheduling.__init__  s    #r9   c           	        t         j                  j                  }||j                  v r|j                  |   S t         j                  j                  rXd|j                          }|}||j                  |<   d| d|z   }t        ||      \  }}|j                  ||| d| d       |S d| j                   }	| xj                  dz  c_        |	|j                  |<   d	}
|j                  |
      t        |
      z   }|j                  d
      }||| }g }|j                  d      D ]G  }|j                         j                  d      r#|j                  |j!                  dd|	 d             I t#        t%        |j&                              }|j(                  j+                  |	 d|	d       |D ]  }|j(                  j+                  |        |j(                  j+                  d| d       |	S )Nmps_lib_zconst char* z
_source = 
F)gpugenerated_kernel_r   r   r*  z#includezkernel void generated_kernel(r#  rL   z = async_compile.metal(z, '''z''', r?   )r   r  rV  src_to_kernelr+  next_kernel_suffixr   define_kernelr  r  rn   rindexsplitstrip
startswithr  r  rn  sortedr1  r=  r   )rF   src_codenode_scheduler   r`  mps_lib_namekernel_nameoriginsdetailed_originsr;  metal_src_startstartend	metal_src
body_linesr  headers_reprs                    r7   r  zMetalScheduling.define_kernel  s    ''&&w,,,((2277%g&@&@&B%CDL&K.9G!!(+%l^:>IH(;M7(S%G%!!h7)26F5G(He "   &d&=&=%>?1$*1h' 3/#o2FFoof%U3'	 
OOD) 	Dzz|&&z2<WIUV>WX	 F6>>23  G9,CG;e!TU 	+DNN$$T*	+  5a!89r9   )r  zScheduler | Noner   rv  )r  r5   r  zlist[SchedulerNode]r   rz  r   r5   )
r   r   r   rz  kernel_typer  r  r  r  r  r  s   @r7   r  r    s=    K$..,?.IT.	.r9   r  )r6   z&float | int | bool | str | CSEVariabler   r5   )I
__future__r   rn  r   loggingr/  pathlibr   typingr   r   r  sympy.printing.precedencer   r3   torch.utils._cpp_embed_headersr   torch.utils._ordered_setr	   torch.utils._sympy.functionsr
   torch.utils._sympy.printersr   r   ExprPrinter_torch.utils._sympy.value_rangesr   r}  r   r   r   virtualizedr   r   r   commonr   r   r   r   r   r   simdr   r   r   ops_handlerr   r    r  r!   r"   r#   	getLoggerr   r   r$   int8int16int32int64uint8r*   r+   r  r   r8   r;   r   _initialize_pointwise_overridesrr  rz  r  r]   r9   r7   <module>r     s'   #      %  0  9 / , O 7 G G , ,  C B 64g! 
JJ	JJ	KK	KK	KK	KK	KK	JJ	NNH
t"| t"nX[ Xv
  . .u 5  & & (N
* N
b5n 5r9   