
    Vh}             -       fp   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	m
Z
mZ d dlmZ d dlZd dlmZ d dlmZmZmZ d dlmZmZmZmZ d d	lmZ d d
lmZmZ d dlmZm Z m!Z!m"Z"m#Z#m$Z$m%Z%m&Z&m'Z'm(Z( d dl)m*Z*m+Z+m,Z,m-Z-m.Z. d dl/m0Z0m1Z1 d dl2m3Z4 d dl5m6Z7  e
d      Z8 ed      Z9ejt                  jv                  Z;ejx                  j{                  ddd      Z> e?d      \  Z@ZAZBdeee9e8f   gee9e8f   f   fdZCde#fdZDd ZEd ZF eCe;j                  e;j                  g       e.       dddej                  ddfd              ZJ eCe;j                  j                  e;j                  j                  g       e.       d               ZN eCe;j                  j                  e;j                  j                  g       e.       dd d!              ZO eCe;j                         e.       d"               ZP eCe;j                  j                  e;j                  j                  e;j                  j                  e;j                  j                  g       e.d#d$      d%               ZS eCe;j                  j                  e;j                  j                  g       e.       d&               ZTd' ZUdd(ed)eVeW   d*eXfd+ZY eCe;j                  j                  e;j                  j                  g       e.       d,               Z[dZ\d)eVeW   fd-Z] eCe;j                  j                  e;j                  j                  g       e.       d.               Z_ eCe;j                  j                        dd/d0       Zb eCe;j                  j                        ej                  dddd1d2       Zd eCe;j                  j                  e;j                  j                  g       e.       ej                  dddd1d3              Zf eCe;j                  j                  e;j                  j                  g       e.       ej                  dddd1d4              Zi eCe;j                  j                  e;j                  j                  g       e.       ddddd1d5              Zk eCe;j                  j                  e;j                  j                  g       e.       d(ed)eVeW   d6eWd7eWfd8              Zm eCe;j                  j                        dd9       Zod: Zp eCe;j                  j                        d;        Zr eCe;j                        	 	 	 dd<ed=ed>ed?e	e   d@e	et   dAe	ej                     fdB       Zv eCe;j                        	 ddCedDedEedAe	ej                     fdF       Zx eCe;j                        dGdGddHd<edCedDedEedAe	ej                     f
dI       Zz eCe;j                        	 	 	 	 	 	 	 ddJej(                  dKej(                  d?e	e   dLe	e   dAe	ej                     dMeXdNeWdOeWdPeXfdQ       Z| eCe;j                  j                        dRdSd(ed)eWdTedUej(                  dVetdWeXdefdX       Z~ eCe;j                  j                        dRdSd(ed)eWdTedUej(                  dVetdWeXdefdY       Z e.        eCe;j                  j                        dZ               Z eCe;j                  j                        dddd ddd[d\edVetd]e	e   d$e	e   d^e	e   d_eWd`eXdefda       Z eCe;j
                  j                  e;j
                  j                  g       e.       db               Z eCe;j
                  j                        ddc       Z eCe;j                  j                  e;j                  j                  g       e.       dd               Z eCe;j                  j                        dde       Z eCe;j                  j                        df        Z eCe;j                  j                        dg        Z eCe;j                   j                        dh        Z eCe;j                   j$                        di        Z eCe;j(                  j                        dj        Z eCe;j,                  j                        ddddddkdl       Z eCe;j0                  j                        ddm       Z eCe;j2                  j                        ddn       Z eCe;j6                  j                        ddo       Z eCe;j8                  j                        dp        Z eCe;j<                  j$                        dq        Zd(edretfdsZd(edteduetfdvZ	 ddwedretdxeXfdyZddtedretdzetfd{Zdted|ed}eXdretfd~Z	 ddetded<edetfdZdetfdZ eCe;jN                  j                  e;jN                  jP                  g       e.dd      ddtedetdeXfd              Z eCe;jT                  j                  e;jV                  j                  g       e.       d<edefd              Z eCe;jZ                  g       e.dd      d<efd              ZdedefdZ eCe;j`                         e.       d(edtedeXdefd              Z eCe;jb                         e.       dd(edtedeXdefd              Z eCe;jd                         e.       dd(edeXdefd              Z eCe;jf                         e.       dd(edeXdefd              Z eCe;jh                  j                        ddtedeXdeXfd       Z eCe;jj                  j                  e;jj                  j                  g       e.       d<ededefd              Z eCe;jl                  j                        ddtedeXfd       Z eCe;jp                  j                  e;jp                  j                  g       e.ddd      dddd(edeXdeXdeeeef   fd              Z eCe;jv                  j                  e;jv                  j                  g       e.       dddeded|edeXdef
d              Z eCe;jz                  j                  e;jz                  j                  g       e.ddd      dRddtedeXdeeeef   fd              Z eCe;j~                  j                  e;j~                  j                  g       e.ddd      dRdddtedeXdeXdeeeef   fd              Z eCe;j                  j                  e;j                  j                  g       e.       dRdddeded|ed}eXdeXdefd              Z eCe;j                         e.ddd      	 	 ddededeXdeXdeeeef   f
d              ZdetdeeXeXf   fdZ eCe;j                  j                  e;j                  j                  g       e.dd      ddtedetdeeef   fd              Z eCe;j                  j                  e;j                  j                  g       e.dddd      dtedeeeeef   fd              Z eCe;j                  j                        	 	 	 ddtedeXdeXde	et   fd       ZdededeeVeW   eVeW   f   fdZdededue	et   deeef   fdZd<ededeXfdZ eCe;j                        dRdddddddted|ed}eXdeXde	e   de	e   de	e   de	e   deeeeef   fd       Z eCe;j                  j                  e;j                  j                  g      dRddddted|edeXd}eXdeXde	e   defdń       Z eCe;j                         e.dddRȫ      	 	 	 dd(edtedeXdeXdeXdeeef   fdʄ              Z eCe;j                  j                        d˄        Z eCe;j                         e.       	 	 dd<ededed}eXdeXdefd̄              Zd̈́ Zd΄ Z eCe;j                         e.       dτ               Z eCe;j                         e.       dЄ               Zdф Z eCe;j                         e.dҫ      dӄ               Z eCe;j                         e.dҫ      dԄ               ZdՄ Z eCe;j                         e.       dք               Z eCe;j                         e.       dׄ               Z eCe;j                  j                  e;j                  j                  e;j                  j                  e;j                  j                  g       e.dҫ      d؄               Zdل Z eCe;j                         e.       dڄ               Z eCe;j                         e.       dۄ               Z eCe;j                  j                  e;j                  j                  e;j                  j                  e;j                  j                  g       e.dҫ      d܄               Z eCe;j                         e.       dd(ededefdބ              Z eCe;j                         e.       ded(edededef
d              Z eCe;j                  j                  e;j                  j                  g       e.       dGdGdd              Z eCe;j                  j                  e;j                  j                  g       e.       dd/d              Z eCe;j                  j                        d d       Z eCe;j                  j                        d d       Z eCe;j                   j                  e;j                   j                  g       e.       dd              Z eCe;j                  j                        	 	 dd       Z eCe;j                         e.       d               Zd Zd!dZ	 ddej(                  d=ej(                  deeVeW   eWf   deeVeW   eWf   deeVeW   eWf   deXdeWde	eeVeW   eWf      fdZd Z	 eCe;j                  j                        dej(                  d=ej(                  d?e	ej(                     de	ej(                     de	ej(                     deXdedefd       Z eCe;j                  j                        dej(                  d=ej(                  d?ej(                  deVeW   deVeW   deVeW   deXdeVeW   deWfd       Zej                  j                  r_ejx                  j{                  ddd      Z eCejt                  j"                  j$                  j                        d        Z eCejt                  j"                  j(                  j                        d        Zej                  j,                  rNejx                  j{                  ddd      Z eCejt                  j0                  j2                        d         Zejx                  j{                  ddd      Z eCejt                  j8                  j:                  j                        d        Z eCejt                  j8                  j:                  j>                        d        Z  eCejt                  j8                  jB                  j                         eCejt                  j8                  jB                  jD                        d               Z# eCejt                  j8                  jB                  j>                         eCejt                  j8                  jB                  jH                        d               Z% eCejt                  j8                  jL                  j                         eCejt                  j8                  jN                  j                        d               Z(ejx                  j{                  ddd      Z) eCejt                  jT                  jV                        	 	 	 	 d"d	       Z, eCejt                  jT                  jZ                        d
        Z.d Z/ eCe;j`                  j                        	 	 	 	 	 d#d       Z1d Z2 eCe;jf                  j                        d        Z4 eCe;jj                         e.       	 	 	 	 	 d#d              Z6 eCe;jn                         e.dҫ      d               Z8 eCe;jr                  j                        d        Z: eCe;jv                  j                        d        Z< eCe;jz                  j                        d        Z> eCe;j~                         e.dҫ      d               Z@dedzetfdZA eCe;j                         e.dd$      d               ZC eCe;j                         e.dҫ      d               ZE eCe;j                         e.dd$      d               ZG eCe;j                         e.dҫ      d               ZI eCe;j                  j(                        dd       ZK eCe;j                  j                  e;j                  j                  g       e.       d               ZM eCe;j                  j                  e;j                  j                  g       e.       dddeWfd              ZN eCejt                  jv                  j                  j                  ejt                  jv                  j                  j                  g       e.       d                ZO eCe;j                  j(                  e;j                  j(                  g      d!        ZR eCe;j                  j                  g      d"        ZT eCe;j                  j                  e;j                  j                  g       e.       dGdGdd#              ZV eCe;j                  j                  e;j                  j                  g      ddd$d%       ZY eCe;j                  j                  g      ddd$d&       Z[ eCe;j                  g       e.       d'               Z] eCe;j                  g      d(        Z_ eCe;j                  g      d)        Za eCe;j                  g      d*        Zc eCe;j                  g      d+        Zed,eWd-eWdeWfd.Zfd/ Zg eCe;j                  g      d?e	e   fd0       Zi eCe;j                  g      d1        Zk eCe;j                  g      d2        Zm eCe;j                  j                        d3        Zo eCe;j                         e.       d4               Zq eCe;j                  j                        	 	 	 	 	 	 d$d5       Zs eCe;j                  j                        d6        Zudd7Zv eCe;j                  j                  e;j                  j                  g       e.       d%dd8d9              Zx eCe;j                  j                  e;j                  j                  g      d:        Z{ eCe;j                  j                  e;j                  j                  e;j                  j                  e;j                  j                  e;j                  j                  e;j                  j                  g       e.d#d$      d&d;              Z eCe;j                   j                        d<        Z eCe;j                  j                        d=        Z eCe;j                  j                        d>        Z eCe;j                  j                  e;j                  j                  e;j                  j(                  e;j                  j(                  e;j                  j                  e;j                  j                  e;j                  j                  g      d?        Z eCe;j                  j                  e;j                  j                  e;j                  j(                  e;j                  j(                  g      dd@       Z eCe;j                   j                  e;j                   j"                  g      dA        ZdB Z eCe;j(                  j(                  e;j(                  j                  g      dC        Z eCe;j,                  j(                  e;j,                  j                  g      dD        Z eCe;j0                  j                        dE        Z eCe;j4                  j(                  e;j4                  j                  g      dF        Z eCe;j8                  j(                  e;j8                  j                  g      dG        Z eCe;j<                  j                        dH        Z eCe;j@                  j(                         e.       ddefdI              Z eCe;jD                  g       e.       	 d'dJ              Z eCe;jH                  g      	 d'dK       Z eCe;jL                  g      	 d'dL       Z eCe;jP                  j                  e;jR                  j                  g      ddM       Z eCe;jV                  j                        dN        Z eCe;jZ                  j                        dO        Z eCe;j^                        dP        Z eCe;jb                         e.       dQ               Z eCe;jf                        dR        Z eCe;jj                  j                        ddS       Z eCe;jn                  j                        dT        ZddUZ eCe;jt                  j                        dV        ZdW ZdX ZdY ZdZ Z	 dd<ed[eWd\eWd]eWd^eWd_eWd`eWdaeWdbeWdceWddeWdeeWdfeWdgeWdheWdieWdjeWdkeWdleWdmeWdetdneXf,doZdp Zd<eded[eWd\eWd]eWd^eWd_eWd`eWdaeWdbeWdceWddeWdheWdieWdjeWdkeWdleWdmeWdetf&dqZdr Z eCe;j                  j                        ds        Z eCe;j                  j                        	 	 	 	 d"dt       Z eCe;j                  j                        du        Z eCe;j                         e.dd$      	 	 	 	 d"dv              Z eCe;j                         e.dҫ      dw               Zd<edxefdyZ G dz d{e      Zd<edxed|eWfd}Z eCe;j                  j                        d~        Z eCe;j                         e.       d               Z eCe;j                         e.dҐd      d               Z eCe;j                  j                  g      d        Z eCe;j                  j                        	 	 	 	 	 d(d       Z eCe;j                  j                        d        Z eCe;j                  j                        d        Z eCe;j                  j                        d)d       Zܐdd)eWdeWdeXfdZݐd Zސd Z eCe;j                  j                        dd       ZddZddZd ZddZd*dZ eCe;j                  j                        d        Z eCe;j                        d        Z eCe;j                  j                  e;j                  j                  e;j                  j                  e;j                  j                  g       e.       dd              Z eCe;j                  j                  e;j                  j                  e;j                  j                  e;j                  j                  g      dd       Z eCe;j                  g      	 	 	 	 d+dededededeXdeXde	e   fd       Z eCe;j                  g      	 	 	 	 d+dededede	e   deXdedeXdeXde	e   fd       Z eCe;j                  g      	 	 	 	 	 d,dededede	e   dedeXdeXde	e   fd       Z eCe;j                  g      	 ddededededededededeWdeWdedeXdedede	e   fd       Z eCe;j                  g      	 	 	 	 d-dededededeXde	e   de	e   fd       Z eCe;j                  g      	 	 ddedededededededeXde	e   de	e   fd       Z eCe;j                  g      	 	 	 d.dededede	e   deXdeXde	e   fd       Z  eCe;j                  g      	 	 d/dedededede	e   dedededededeVeX   deXde	e   fd       Z eCe;j                  g      	 ddedededededededededededeWdeWdedeXde	e   f d       Z eCe;j
                  g      	 	 	 	 	 d(dededede	e   de	e   deWdeWdedeXdeXde	e   de	eW   de	eW   de	e   de	e   fd       Z eCe;j                  g      	 	 	 ddededededededededeWdeWdedeXdedede	e   de	eW   de	eW   f"d       Z eCe;j                  g      	 	 	 	 	 d0dededed?e	e   de	e   de	e   de	eW   de	eW   dedeWdeXde	e   de	e   de	e   de	eW   fd       Z
 eCe;j                  g      	 	 	 d*dedededed?e	e   de	e   de	e   dej                  dej                  dededededeWdeXde	e   de	eW   deXf$dĄ       Z eCe;j                  j                  g      	 	 	 	 d1d(ej(                  dEej(                  dej(                  dej(                  d?e	ej(                     de	ej(                     dAe	ej                     deXfdɄ       Z eCe;j                   j"                  e;j                   j$                  g       e.       ddʄ              Z eCe;j(                  j"                        dd˄       Z eCe;j,                  j                  e;j,                  j                  g       e.       ddd/d̄              Zd̈́ Zd΄ Z eCe;j4                  j                  e;j6                  j                  g      ddτ       Z eCe;j8                  j                  e;j:                  j                  g      ddЄ       Z eCe;j<                  j                  e;j>                  j                  g      	 	 ddedeeeWej                  f      deeeWej                  f      de	e   de	e   f
dՄ       Z eCe;j@                  j                  e;jB                  j                  g      ddք       Z  eCe;jD                  j                  e;jD                  jF                  e;jD                  j                  e;jD                  jH                  g      d2dׄ       Z%d؄ Z& eCe;jN                  j                        	 	 ddل       Z( eCe;jR                  j                        dڄ        Z) eCe;jT                  j                        dۄ        Z*d܄ Z+d݄ Z, eCe;jZ                  j                  e;j\                  j                  g      d%dބ       Z/ eCe;j`                  j                        d3d߄       Z0 eCe;jb                  j                        d4d       Z2 eCe;jf                         e.       	 d5d              Z4 eCe;jj                  j                  e;jj                  j                  g       e.d#d$      d&d              Z6ejn                  Z8d Z9 eCe;jt                  j                        d        Z: eCe;jv                  j                        d        Z; eCe;jx                  j                        d        Z= eCe;j|                  j                        d        Z> eCe;j~                  j(                  e;j~                  j                  g       e.       dddd              ZA eCe;j                  g       e.       d6d              ZC eCe;j                  j                  e;j                  j                  g      	 	 dd       ZF eCe;j                  j                  g      	 	 dd       ZH eCe;j                  j                        d        ZI eCe;j                  j                  e;j                  j                  g       e.       dd              ZJ eCejt                  jv                  j                        d        ZK eCejt                  jv                  j                        d        ZL eCe;j                         e.       dddddd              ZNd ZO eCe;j                        	 d7d       ZQ eCe;j                        	 d7d       ZS eCe;j                        	 d7d       ZU eCe;j                         e.       dddd              ZW eCe;j                         e.       deWd(edefd              ZY eCe;j                        d(efd       Z[ eCe;j                         e.dRȫ      d(edefd              Z\ eCe;j                         e.       d(edefd              Z] eCe;j                         e.       ded)eWdeXdefd               Z_ eCe;j                         e.       	 	 	 d8d=ed$edeWdeXdeXdefd              Z` eCe;j                  j                        	 d9d#ed^eVe   deVeW   defd       Zbd Zcd	 Zd ece;j                          ece;j                          ece;j                          ece;j                          ece;j                          ece;j                          ece;j                          ece;j                          ece;j                          ede;j                          ede;j                          ede;j                          ede;j                          ede;j                          ede;j                          ede;j                          ede;j                          ede;j                          ede;j                          ede;j                          ede;j                         d
 Zz eCe;j                         e.       d               Z{ eCe;j                         e.       dGdd              Z| eCe;j                         e.       dGdd              Z} eze;j                        Z~ eze;j                        Z eze;j                        Zd dl/Zd dlZd dlZd Z e        y(:      N)Sequence)Enum)wraps)CallableOptionalTypeVarUnion)	ParamSpec)SymBoolSymFloatTensor)_add_op_to_registry_convert_out_paramsglobal_decomposition_table
meta_table)
OpOverload)_prim_elementwise_meta$ELEMENTWISE_PRIM_TYPE_PROMOTION_KIND)
BoolLikecorresponding_complex_dtypecorresponding_real_dtypeelementwise_dtypesELEMENTWISE_TYPE_PROMOTION_KIND	FloatLikeIntLikemake_contiguous_strides_forNumber
TensorLike)_maybe_convert_to_dtype_maybe_resize_out_resize_output_check_safe_copy_outout_wrapper)_broadcast_shapes_maybe_broadcast)_config)_pytree_T_PatenIMPLMeta   returnc                       fd}|S )Nc                 V     t                 fd}t        j                  |        S )Nc                 (    t        t        |        y N)r   r   )opfns    I/home/dcms/DCMS/lib/python3.12/site-packages/torch/_meta_registrations.pyregisterz0register_meta.<locals>.wrapper.<locals>.register9   s    
B3    )r   pytree	tree_map_)r4   r6   r3   s   ` r5   wrapperzregister_meta.<locals>.wrapper6   s)     $	4 	2&	r7    )r3   r:   s   ` r5   register_metar<   5   s     Nr7   type_promotionc                     t        j                  |d| i\  }}|D cg c]  }t        ||       }}t        | }t	        |dt
        j                  iS c c}w )Ntype_promotion_kindr=   )utilsr   r   r%   r   r   DEFAULT)r=   args_result_dtypexs        r5   elementwise_metarF   B   ss    
 ..	*OA| ?CC#A|4CDC T"D "	BJJ  Ds   Ac                     t         j                  t         j                  t         j                  t         j                  t         j
                  t         j                  i}|j                  | |       S r2   )torch	complex32halfcfloatfloatcdoubledoubleget)dtypefrom_complexs     r5   toRealValueTyperR   V   sE    ekku||L
 E5))r7   c                 l     t        t         g|       t        j                   k(   fd       y )Nc                      d d  S )Nzoutput with shape z# doesn't match the broadcast shape r;   )broadcasted_shape
self_shapes   r5   <lambda>z)check_inplace_broadcast.<locals>.<lambda>c   s    $ZL0STeSfg r7   )tupler$   rH   _check)rV   
args_shaperU   s   ` @r5   check_inplace_broadcastr[   _   s0    /
HZHI	LLZ'gr7   Fc	                 <   	 t         t        j                        r(t        j                   j	                         dk(  d        t        t        j                        r(t        j                  j	                         dk(  d        t        d  fD              rZt        j                  t        j                               		nFt        j                  t        j                        	fd       nxs t        j                         t        t        j                        sJ t        j                  t        t               fd       t        t              sJ t        j                  dk\  d        t        j                  f|d||	      S )
Nr   c                       yNz:linspace only supports 0-dimensional start and end tensorsr;   r;   r7   r5   rW   z(meta_linspace_logspace.<locals>.<lambda>w       r7   c                       yr^   r;   r;   r7   r5   rW   z(meta_linspace_logspace.<locals>.<lambda>|   r_   r7   c              3   <   K   | ]  }t        |t                y wr2   )
isinstancecomplex).0args     r5   	<genexpr>z)meta_linspace_logspace.<locals>.<genexpr>   s     
C:c7#
Cs   c                      d  d S )Nzlinspace(): inferred dtype z& can't be safely cast to passed dtype r;   )default_complex_dtyperP   s   r5   rW   z(meta_linspace_logspace.<locals>.<lambda>   s    56K5LLrsxryz r7   c                      dt              j                   dt               j                   dt              j                   dS )Nz4received an invalid combination of arguments - got (, ))type__name__)endstartstepss   r5   rW   z(meta_linspace_logspace.<locals>.<lambda>   sD     u+r$s),,-RU0D0D/EQH r7   c                       y)Nz$number of steps must be non-negativer;   r;   r7   r5   rW   z(meta_linspace_logspace.<locals>.<lambda>   r_   r7   meta)rP   layoutdevice
pin_memoryrequires_grad)rb   rH   r   rY   dimanyr@   r   get_default_dtypeis_complex_dtyperP   _check_typer   empty)
ro   rn   rp   baserP   rt   rs   ru   rv   rh   s
   ``` `    @r5   meta_linspace_logspacer~   g   sL    %&IIK1P	
 #u||$GGINP	

 
CsE/B
CC % A A##%!
 =)ELL&&u-z
 2002eU[[))) 
5'"	H
 eW%%%	LL!KL;;	# r7   c                    t        j                  j                  t         j                  k(  fd       t        j                  | j                         dk(  xr j                         dk7   d        | j                  j                        S )Nc                  "    d j                    S )Nz2take(): Expected a long tensor for index, but got rP   indexs   r5   rW   zmeta_take.<locals>.<lambda>   s    DU[[MR r7   r   c                       y)Nz*take(): tried to take from an empty tensorr;   r;   r7   r5   rW   zmeta_take.<locals>.<lambda>   r_   r7   )rH   rY   rP   long_check_indexnumel	new_emptyshape)selfr   s    `r5   	meta_taker      sm     
LLuzz!R
 
ZZ\Q55;;=A#56< >>%++&&r7   rw   c                T     j                   }j                   }t        j                  ||k(  d        t        j                   j                        dk(  xr j                        dk(   fd       t	         j
                  j
                        } j                  |      S )Nc                       y)Nz=linalg.cross: inputs must have the same number of dimensions.r;   r;   r7   r5   rW   zlinalg_cross.<locals>.<lambda>   r_   r7   r-   c                  V    d  dj                          dj                          S )Nzlinalg.cross: inputs dimension z must have length 3. Got  and size)rw   otherr   s   r5   rW   zlinalg_cross.<locals>.<lambda>   s6    -cU 399S>"%

3'8: r7   )ndimrH   rY   r   r$   r   r   )r   r   rw   x_dy_d	out_shapes   ```   r5   linalg_crossr      s     ))C
**C	LLs
O 
LL		#!4

31 4	
 "$**ekk:I>>)$$r7   c                 |    t        | d       t        | d       t        j                  | t        j                        S )Nzlinalg.matrix_expmemory_format)squareCheckInputscheckFloatingOrComplexrH   
empty_likecontiguous_formatr   s    r5   linalg_matrix_expr      s3     d/04!45D0G0GHHr7   valuesindicesc                 Z   t        j                  | j                  | j                  | j                        }t        j                  | j                  | j                  t         j
                        }| j                         dk7  r%| j                  dk7  rt        || j                         ||fS )Nrt   rP   r   )	rH   r|   r   rt   rP   int64r   r   maybe_wrap_dim)r   rw   r   r   s       r5   	cummaxminr      sp    
 [[DKKtzzJFkk$**T[[LGzz|qTYY!^sDII&7?r7   c                 x    t        || j                         t        j                  | t        j                        S Nr   )r   r   rH   r   r   )r   rw   s     r5   logcumsumexpr      s+     3		"D0G0GHHr7   c                D   |j                   }t        |      }||z
  }t        t        |            }t        |      D 	cg c]  }	d }
}	|D ]  }d|
|<   	 g g }}|D ]*  }|
|   s|j	                  |       |j	                  |       , ||z   }t        |      }|j                         |d | }|j                  fdd       |||d  z   }|j                  |      }dgt        |j                  |d        z   }|j                  |      }|j                  d      }||d<   t        |      }t        t        |            D ]  }|||      ||dz   <    | j                  |t        j                         t        |      D 	cg c]  }	d }}	d}|dz
  }|dk\  r0|| j                  d      z  |||   <   ||||      z  }|dz  }|dk\  r0t        ||      D ]  }| j                  d||z
  z         |||   <   ! | j                  ||| j                                | S c c}	w c c}	w )	NFTc                     |    S r2   r;   )rE   self_stridess    r5   rW   z_exec_fft.<locals>.<lambda>   s    <? r7   keyreverser   r      r   )r   lenlistrangeappendstridesortpermuter   reshaper   resize_rH   r   as_strided_storage_offset)outr   	out_sizesrw   forwardr   signal_ndim
batch_dimsdim_permuterC   is_transformed_dimdleftright	batch_endtmpinputbatched_sizes
batch_sizebatched_out_sizesiout_stridesbatch_numelr   s                          @r5   	_exec_fftr      sZ   99Dc(K#J uT{#K).t5A%55 % $1% b%D !!$KKNLLO	
 ,KD	I;;=L
jy
!CHH*DH9IJ//KLL%E D4JK 899MMM-(EAJ!M!]+3s8_ 5#,SV#4!a% 5KK!1H1HKI $Dk*1*K*KQA
q&&1CJJqM&AKN#yQ00	Q q& :t$ G&)jja*n1E&FKN#GOOI{C,>,>,@AJW 6@ +s   	H 	Hr   rw   exclude_lastc                     t        |      }| j                         |d t        |      t        |      z
   j	                  fd       |S )Nc                     |    S r2   r;   )r   r   s    r5   rW   z_sort_dims.<locals>.<lambda>!  s    l1o r7   )r   )r   r   r   intr   )r   rw   r   sorted_dimsr   s       @r5   
_sort_dimsr     sL    s)K;;=L6#k"S%667<<% =  r7   c                 
   t        j                  | j                  j                         |s| j	                         S t        | |      }| j                  | j                               }t        || | j                         ||      S )Nr   )	rH   rY   rP   
is_complexcloner   r   r   r   )r   rw   normalizationr   r   r   s         r5   meta_fft_c2cr   (  sb     
LL&&'zz|T3'K
..
%CS$		['JJr7   c                 f    t        |       t        kD  st        |       dk\  r| d   dk(  r	| d   dk(  ryy)N   r   r   FT)r   cufft_max_ndimr   s    r5   use_optimized_cufft_pathr   7  s3    
3x. SX]s1v{s1vQR{r7   c                    t        j                  | j                  j                         t	        | j                               }t	        |      }|d   }||   dz  dz   }t	        |      }|||<   |r|||<   t        |       dk(  rl| j                  |t        j                  | j                              }	| }
t        |      rt        |	|
||d       nt        |      dk(  r|n|}t        |	|
||gd       t        |      dkD  r0| j                  |t        j                  | j                              }
|d d }|rx|
|	}
}	|
j                         |j                  fdd	       t        t         t        |            }|t        |      |z
  d  }t        |	|
||d       |d t        |      |z
   }|rx|s:|	j                  |      ||   k7  r#|
j#                  |t         j$                  
       |
}	|	S t        |       dk(  rNt'        | |d      }| j                  |t        j                  | j                              }t        || ||d      S | j                  |t        j                  | j                              S )Nr   r   r   cudar   Tr   c                     |    S r2   r;   )r   stridess    r5   rW   zmeta_fft_r2c.<locals>.<lambda>d  s    '!* r7   r   r   xpu)r   )rH   rY   rP   is_floating_pointr   r   device_hintr   r@   r   r   r   r   r   r   minr   r   r   r   )r   rw   r   onesidedinput_sizesr   last_dimlast_dim_halfsizeonesided_sizesoutputworking_tensortarget_sizesr   max_dims	last_dimsr   r   s                   @r5   meta_fft_r2cr   >  sx    
LL--.tyy{#K[!I2wH#H-2Q6+&N0N8/	(4F"U>>tzzJ   
 #C(fnidK ),CA9>LfnlXJPTU3x!|!%U%F%Ftzz%R "0 "
 cr(K)7(//1  ,d !  ~s;/?@'K(88(C(EF	NNIt **GC,<x,GH  {{8$	((;;&&y@W@W&X'	T	e	# s>nnU>>tzzJ  
 dI{DII~~U>>tzzJ  
 	
r7   )	generatorc                B    t        |t        j                  | g            S r2   )r    rH   Size)nr   r   s      r5   meta_randpermr     s    S%**aS/22r7   rP   rs   rt   ru   c                6    t        j                  | ||||      S Nr   rH   r|   )r   rP   rs   rt   ru   s        r5   meta_randperm_defaultr     s      ;;	vf r7   c                x     dt        j                   kD   fd       t        j                  |||||      S )Nr   c                      d d  S Nz:random_ expects 'from' to be less than 'to', but got from=z >= to=r;   highlows   r5   rW   zmeta_randint.<locals>.<lambda>      LSEQXY]X^_ r7   r   rH   rY   r|   )r  r   rP   rs   rt   ru   r  s   `     @r5   meta_randintr    s>     C	LLs
_ ;;E&J r7   c                t     t        j                   kD   fd       t        j                  |||||      S )Nc                      d d  S r  r;   r  s   r5   rW   z"meta_randint_low.<locals>.<lambda>  r  r7   r   r  )r  r  r   rP   rs   rt   ru   s   ``     r5   meta_randint_lowr    s9     
LLs
_ ;;E&J r7   c                6    t        j                  | ||||      S r   r   )r   rP   rs   rt   ru   s        r5   meta_rand_defaultr    s      ;;E&J r7   r   lastdimc                    t        j                  | j                  j                         t	        |       dk(  rt        | j                               }|||d   <   | j                  |t        | j                              }t        |      r.t        || j                  t         j                        ||d      S t        |      dkD  rt        | |d d d|      }n | j                  t         j                        }t        ||||d   gd      S | }t        |      dkD  r|d d }t        | ||d      }|dd  }t        |j                               }|||d   <   | j                  |t        | j                              }	t        |	|||d      S )	Nr   r   r   r   Fr   r   r   )rH   rY   rP   r   r   r   r   r   rR   r   r   r   r   r   r   )
r   rw   r   r  r   r   tempr   c2c_dimsr   s
             r5   meta_fft_c2rr    sk    
LL&&'4F"%	$	#b'	1LM#C(

)@)@
A  3x!|#D#cr(Aw?zz0G0GzHVT9s2wiOO s8a<3BxH xNEbc(C&	$	#b'nnYodjj.InJeYUCCr7   c                 J   ddl m}  ||       s#t        j                  |       dk(  rt	        d      t        |t              ra|j                  | |      }| j                         |j                         k7  r.t        j                  j                  || j                                | S )Nr   )free_unbacked_symbolsr   zQmore than one element of the written-to tensor refers to a single memory location)%torch.fx.experimental.symbolic_shapesr  rH   _debug_has_internal_overlapRuntimeErrorrb   r   tor   r*   expand_copydefault)r   srcnon_blockingr  intermediates        r5   
meta_copy_r    s     L "$'E,M,Md,SWX,X_
 	
 #vvvdL199;,++--$$\499;?Kr7   c                     t        | j                               }t        | j                               }|| j                         k\  rdn
||   ||   z  }|j	                  |d       |j	                  ||       ||fS Nr   )r   r   r   rw   insert)tensorrw   result_sizesresult_strides
new_strides        r5   inferUnsqueezeGeometryr&    sq    &L&--/*NVZZ\)|C/@>RUCV/VJQ#z*''r7   c                     t        || j                         dz         }t        | |      \  }}| j                  ||       | S r   )r   rw   r&  r   )r   rw   g_sizes	g_stridess       r5   meta_unsqueeze_r*    s>    
dhhj1n
-C/c:GYWi(Kr7   r   weight_metabias_activation_opt	out_dtypec                 8   t        | j                        }|*|j                  d      |j                  d      k(  sJ d       |j                  d      | j                  d      dz  k(  sJ |j                  d      |d<   t        | j                        dk(  sJ d       d| j                  d      f}|7| j                  t
        j                  k(  r|t
        j                  k(  sJ d       | j                  ||| j                  n|      j                  ||      }|S )	Nr   zoutput size mismatchr   r   r   z*we can only handle the squashed input case9out_dtype is only supported for i8i8->i32 linear operatorr   )
r   r   r   r   rP   rH   int8int32r   
as_strided)	r   r+  r,  r-  r.  r/  output_sizestransposed_stridesr   s	            r5   meta_sparse_structured_linearr7    s    $L{{1~1-E/EE-;;q>UZZ^a////{{1~L u{{q N"NN UZZ]+{{ejj(Y%++-E 	
G	
E __&.ekkI   j12 
 Mr7   mat1	mat1_metamat2c                    t        | j                        dk(  sJ t        |j                        dk(  sJ t        |j                        dk(  sJ | j                  d      |j                  d      dz  k(  sJ | j                  d      |j                  d      g}|7|j                  t        j
                  k(  r|t        j                  k(  sJ d       |j                  |||j                  n|      }|S )Nr   r   r   r1  r   r   r   r   rP   rH   r2  r3  r   )r8  r9  r:  r/  r5  r   s         r5   meta_sparse_structured_mmr=  :  s     tzz?ay1$$$tzz?a99Q<499Q<!++++IIaL$))A,/LzzUZZ'I,D 	
G	
D ^^%-djj9  F
 Mr7   r   )alphabetar/  c                |   t        | j                        dk(  sJ d       t        |j                        dk(  sJ t        |j                        dk(  sJ t        |j                        dk(  sJ | j                  d      |j                  d      k(  sJ d       |j                  d      |j                  d      dz  k(  sJ |j                  d      |j                  d      g}|7|j                  t        j
                  k(  r|t        j                  k(  sJ d       |j                  |||j                  n|      }|S )Nr   zEonly input broadcasted to columns of mat1 * mat2 product is supportedr   r   r1  r   r<  )	r   r8  r9  r:  r>  r?  r/  r5  r   s	            r5   meta_sparse_structured_addmmrA  S  s/    u{{q  O  tzz?ay1$$$tzz?a::a=DIIaL( O( 99Q<499Q<!++++IIaL$))A,/LzzUZZ'I,D 	
G	
D ^^%-djj9  F
 Mr7   compressed_Adense_Br>  transpose_resultalg_idsplit_ksplit_k_one_kernelc	                 L   |j                   t        j                  t        j                  t        j                  t        j
                  t        j                  hv sJ d       | j                   |j                   k(  sJ d       t        |j                        dk(  sJ d       | j                   t        j
                  t        j                  fv }	|	rdnd}
|	r|j                         rJ d       |j                  d      }|j                  d	      }| j                         d
z  |
|z  z  }|||j                  d      k(  sJ |I|	r@|t        j                  t        j                  t        j                  t        j                  hv sJ d       |r||fn||f}|j                  ||      S )Nz;_cslt_sparse_mm only supports fp16, bf16, int8, and fp8e4m3zinputs must have the same dtyper   z'_cslt_sparse_mm only supports 2d inputs
   	   z.dense input must be transposed for 8bit dtypesr   r      z\out_dtype is not supported for {compressed_A.dtype} x {dense_B.dtype} -> {out_dtype} matmul!r   )rP   rH   float32float16bfloat16r2  float8_e4m3fnr   r   is_contiguousr   r   r3  r   )rB  rC  r-  r>  r/  rD  rE  rF  rG  is_8bit_input_typecompression_factorkr   moutput_shapes                  r5   meta__cslt_sparse_mmrV  v  s    ==

  E EE  .Q0QQ.w}}"M$MM"%++

E<O<O/PP1q((* 	
<	
* 	QAQA					"(:Q(>?ADIIaL   !iMMNNKK	4
 '
 	
 k	
 
 .Aq6Aq6L\;;r7   T)include_selfr   sourcereducerW  c                L    t        j                  | t         j                        S r   rH   r   r   r   rw   r   rX  rY  rW  s         r5   meta_index_reducer]    s     D0G0GHHr7   c                    | S r2   r;   r\  s         r5   meta_index_reduce_r_    s	     Kr7   c                     t        | j                               }| j                         dkD  r|j                         ||<   | j	                  |      S Nr   )r   r   rw   r   r   )r   rw   r   result_sizes       r5   meta_index_selectrc    s@     tyy{#KxxzA~ ;;=C>>+&&r7   )lengthsr   offsetsaxisunsafeinitialdatard  re  rf  rg  c                     |t        d       fd}| ||j                        S |+|j                  d d |j                  d   dz
  fz   }	 ||	      S t        d      )Nz?segment_reduce(): indices based reduction is not supported yet.c                     t        j                  | j                  dz   d  z   j                  dt         j                        S )Nr   rr   rP   rt   r   )rH   r|   r   rP   r   )lengths_shaperf  ri  s    r5   segment_reduce_lengths_tensorz:meta_segment_reduce.<locals>.segment_reduce_lengths_tensor  s>    {{DJJtaxz22**11	
 	
r7   r   r   z<segment_reduce(): Either lengths or offsets must be defined.)NotImplementedErrorr   r  )
ri  rY  rd  r   re  rf  rg  rh  rn  rm  s
   `    `    r5   meta_segment_reducerp    s|     !M
 	

 ,W]];; cr*gmmB.?!.C-EE,];;
U
VVr7   c                 $    | j                  d      S Nr;   r   r   s    r5   meta_maxrt         >>"r7   c                     t        j                  | j                  |f      }t        | ||      }| j	                  |      | j	                  |t
        j                        fS Nr   r@   reduction_dimsr   _compute_reduction_shaper   rH   r   r   rw   keepdimrU  s       r5   meta_max_dimr}    R    


tzzC6
2C+D#w?L|$|5::6 r7   c                 $    | j                  d      S rr  rs  r   s    r5   meta_minr    ru  r7   c                     t        j                  | j                  |f      }t        | ||      }| j	                  |      | j	                  |t
        j                        fS rw  rx  r{  s       r5   meta_min_dimr    r~  r7   c                     | j                         rt        | j                        }nt        | t        j
                        \  }}t        j                  | |      S Nr?   r   )r   r   rP   r   r   INT_TO_FLOATrH   r   )r   rD   rC   s      r5   
meta_angler    sI    /

;, ? L L
< D55r7   c                     t        j                  || j                         | j                         |j	                  t        j
                  |             S r2   )rH   _resize_output_r   rt   copy_angle)r   r   s     r5   meta_angle_outr    s6    	#tyy{DKK899U[[&''r7   c                      y r2   r;   )vals    r5   assert_asyncr  !      
r7   c                      y r2   r;   )r  
assert_msgs     r5   assert_async_metar  &  r  r7   c                      y r2   r;   )ss    r5   
print_metar  +  r  r7   rP   rs   rt   ru   r   c                 0    t        j                  dd      S )Nr   rr   rt   r   r  s        r5   make_dep_tokenr  0  s     ;;q((r7   c                 h    ddl m} t        | t        t        f      rt        d       || ||       y )Nr   )constrain_range'Constraining SymFloat or Symbool is nyir   max)r  r  rb   r   r   
ValueError)r   r   r  r  s       r5   sym_constrain_ranger  <  s/     F$7+,BCCDcs+r7   c                 6    t         j                  | ||       |S Nr  )r*   r  r   r   r  	dep_tokens       r5   functional_sym_constrain_ranger  F  s    Ts4r7   c                 (   ddl m} ||t        j                  |        y t	        | t
        t        f      rt        d      t        |       t        u r5|t        j                  | |k\         |t        j                  | |k         y  || ||       y )Nr   )_constrain_range_for_sizer  r  )r  r  rH   _check_is_sizerb   r   r   r  rl   r   rY   )r   r   r  r  s       r5   sym_constrain_range_for_sizer  L  s     P
{s{T"$7+,BCCDzS?LL%?LL%d5r7   c                 6    t         j                  | ||       |S r  )r*   r  r  s       r5   'functional_sym_constrain_range_for_sizer  `  s    %%d%=r7   c                     |S r2   r;   )r  r  r  s      r5   functional_assert_async_metar  f  s    r7   f_namec                     | j                         dk\  s
J | d       | j                  d      | j                  d      k(  s.J | d| j                  d       d| j                  d       d       y )Nr   z3: The input tensor must have at least 2 dimensions.r   z5: A must be batches of square matrices, but they are  by 	 matrices)rw   r   )r   r  s     r5   r   r   l  s}    88:? (EF? 99R=DIIbM) (G		RTVZ[_[d[deg[hZiirs)r7   Anamec                     t        j                   j                  j                  k(   fd       t        j                   j                  j                  k(   fd       t        j                  j	                  d      j	                  d      k(  fd       t        j                  j	                  d       j	                  d      k(   fd       y )Nc                  >    dj                    d j                    dS )Nz:Expected b and A to be on the same device, but found b on z
 and A on 	 instead.r  r  r   s   r5   rW   z(linearSolveCheckInputs.<locals>.<lambda>{  s%    H{{m:ahhZy: r7   c                  >    dj                    d j                    dS )Nz=Expected b and A to have the same dtype, but found b of type z and A of type r  r   r  s   r5   rW   z(linearSolveCheckInputs.<locals>.<lambda>  s%    Kzzl/!'')= r7   r   r  c                  R    d j                  d       d j                  d       dS )Nz3A must be batches of square matrices, but they are r  r  r   r  r   r  s   r5   rW   z(linearSolveCheckInputs.<locals>.<lambda>  s0    FF2J<tAFF2J<yB r7   c                      d d j                  d       d j                  d       dj                  d       dj                  d       
S )NzIncompatible matrix sizes for z: each A matrix is r   r  z but each b matrix is r  r   )r  r  r   s   r5   rW   z(linearSolveCheckInputs.<locals>.<lambda>  sR    ,TF 3D$TYYr]O4		"H r7   )rH   rY   rt   rP   r   )r   r  r  s   ```r5   linearSolveCheckInputsr  x  s    	LLqxx	
 
LL

agg	
 
LL	r
affRj 	
 
LL	r
diim#	
r7   tallow_low_precision_dtypesc                 J   | j                   t        j                  | j                         xs | j	                         fd       |sYt        j                  t        j
                  t        j                  t        j                  t        j                  fv fd       y y )Nc                       d  S )Nz<: Expected a floating point or complex tensor as input. Got r;   rP   r  s   r5   rW   z(checkFloatingOrComplex.<locals>.<lambda>  s    6(VW\V]^ r7   c                       d  S )Nz*: Low precision dtypes not supported. Got r;   r  s   r5   rW   z(checkFloatingOrComplex.<locals>.<lambda>  s    vhHP r7   )	rP   rH   rY   r   r   rL   rN   rK   rM   )r  r  r  rP   s    ` @r5   r   r     sn    
 GGE	LL	/^ &ekk5<<u}}MMP	
 &r7   arg_namec                 ^    t        j                  | j                         dk\  fd       y )Nr   c                       d  dS )Nz: The input tensor z! must have at least 2 dimensions.r;   )r  r  s   r5   rW   zcheckIsMatrix.<locals>.<lambda>  s    6(-hZ7XY r7   )rH   rY   rw   )r  r  r  s    ``r5   checkIsMatrixr    s    	LL	1Yr7   Br   c                      t                t               t        j                  r# j	                  d      j	                  d      k(  n" j	                  d      j	                  d      k(   fd       y )Nr  r   c                       drdnd d j                  d       d j                  d       dj                  d       dj                  d       d	S )
Nz2: Incompatible shapes of A and B for the equation zAX = BzXA = Bz (r  rE   r   r   rk   r   )r  r  r  r   s   r5   rW   z#checkInputsSolver.<locals>.<lambda>  s[    hHxX.AaffRj\qvvbzl!AFF2J<qJ r7   )r   r  rH   rY   r   )r  r  r   r  s   ````r5   checkInputsSolverr    sY    a !V	LL$(r
affRj affRjAFF2J.F	
r7   resultfn_nameresult_namec                 r     t        j                  j                  j                  k(   fd       y )Nc            	      L      d d dj                    dj                    	S )Nz: Expected z5 and input tensors to be on the same device, but got z on z and input on r  )r  r   r  r  s   r5   rW   z!checkSameDevice.<locals>.<lambda>  s5    i{;-/dm4nU\\NL r7   )rH   rY   rt   )r  r  r   r  s   ````r5   checkSameDevicer    s&     
LL%	
r7   UPLOc                       j                         }t        j                  t               dk(  xr |dk(  xs |dk(   fd       y )Nr   ULc                      d  S )Nz1Expected UPLO argument to be 'L' or 'U', but got r;   )r  s   r5   rW   zcheckUplo.<locals>.<lambda>  s    CD6J r7   )upperrH   rY   r   )r  UPLO_uppercases   ` r5   	checkUplor    s<    ZZ\N	LLD	QKNc1J^s5JJr7   eigenvalueseigenvectorsr  	compute_vc                 T   t        | d       t        |       t        | j                        }|r/| j	                  |      }|j                  |t        |d             n| j	                  dg      }|j                          | j	                  |t        | j                              }||fS )Nzlinalg.eighF	row_majorr   r   )
r   r  r   r   r   r   r   poprR   rP   )r  r  r  r   vecsvalss         r5   meta__linalg_eighr    s     a'dOME{{5! ;EU ST{{A3	IIK;;uOAGG$<;=D:r7   c                     t        | d       t        j                  | j                        r| j                  nt        j                  | j                        }| j                  | j                  d d |      S )Nzlinalg.eigvalsr   r   r   r@   rz   rP   r   r   r   )r   complex_dtypes     r5   meta__linalg_eigvalsr    sc     e-. !!%++. 	..u{{; 
 ??5;;s+=?AAr7   c                 0   t        | d       t        j                  | j                        r| j                  nt        j                  | j                        }| j                  | j                  d d |      }| j                  | j                  |      }||fS )Nz
linalg.eigr   r   r  )r   r  r   vectorss       r5   meta_linalg_eigr    s     e\* !!%++. 	..u{{; 
 __U[["-]_CFooekko?G7?r7   r  c                 v    | j                   j                  t        j                        j	                  dd      S )Nr   r  r   )mTr   rH   r   	transpose)r  s    r5   cloneBatchedColumnMajorr    s*    66<<e&=&=<>HHRPPr7   r  c                     t        |       S r2   )r  )r   r  r  s      r5   _cholesky_solve_helperr  
  s     #4((r7   c                      t        j                   j                  dk\   fd       t        j                  j                  dk\  fd       t         d      \  }}t	        |||      S )Nr   c                  $    d j                    dS )Nz-b should have at least 2 dimensions, but has  dimensions insteadr   r   s   r5   rW   z cholesky_solve.<locals>.<lambda>  s    ?		{J]^ r7   c                  $    d j                    dS )Nz-u should have at least 2 dimensions, but has r  r  r  s   r5   rW   z cholesky_solve.<locals>.<lambda>  s    ?xGZ[ r7   cholesky_solve)rH   rY   r   !_linalg_broadcast_batch_dims_namer  )r   r  r  self_broadcastedA_broadcasteds   ``   r5   r  r    sh     
LL		Q^ 
LL	![ 'Ha!'#m ""2M5IIr7   c                     | j                         dk(  r%t        j                  | t        j                        S t	        | d       t        |       S )Nr   r   cholesky)r   rH   r   legacy_contiguous_formatr   r  r   r  s     r5   r  r  !  s@     zz|qE4R4RSSdJ'"4((r7   c                 0    t        | d       t        |       S )Ncholesky_inverse)r   r  r  s     r5   r   r   *  s     d./"4((r7   check_errorsc                    t        | d       t        | d       | j                  }t        |      }t	        |d      }| j                  |      }|j                  ||       | j                  |d|dz
   t        j                        }||fS )Nzlinalg.choleskyFr   r   r   )	r   r   r   r   r   r   r   rH   r3  )r  r  r  A_shaper   	L_stridesr  infoss           r5   linalg_cholesky_exr  2  s    a*+1/0ggGw<D ,GU;I	GAMM'9% KKD1H-U[[KAEe8Or7   tauc                 @    t        j                   j                  dk\  d        t        j                   j                  d       j                  d      k\  d        t        j                   j                  d      j                  d      k\  d        t        j                   j                  j                  z
  dk(   fd        j                  dkD  r: j                  d d }j                  d d t        j                  |k(  fd	       t        j                  j
                   j
                  k(   fd
       t        d d       t        j                   j                  t         j                  d       j
                   j                        S )Nr   c                       y)NzHtorch.linalg.householder_product: input must have at least 2 dimensions.r;   r;   r7   r5   rW   z,linalg_householder_product.<locals>.<lambda>K  r_   r7   r  r   c                       y)Nzbtorch.linalg.householder_product: input.shape[-2] must be greater than or equal to input.shape[-1]r;   r;   r7   r5   rW   z,linalg_householder_product.<locals>.<lambda>O  r_   r7   c                       y)Nz`torch.linalg.householder_product: input.shape[-1] must be greater than or equal to tau.shape[-1]r;   r;   r7   r5   rW   z,linalg_householder_product.<locals>.<lambda>S  r_   r7   r   c                  <    dj                    d j                    S )Nzptorch.linalg.householder_product: Expected tau to have one dimension less than input, but got tau.ndim equal to  and input.ndim is equal to r  r   r  s   r5   rW   z,linalg_householder_product.<locals>.<lambda>X  '    )),
2Nuzzl\ r7   c                      d  S )Nzltorch.linalg.householder_product: Expected batch dimensions of tau to be equal to input.shape[:-2], but got r;   actual_batch_tau_shapes   r5   rW   z,linalg_householder_product.<locals>.<lambda>b      66L5MO r7   c                  <    dj                    d j                    S )Nz,torch.linalg.householder_product: tau dtype z does not match input dtype r   r  s   r5   rW   z,linalg_householder_product.<locals>.<lambda>j  s#    :399+*5;;-9 r7   z torch.linalg.householder_productr  Fr  r   r   rP   rt   )
rH   rY   r   r   r   rP   r  empty_stridedr   rt   )r   r  expected_batch_tau_shaper  s   `` @r5   linalg_householder_productr  D  sK   
 
LL

aZ 
LL

2%**R.(t 
LL

2#((2,&r
 
LL

SXX"	
 zzA~#(;;s#3 !$3B"&>>	
 
LL		U[[ 	
 6UEJ[[*5;;%Hkk||	 r7   c                 2   t        | d       t        | dd       | j                  | j                        }|j	                  | j                  t        | j                  d             | j                  | j                  d d t        j                        }||fS )Nzlinalg.inv_exF)r  r  r  r   r   r   r   r   r   r   rH   r3  )r  r  r  r  s       r5   linalg_inv_ex_metar  z  sq    a)1o%P	AGGAMM!''6qww%PQKKEKKK8Ee8Or7   LDpivotsinfo)	hermitianr  r  c                   t        | d       t        | d       t        j                  | j                  t        | j                  d      | j                  | j                        }| j                  | j                  d d t        j                        }| j                  | j                  d d t        j                        }|||fS )Nztorch.linalg.ldl_factor_exFr  r  r   r   r  )
r   r   rH   r  r   r   rP   rt   r   r   )r   r  r  r  r  r  s         r5   linalg_ldl_factor_ex_metar!    s     d894!=>			ZZ*4::Gjj{{	
B ^^DJJsO599^=F>>$**Sb/>;Dvtr7   )r  c                d    t         d       t         d       t         d       t        j                  j
                  dk\  fd        j                  d d }t        j                  |j                  k(  fd       t        j                  t        j                  j                        fd       t        j                   j                  j                  k(   fd       t               \  }}t        j                  |t        |d	      j                  j                  
      S )Nztorch.linalg.ldl_solver   c                  $    d j                    dS )NzMtorch.linalg.ldl_solve: Expected B to have at least 2 dimensions, but it has r  r  )r  s   r5   rW   z'linalg_ldl_solve_meta.<locals>.<lambda>      &&!46 r7   r   c                  $    d j                    dS )Nzjtorch.linalg.ldl_solve: Expected LD.shape[:-1] and pivots.shape to be the same, but got pivots with shape  insteadr   r  s   r5   rW   z'linalg_ldl_solve_meta.<locals>.<lambda>      ))/h@ r7   c                  "    d j                    S )Nz<torch.linalg.ldl_solve: Expected pivots to be integers. Got r   r(  s   r5   rW   z'linalg_ldl_solve_meta.<locals>.<lambda>  s    Nv||n] r7   c                  <    dj                    d j                    S )Nz!torch.linalg.ldl_solve: LD dtype z does not match b dtype r   )r  r  s   r5   rW   z'linalg_ldl_solve_meta.<locals>.<lambda>  s"    3BHH:=UVWV]V]U^_ r7   Fr  r  )r   r   r  rH   rY   r   r   r@   is_integer_dtyperP   _linalg_broadcast_batch_dimsr  r   rt   )r  r  r  r  expected_pivots_shapeB_broadcast_sizerC   s   ```    r5   linalg_ldl_solve_metar0    s     b232781b":;	LL	!	
 HHSbM	LL-	
 
LLv||,] 
LL
AGG_ 7q"=a*+;uMggxx	 r7   Pr  )pivotr2  c                h    t        j                   j                  dk\   fd       t         j                        }|d   }|d   }t        ||      }||d<   |r j                  |      }n j                  dg      }||d<    j                  |      }||d<   ||d<    j                  |      }|||fS )Nr   c                  $    d j                    dS )Nz@linalg.lu: Expected tensor with 2 or more dimensions. Got size: r&  r'  r  s   r5   rW   z linalg_lu_meta.<locals>.<lambda>  s    RSTSZSZR[[cd r7   r  r   r   )rH   rY   r   r   r   r   r   )	r  r2  sizesrT  r   rS  r1  r  r  s	   `        r5   linalg_lu_metar6    s     
LL	!d
 MEb	Ab	AAq	AE"IKKKKE"I	EAE"IE"I	EAa7Nr7   LU)r2  r  c                    t        j                   j                  dk\   fd       t         j                        }|d   }|d   }t        j
                  |t        |d       j                   j                        }|j                          t        ||      |d<    j                  |t         j                        }|j                           j                  |t         j                        }|||fS )	Nr   c                  $    d j                    dS )NzFtorch.lu_factor: Expected tensor with 2 or more dimensions. Got size: r&  r'  r  s   r5   rW   z*linalg_lu_factor_ex_meta.<locals>.<lambda>  s    XYZY`Y`Xaaij r7   r  r   Fr  r  r   )rH   rY   r   r   r   r  r   rP   rt   r  r   r   r   )	r  r2  r  r5  rT  r   r7  r  r  s	   `        r5   linalg_lu_factor_ex_metar:    s     
LL	!j
 MEb	Ab	A			*5EBggxx	
B 
IIKAq	E"I[[eii[0F 
IIK;;uEII;.Dvtr7   )r   adjointr;  c                    t         d       t        j                   j                  j                  k(   fd       t        j                  j                  t        j                  k(  d        t         d       t         |d       t        j                   j                  d      j                  d      k(  d        t        j                   j                  d d j                  k(  fd       t               \  }}t        j                  |t        ||       j                  j                  	      }|j                         d
k7  r"|s |j                         r|j                         }|S )Nztorch.linalg.lu_solvec                  >    dj                    d j                    dS )NzPlinalg.lu_solve: Expected LU and B to have the same dtype, but found LU of type  and B of type r&  r   )r  r7  s   r5   rW   z&linalg_lu_solve_meta.<locals>.<lambda>  s(    $$&HH:_QWWIXO r7   c                       y)NzElinalg.lu_solve: pivots should be a Tensor of scalar type torch.int32r;   r;   r7   r5   rW   z&linalg_lu_solve_meta.<locals>.<lambda>  r_   r7   zlinalg.lu_solver   c                       y)NzYlinalg.lu_solve: Number of pivots per batch should be same as the dimension of the matrixr;   r;   r7   r5   rW   z&linalg_lu_solve_meta.<locals>.<lambda>$  r_   r7   c                  $    d j                    dS )Nzclinalg.lu_solve: Expected LU.shape[:-1] and pivots.shape to be the same, but got pivots with shape r&  r'  r(  s   r5   rW   z&linalg_lu_solve_meta.<locals>.<lambda>*  r)  r7   r  r  r   )r   rH   rY   rP   r   r   r  r   r   r-  r  r   rt   r   r   conj)r7  r  r  r   r;  r/  rC   r  s   ```     r5   linalg_lu_solve_metarC    s.    267	LL
AGG	
 
LL		!W b12b!T#45	LL
v{{2&k 
LL
"%	
 7q"=a  *+;4xPggxx	F ||~4[[]FMr7   unpack_dataunpack_pivotsc                     t        j                   j                  dk\   fd       |r2t        j                  |j                  t         j                  k(  d        t         j                        }|d   }|d   }t        ||      }||d<   |r j                  |      }n j                  dg      }|r2||d<    j                  |      }	||d<   ||d<    j                  |      }
n$ j                  dg      }	 j                  dg      }
||	|
fS )Nr   c                  $    d j                    dS )NzFtorch.lu_unpack: Expected tensor with 2 or more dimensions. Got size: r&  r'  )r7  s   r5   rW   z lu_unpack_meta.<locals>.<lambda>J  s    XY[YaYaXbbjk r7   c                       	 y)Nztorch.lu_unpack: LU_pivots is expected to be a contiguous tensor of torch.int32 dtype.
Note: this function is intended to be used with the output produced by torch.linalg.lu_factorr;   r;   r7   r5   rW   z lu_unpack_meta.<locals>.<lambda>O  s    p r7   r  r   r   )	rH   rY   r   rP   r3  r   r   r   r   )r7  r  rD  rE  r5  rT  r   rS  r1  r  r  s   `          r5   lu_unpack_metarI  @  s     
LL
1k LLEKK'	
 NEb	Ab	AAq	AE"ILLLL!b	LLb	b	LLLL!LL!a7Nr7   modec                       dk(  rd}d}||fS  dk(  rd}d}||fS  dk(  rd}d}||fS t        j                  d fd       fS )NreducedTcompleteFrc                      d  dS )Nzqr received unrecognized mode 'z=' but expected one of 'reduced' (default), 'r', or 'complete'r;   )rJ  s   r5   rW   z _parse_qr_mode.<locals>.<lambda>w  s    1$ 8N O r7   rH   rY   )rJ  	compute_qrL  s   `  r5   _parse_qr_moderR  j  s    y	 g 
		 g 
	 g 		
 gr7   QRc                    t        | d       t        | d       t        |      \  }}| j                  d   }| j                  d   }t	        ||      }|rMt        | j                        }|r|n||d<   | j                  |      }|j                  |t        |d             n| j                  dg      }t        | j                        }	|s|s|n||	d<   | j                  |	      }
|
j                  |	t        |	d             ||
fS )Nz	linalg.qrr  r   Fr  r   )	r  r   rR  r   r   r   r   r   r   )r  rJ  rQ  reduced_moderT  r   rS  Q_shaperS  R_shaperT  s              r5   linalg_qr_metarY    s     ![!1k*,T2I|	A	AAq	Aqww-'aQKK 	g:7eTUKK 177mG#9!!GBK	GAMM'6w%PQa4Kr7   sign	logabsdetc                    t        | d       t        | dd       | j                  }| j                  |d d       }| j                  |d d t	        | j
                              }t        j                  |t        |d      | j
                  | j                        }| j                  |d d t        j                        }||||fS )Nzlinalg.slogdetFr  r   r  r   )r   r   r   r   rR   rP   rH   r  r   rt   r3  )r  r   rZ  r[  r7  r  s         r5   _linalg_slogdetr]    s     a)*1.6GGE;;uSbz"DE#2Joagg.FGI			*5%8ggxx	
B [[s5;;[7FB&&r7   full_matrices
compute_uvdriverc                 b   t        | d       t        | d       t        | j                  d d       }| j                  d   }| j                  d   }t	        ||      }|r|||r|n|gz   }| j                  |      }	|	j                  |t        |d             ||r|n||gz   }
| j                  |
      }t        |       dk(  }|j                  |
t        |
|             n$| j                  dg      }	| j                  dg      }| j                  ||gz   t        | j                              }|	||fS )	Nz
linalg.svdr  r   Fr  r   r   r   )r  r   r   r   r   r   r   r   r   rR   rP   )r  r^  r_  r`  r   rT  r   rS  U_shaper  V_shapeVis_cudaSs                 r5   _linalg_svd_metarg    s#    !\"1l+aggcrl#J	A	AAq	A11==KK 	g:7eTU]1==KK 
 a.F*	g:7gVW KKKK 	
J!$OAGG,DEAa7Nr7   arg1arg2c                    | j                   d d }|j                   d d }t        ||      }t        |      }|| j                  d      | j                  d      gz  }t        |      }||j                  d      |j                  d      gz  }||fS )Nr  r   )r   r$   r   r   )rh  ri  arg1_batch_sizesarg2_batch_sizesexpand_batch_portionarg1_expand_sizearg2_expand_sizes          r5   r-  r-    s    
 zz#2zz#2,-=?OP012		"66012		"66---r7   c                     |rt        | ||       t        | |      \  }}|| j                  k(  r| n| j                  |      }||j                  k(  r|n|j                  |      }||fS r2   )r  r-  r   expand)rh  ri  r  rn  ro  arg1_broadcastedarg2_broadcasteds          r5   r  r    sv     tT40)EdD)Q&& !DJJ.DKK@P4Q  !DJJ.DKK@P4Q  ---r7   r   c                     | j                   d d }|j                  dk(  xs- | j                  dz
  |j                  k(  xr |j                   |k(  }|S )Nr   r   )r   r   )r   r   expected_batched_rhs_shapevector_cases       r5   linalg_solve_is_vector_rhsrw    sS    !&Sb!1**/ 

Q%**$R8R)R  r7   )r   r  r  r7  r  r  c                    t         d       t        j                   j                  j                  k(   fd       t	               }|rj                  d      n}	t         |	|d       t        |	       \  }
}t        j                  |xs | d        |r|
d d n|
}t        j                  |t        ||       j                  j                        } j                  }t        j                  |t        |d       j                   j                        } j                  |d d t        j                        } j                  |d d t        j                        }||||f}||||f}t        d	 |D              rbt        ||      D ]S  \  }}t!        ||j                         |j#                  |j                  |j%                                t'        ||d
       U |S )Nzlinalg.solvec                  >    d j                    dj                    dS )NzKlinalg.solve: Expected A and B to have the same dtype, but found A of type r>  r&  r   )r  r  s   r5   rW   z"_linalg_solve_ex.<locals>.<lambda>  s%    Ywwiqwwix9 r7   r   c                       	 y)Nzlinalg.solve: Vector broadcasting of the left hand side is not supported for left=False. In this case linalg.solve is equivalent to B / A.squeeze(-1)r;   r;   r7   r5   rW   z"_linalg_solve_ex.<locals>.<lambda>  s    K r7   r  Fr   r  c              3   $   K   | ]  }|d u 
 y wr2   r;   rd   rE   s     r5   rf   z#_linalg_solve_ex.<locals>.<genexpr>3  s     
&Q1D=
&s   )	copy_fromcopy_toexact_dtype)r   rH   rY   rP   rw  	unsqueezer  r-  r  r   rt   r   r   r3  allzipr    r   r   r"   )r  r  r   r  r  r7  r  r  rv  B_B_broad_shaperC   result_shaperesult_r   LU_pivots_info_r   resrN  os   ``                    r5   _linalg_solve_exr    s    1n-	LL	177	
 -Q2K'RQBaT>23B:M1	LLK	
 *5="%-L!!*<TBggxx	G GGE


*5%8ggxx	C kk%*EKKk8GKKcr
%++K6E2vt
$CC%
(C

&#
&&SM 	FDAqa)MM!''188:.QuE	F Jr7   )r   unitriangularr   r  r   c                   || j                  dg      }t        |t              sJ t        | ||d       t	        || d       \  }}|j                  dd      j                         xr |j                         }|rt        ||j                        }|S t        ||j                        r=|j                  |j                  dd      j                         |j                  dd       |S )Nr   zlinalg.solve_triangularr  r   )r   rb   r   r  r  r  rP  is_conjr    r   r!   r   
transpose_)	r  r  r  r   r  r   r  A_avoid_copy_As	            r5   linalg_solve_triangular_metar  =  s     {kk1#c:&&&aD";<.q!T:FB<<B'557HBJJLLRXX. J  RXX.KKR,223NN2r"Jr7   XM)r  r  c                     t        j                   j                  dk\   fd       t        j                  j                  dk\  fd       t         d       j                  t         j
                  k(  rt               \  }}t        j                  |t        |d       j                   j                        }t        j                  |t        |d      j                  j                        }||fS j                  t         j                  k(  sj                  t         j                  k(  r+t        j                         } j                  dg      }||fS t        j                  dd	        fS )
Nr   c                  $    d j                    dS )NzMtorch.triangular_solve: Expected b to have at least 2 dimensions, but it has r  r  r   s   r5   rW   z'triangular_solve_meta.<locals>.<lambda>b  s    ))$79 r7   c                  $    d j                    dS )NzMtorch.triangular_solve: Expected A to have at least 2 dimensions, but it has r  r  r  s   r5   rW   z'triangular_solve_meta.<locals>.<lambda>i  r$  r7   triangular_solveFr  r  r   c                       y)Nz+triangular_solve: Got an unexpected layout.r;   r;   r7   r5   rW   z'triangular_solve_meta.<locals>.<lambda>  r_   r7   )rH   rY   r   r  rs   stridedr-  r  r   rP   rt   
sparse_csr
sparse_bsrr   r   )	r   r  r  r  r  self_broadcast_sizeA_broadcast_sizesolutioncloned_coefficients	   ``       r5   triangular_solve_metar  W  sL    
LL		Q	
 
LL	!	
 4$67xx5== 0LTST0U--&&$./BeT**;;	
 #00!./?5Q''88	
 ''' 
U%%	%U5E5E)E##D)!^^QC0 ''' 	UQR'''r7   c                 l   t        | d       t        | d       | j                  | j                  d d       }| j                  | j                        }|j	                  | j                  t        | j                  d             | j                  | j                  d d t        j                        }|||fS )Nz
linalg.detr  Fr  r   r   r  )r  detr7  r  s       r5   _linalg_det_metar    s    a&1l+
++aggcrl
#C	
QWW	BNN17775QR[["U[[[9FF?r7   c                 0    t        j                   j                  dk\  d        t        j                  j                  dk\  d        |rdndt        j                  j                     j                  d   k\  fd       t        j                  j                      j                  d   k(  fd       t        j                  j                  d    j                  d   k  d        t        j                   j                  j                  z
  d	k(   fd
       t        j                   j                  j                  k(   fd        j                  dkD  re j                  d d }j                  d d t        j                  |k(  fd       j                  d d t        j                  |k(  fd       t        j                  j                   j                  k(   fd       t        j                  j                   j                  k(   fd       t        d d       t        d d       t        j                  j                  t        j                  d      j                  j                        S )Nr   c                       y)Nz3torch.ormqr: input must have at least 2 dimensions.r;   r;   r7   r5   rW   zormqr.<locals>.<lambda>  r_   r7   c                       y)Nz3torch.ormqr: other must have at least 2 dimensions.r;   r;   r7   r5   rW   zormqr.<locals>.<lambda>  r_   r7   r  r   c                      d  dS )Ntorch.ormqr: other.shape[z0] must be greater than or equal to tau.shape[-1]r;   left_size_conditions   r5   rW   zormqr.<locals>.<lambda>  s    +,?+@@pq r7   c                      d  dS )Nr  z"] must be equal to input.shape[-2]r;   r  s   r5   rW   zormqr.<locals>.<lambda>  s    +,?+@@bc r7   c                       y)NzHtorch.ormqr: tau.shape[-1] must be less than or equal to input.shape[-1]r;   r;   r7   r5   rW   zormqr.<locals>.<lambda>  r_   r7   r   c                  <    dj                    d j                    S )Nz[torch.ormqr: Expected tau to have one dimension less than input, but got tau.ndim equal to r  r  r  s   r5   rW   zormqr.<locals>.<lambda>  r  r7   c                  <    dj                    d j                    S )Nzhtorch.ormqr: Expected other to have the same number of dimensions as input, but got other.ndim equal to r  r  r   r   s   r5   rW   zormqr.<locals>.<lambda>  s+    ++0::,6RSXS]S]R^` r7   c                      d  S )NzWtorch.ormqr: Expected batch dimensions of tau to be equal to input.shape[:-2], but got r;   r  s   r5   rW   zormqr.<locals>.<lambda>  r  r7   c                      d  S )NzYtorch.ormqr: Expected batch dimensions of other to be equal to input.shape[:-2], but got r;   )actual_batch_other_shapes   r5   rW   zormqr.<locals>.<lambda>  s    66N5OQ r7   c                  <    d j                    dj                    S )NzPtorch.ormqr: Expected input and tau to have the same dtype, but input has dtype z and tau has dtype r   r  s   r5   rW   zormqr.<locals>.<lambda>  s'    ##(;;-/B399+O r7   c                  <    d j                    dj                    S )NzRtorch.ormqr: Expected input and other to have the same dtype, but input has dtype z and other has dtype r   r  s   r5   rW   zormqr.<locals>.<lambda>  s'    ##(;;-/DU[[MS r7   ztorch.ormqrr  r   Fr  r  )	rH   rY   r   r   rP   r  r  r   rt   )	r   r  r   r   r  expected_batch_shaper  r  r  s	   ```   @@@r5   ormqrr    s    
LL

aV 
LL

aV !%""	LL'(CIIbM9q 
LL'(EKKO;c
 
LL		"R(Z
 
LL

SXX"	
 
LL

ejj 	
 zzA~${{3B/!$3B"&::	
 $);;s#3 $(<<	
 
LL		U[[ 	
 
LLu{{"	
 M3u5M5%9[[*5;;%Hkk||	 r7   c                t    t        j                  t              dz  k(  fd        j                  }|dz   k(  }|}| }|r*t	        d|      D ]  }|xr  j                  |      dk7  } n)t	        d|      D ]  }|xr  j                  |      dk7  } t        j                  |xs | fd       y )Nr   c                  ,    dd z   dt               S )Nzpadding size is expected to be r   z, but got: r   )rw   paddings   r5   rW   z,_padding_check_valid_input.<locals>.<lambda>  s    1!c'+c'l^T r7   r   r   c                  :    d dz    d dz    dj                    S )Nz	Expected r   zD or r   zcD (batch mode) tensor with possibly 0 batch size and other non-zero dimensions for input, but got: r'  )rw   r   s   r5   rW   z,_padding_check_valid_input.<locals>.<lambda>
  s2    aycAgY /AAFO r7   )rH   rY   r   r   r   r   )r   r  rw   	input_dimis_batch_modevalid_batch_modevalid_non_batch_moder   s   ```     r5   _padding_check_valid_inputr    s    	LLGCT
 

I#'*M$,,q)$ 	GA/FEJJqMQ4F	G q)$ 	OA#7#NEJJqMQ<N 	O 
LL00	
r7   c                   	
 d}dd} j                   dk(  r j                  d      }dz  |dz  }t         |d       |\  	
 j                  |      } j                        	z   
z   |r&t        j                  	k  xr 
k   	
fd       t        j                  dk\  fd        j                   dk(  r j                  |f      S  j                  ||f      S )Nr   r   r-   r   c                  4    d d d  dj                    S NzcArgument #4: Padding size should be less than the corresponding input dimension, but got: padding (rj   ) at dimension 
 of input r'  dim_wr   pad_lpad_rs   r5   rW   z_pad1d_common.<locals>.<lambda>&  2    %%*G2eWOE7*UZU`U`Tac r7   c                      d  d S )Nz
input (W: z%) is too small. Calculated output W: r;   )input_woutput_ws   r5   rW   z_pad1d_common.<locals>.<lambda>.  s    *WI%J8*U r7   r   )r   r   r  rH   rY   r   )r   r  is_reflection	dim_planenbatchnplaner  r  r  r  r  s   `     @@@@@r5   _pad1d_commonr    s    IEFzzQA
Q	ug15LE5ZZ	"FjjG&HGO/	
 
LLAU
 zzQ1229::r7   c                     t        | |d      S NTr  r  r   r  s     r5   meta_reflection_pad1dr  7       t<<r7   c                     t        | |d      S NFr  r  r  s     r5   meta_replication_pad1dr  =       u==r7   c                    d|s#t        j                  t        |      dk(  d        j                  dk(  rdz  |\  j	                        }|z   z   |r&t        j                  |k  xr |k  fd       t        j                   j	                        k(   fd       j                  j                        S )Nr   r   c                       y)Nz padding size is expected to be 2r;   r;   r7   r5   rW   z(_pad1d_backward_common.<locals>.<lambda>F  r_   r7   r-   c                  4    d d d  dj                    S r  r'  r  s   r5   rW   z(_pad1d_backward_common.<locals>.<lambda>S  r  r7   c                  2    d dj                          S Nz(grad_output width unexpected. Expected: , Got: r   r  grad_outputr  s   r5   rW   z(_pad1d_backward_common.<locals>.<lambda>[  "    :8*GKL\L\]bLcKde r7   rH   rY   r   r   r   r   r   )	r  r   r  r  r  r  r  r  r  s	   ``   @@@@r5   _pad1d_backward_commonr  C  s    ES\Q&(RSzzQ
LE5jjG&HGO/	
 
LLK$$U++e
 ??5;;''r7   
grad_inputc                      t        | ||d      S r  r  r  r   r  s      r5   meta_reflection_pad1d_backwardr  a  s     "+ugTRRr7   c                      t        | ||d      S r  r  r  s      r5   meta_replication_pad1d_backwardr  g  s     "+ugUSSr7   c                   	
 ddd}d}t         |d        j                  }|dk(  r  j                  d      }dz  dz  |dz  }|\   j                  |      } j                        	 j                        
	z   z   
z   z   |rLt        j                  
k  xr 
k   fd       t        j                  	k  xr 	k   fd       t        j                  dk\  xs dk\  	
fd        j                  d	k(  r j                  |f      S  j                  ||f      S )
Nr   r   r   r      c                  4    d d d  dj                    S r  r'  r  s   r5   rW   z_pad2d_common.<locals>.<lambda>  r  r7   c                  4    d d d  dj                    S NzcArgument #6: Padding size should be less than the corresponding input dimension, but got: padding (rj   r  r  r'  dim_hr   pad_bpad_ts   r5   rW   z_pad2d_common.<locals>.<lambda>  r  r7   c                       d  d d d S )Nz
input (H:  W: z%) is too small. Calculated output H: r;   )input_hr  output_hr  s   r5   rW   z_pad2d_common.<locals>.<lambda>  s*    	gY /$$,:T(= r7   r-   r  r   r   rH   rY   r   )r   r  r  
dim_slicesr  r   r  r  r  r  r  r  r  r  r  r  r  s   `      @@@@@@@@@@r5   _pad2d_commonr  m  sU   EEJFug15::DqyA

a
!(E5%ZZ
#FjjGjjG&H&HGO/	
 	GO/	
 
LLA&Q	
 zzQ(;<<(CDDr7   c                     t        | |d      S r  r  r  s     r5   meta_reflection_pad2dr    r  r7   c                     t        | |d      S r  r  r  s     r5   meta_replication_pad2dr    r  r7   c                     ddd}|j                   }|j                         dk(  rdz  dz  |dz  }|\  }}}}|   }	|   }
|	|z   |z   |
|z   |z   t        j                   j	                        k(   fd       t        j                   j	                        k(   fd       |j                  |j                         S )Nr   r   r   r  c                  2    d dj                          S r  r   r  s   r5   rW   z%meta_pad2d_backward.<locals>.<lambda>  r  r7   c                  2    d dj                          S Nz)grad_output height unexpected. Expected: r  r   r  r  r  s   r5   rW   z%meta_pad2d_backward.<locals>.<lambda>  "    ;H:W[M]M]^cMdLef r7   )r   rw   rH   rY   r   r   )r  r   r  r  rV   r  r  r  r  r  r  r  r  r  r  s   `          @@@@r5   meta_pad2d_backwardr    s     EEIJxxzQ

Q	!(E5%GG&H&H	LLK$$U++e 
LLK$$U++f >>$**%%r7   c          	      $   	
 d	ddd}t         |d        j                  dk(  }|r% j                  d      }	dz  	dz  dz  |dz  }|\   j                  |      } j                        
 j                         j                  	      
z   z   z   z   z   z   |rrt        j                  k  xr k  	 fd       t        j                  k  xr k   fd       t        j                  
k  xr 
k   fd	       t        j                  dk\  xs dk\  xs dk\  
fd
       |r j                  |f      S  j                  |f      S )Nr-   r   r   r   r      c                  4    d d d  dj                    S r  r'  r  s   r5   rW   z_pad3d_common.<locals>.<lambda>  r  r7   c                  4    d d d  dj                    S r  r'  r  s   r5   rW   z_pad3d_common.<locals>.<lambda>  r  r7   c                  4    d d d  dj                    S )NzcArgument #8: Padding size should be less than the corresponding input dimension, but got: padding (rj   r  r  r'  )dim_dr   pad_bkpad_fs   r5   rW   z_pad3d_common.<locals>.<lambda>  s2    %%*G2fX_UG:V[VaVaUbd r7   c                  ,    d  d d d d d S )Nz
input (D:  H: r  z%) is too small. Calculated output D: r;   )input_dr  r  output_dr  r  s   r5   rW   z_pad3d_common.<locals>.<lambda>  s7    	gYd7) <$$,:T(4zK r7   r  )r   r  r  r  
batch_moder  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  s   `      @@@@@@@@@@@@@@@r5   _pad3d_commonr    s   EEEIug15qJA


Q	07-E5%vZZ	"FjjGjjGjjG'H&H&HGO/	
 	GO/	
 	GO0 0	
 
LLA7Q7(a-	
 	
 (HMNN(HEFFr7   c                     t        | |d      S r  r  r  s     r5   meta_reflection_pad3dr    r  r7   c                     t        | |d      S r  r  r  s     r5   meta_replication_pad3dr     r  r7   c                     t        j                  t        |      dk(  d        |j                  dkD  sJ  j                  |j                  k(  sJ ddd|j                  dk(  rdz  dz  dz  |\  }}}}}}|j	                        }	|j	                        }
|j	                        }|	|z   |z   |
|z   |z   ||z   |z   t        j                   j	                        k(   fd       t        j                   j	                        k(   fd       t        j                   j	                        k(   fd	       |j                  |j                        S )
N   c                       y)Nz padding size is expected to be 6r;   r;   r7   r5   rW   z%meta_pad3d_backward.<locals>.<lambda>)  r_   r7   r-   r   r   r  c                  2    d dj                          S r  r   r  s   r5   rW   z%meta_pad3d_backward.<locals>.<lambda>A  r  r7   c                  2    d dj                          S r
  r   r  s   r5   rW   z%meta_pad3d_backward.<locals>.<lambda>E  r  r7   c                  2    d dj                          S )Nz(grad_output depth unexpected. Expected: r  r   )r  r  r  s   r5   rW   z%meta_pad3d_backward.<locals>.<lambda>I  r  r7   r  )r  r   r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  s   `           @@@@@@r5   meta_pad3d_backwardr'    s_    
LLW"$NO::>>uzz)))EEEzzQ


07-E5%vjjGjjGjjG'H&H&H	LLK$$U++e 
LLK$$U++f 
LLK$$U++e
 ??5;;''r7   pc                 J   t        j                  | j                         d        | j                  d      }|dk  r0| j	                  dg      j                  t         j                        S | j	                  ||dz
  z  dz  f      j                  t         j                        S )Nc                       y)Nz(_pdist_forward requires contiguous inputr;   r;   r7   r5   rW   z%meta__pdist_forward.<locals>.<lambda>S  r_   r7   r   r   r   r   )rH   rY   rP  r   r   r  r  )r   r(  r   s      r5   meta__pdist_forwardr+  O  s     
LLP 			!AAv~~qc"%%E4R4R%SS~~qAE{a/125588 6 
 	
r7   gradpdistc                     t        j                  |j                         d        t        j                  |j                         d        t        j                  |t         j                        S )Nc                       y)Nz._pdist_backward requires self to be contiguousr;   r;   r7   r5   rW   z&meta__pdist_backward.<locals>.<lambda>b  r_   r7   c                       y)Nz/_pdist_backward requires pdist to be contiguousr;   r;   r7   r5   rW   z&meta__pdist_backward.<locals>.<lambda>e  r_   r7   r   )rH   rY   rP  r   r  )r,  r   r(  r-  s       r5   meta__pdist_backwardr1  ^  sW     
LLV 
LLX D0N0NOOr7   )r?  r>  c                   	
 j                  d      }j                  d      }j                  d      } j                  |||f       t        j                  j	                         dk(  d        t        j                  j	                         dk(  d        t        j                   j
                  j
                  cxk(  xr j
                  k(  nc  fd       j                  }j                  	|d   
|d   t        j                  	d   
k(  xr 	d   k(  	
fd        j                   j                               S )	Nr   r   r   r-   c                       yNzbatch1 must be a 3D tensorr;   r;   r7   r5   rW   zmeta_baddbmm.<locals>.<lambda>q  r_   r7   c                       yNzbatch2 must be a 3D tensorr;   r;   r7   r5   rW   zmeta_baddbmm.<locals>.<lambda>r  r_   r7   c                  V    dj                    d j                    dj                    S )Nz+Input dtypes must be the same, got: input: z
, batch1: z
, batch2: r   )batch1batch2r   s   r5   rW   zmeta_baddbmm.<locals>.<lambda>u  s0    =djj\TZT`T`Saaklrlxlxkyz r7   c            	      .    d d d d    d d    d	S Nz@Expected size for first two dimensions of batch2 tensor to be: [rj   z] but got: [r   r   ].r;   batch2_sizesbscontraction_sizes   r5   rW   zmeta_baddbmm.<locals>.<lambda>}  s:    t2&'|LO3DB|TUFWWY[ r7   )r   rq  rH   rY   rw   rP   r   r   )r   r8  r9  r?  r>  dim1dim2dim3batch1_sizesr>  r?  r@  s   ```      @@@r5   meta_baddbmmrE  j  s    ;;q>D;;q>D;;q>D;;dD)*D	LL"$HI	LL"$HI	LL

fll2fll2z <<L<<L	aB#A	LLQ2E,q/5E"E	
 >>$))+&&r7   c                L    t        j                  | t         j                        S r   r[  r   r   s     r5   meta_bernoullirH    s     D0G0GHHr7   c                     | S r2   r;   r   r(  r   s      r5   meta_bernoulli_rK        Kr7   c                 L    t        j                  | t         j                        S r   r[  rJ  s      r5   meta_bernoulli_prN    s     D0G0GHHr7   c                 ,    t        j                  |       S r2   rH   r   rG  s     r5   meta_poissonrQ         D!!r7   c                     t        j                  |
| j                         k  d        t        j                  | t         j                        }t        j                  |       |fS )Nc                       y)NzJError in fused_moving_avg_obs_fake_quant_cpu: ch_axis must be < self.dim()r;   r;   r7   r5   rW   z6meta__fused_moving_avg_obs_fq_helper.<locals>.<lambda>  r_   r7   r   )rH   rY   rw   r   bool)r   observer_onfake_quant_onrunning_minrunning_maxscale
zero_pointaveraging_const	quant_min	quant_maxch_axisper_row_fake_quantsymmetric_quantmasks                 r5   $meta__fused_moving_avg_obs_fq_helperrc    sO      
LL$((*\ D

3DT"D))r7   c                 H   t        j                  | j                         dk(  d        t        j                  |j                         dk(  d        | j                  \  |j                  \  t        j                  k(  fd       | j	                        S )Nr   c                       y)Nza must be 2Dr;   r;   r7   r5   rW   zmeta_mm.<locals>.<lambda>  r_   r7   c                       y)Nzb must be 2Dr;   r;   r7   r5   rW   zmeta_mm.<locals>.<lambda>  r_   r7   c            	      "    d d  d d d	S )Nz/a and b must have same reduction dim, but got [rj   z] X [r<  r;   )M1M2Nr1  s   r5   rW   zmeta_mm.<locals>.<lambda>  s(    A!Brd%PRtSUVWUXXZ[ r7   )rH   rY   rw   r   r   )abrh  ri  rj  r1  s     @@@@r5   meta_mmrm    sz     
LLA56	LLA56GGEArGGEB	LL
b[ ;;q!r7   c                      |r(t         fdt         j                        D              S t        j                   j
                        S )Nc              3   H   K   | ]  }|vrj                   |   nd   yw)r   Nr'  )rd   r   dimsr   s     r5   rf   z+_compute_reduction_shape.<locals>.<genexpr>  s$     UqatmTZZ]:Us   ")rX   r   r   r@   compute_reduction_output_shaper   )r   rp  r|  s   `` r5   rz  rz    s7    UE$))DTUUU//

DAAr7   c                    t        | t        j                  j                        r| j                  j
                  S t        | d      rEt        | j                  d      r/| j                  j
                  dk7  r| j                  j
                  S y)Nrt   rl   rr   r   )rb   rH   _subclasses
FakeTensorfake_devicerl   hasattrrt   )r"  s    r5   r   r     sg    &%++667!!&&&!FMM6*MM&(}}!!!r7   input_tensorr   r  dilationis_transposedgroupsoutput_paddingc                    dt         dt         dt         dt         dt         dt         fd}dt         dt         dt         dt         dt         dt         dt         fd	}	|j                  d
d  }
| j                  d
d  |r||j                  d   z  }n<|j                  d   }|j                  d   |z  | j                  d   k7  rt        d      | j                  d   |gt        |t              r|gt              z  }n t        |      dk(  r|d   gt              z  }t        |t              r|gt              z  }n t        |      dk(  r|d   gt              z  }t        |t              r|gt              z  }n t        |      dk(  r|d   gt              z  }d }|rCt        |t              r|gt              z  }n#t        |      dk(  r|d   gt              z  }n|}t        t                    D ]]  }|r/j                   |	|   ||   ||   |
|   ||   ||                4j                   ||   ||   ||   |
|   ||                _ t        j                  t        d d
d  D              fd       S )Nlnr(  r   rS  r  r.   c                 6    | d|z  z   ||dz
  z  z
  dz
  |z  dz   S )a  
        Formula to apply to calculate the length of some dimension of the output

        See: https://pytorch.org/docs/stable/generated/torch.nn.Conv2d.html

        Args:
            ln: length of the dimension
            p: padding in that dim
            d: dilation in that dim
            k: kernel size in that dim
            s: stride in that dim
        Returns:
            The output length
        r   r   r;   )r}  r(  r   rS  r  s        r5   _formulaz+calc_conv_nd_return_shape.<locals>._formula  s.     QU
Q!a%[(1,2Q66r7   r3   c                 <    | dz
  |z  d|z  z
  ||dz
  z  z   |z   dz   S )a  
        Formula to apply to calculate the length of some dimension of the output
        if transposed convolution is used.
        See: https://pytorch.org/docs/stable/generated/torch.nn.ConvTranspose2d.html

        Args:
            ln: length of the dimension
            p: padding in that dim
            d: dilation in that dim
            k: kernel size in that dim
            s: stride in that dim
            op: output padding in that dim

        Returns:
            The output length
        r   r   r;   )r}  r(  r   rS  r  r3   s         r5   _formula_transposedz6calc_conv_nd_return_shape.<locals>._formula_transposed  s2    " Q!|a!e#a1q5k1B6::r7   r   r   r   zInvalid channel dimensionsc              3   &   K   | ]	  }|d kD    ywr   Nr;   r|  s     r5   rf   z,calc_conv_nd_return_shape.<locals>.<genexpr>=	  s     )aAE)   c                  .    dt                ddd   dS )NzGiven input size per channel: z&. Calculated output size per channel: r   z. Output size is too small)r   )rp  	ret_shapes   r5   rW   z+calc_conv_nd_return_shape.<locals>.<lambda>>	  s*    0d =//8}o >#$ r7   )r   r   r  rb   r   r   r   r   rH   rY   rx   )rw  r+  r   r  rx  ry  rz  r{  r  r  kernel_sizeout_channelsoutput_padding_listr   rp  r  s                 @@r5   calc_conv_nd_return_shaper    s   7S 7S 7S 7S 7S 7S 7"; ; ; ; ; ; ;QT ;& ,,qr"Kab!DQ/||A<<?V#|'9'9!'<<;<<##A&5I&'"CI%	V	)s4y('7#)c$i'	W	1:,T*(G$:D	)	X!	QK=3t9,/3ng.#1"2SY"> A%#1!#4"5D	"A"03t9 #GAJQKN1I'*	 a'!*hqk;q>6RS9U" 
LL)9QR=))	$ r7   c                 b    t         j                  j                  |       t         j                  k(  S r2   rH   _prims_commonsuggest_memory_formatchannels_lasttens    r5   is_channels_lastr  F	  s$    44S9U=P=PPPr7   running_meanrunning_vartrainingexponential_average_factorepsilonc                 r     j                   }||j                   n|j                   }	||j                   n|j                   }
 fd} j                  |      j                   |             }|r# j                  |	      } j                  |
      }n" j                  d      } j                  d      }|||fS )Nc                      t               rt        j                  S  j                  t        j                        rt        j                  S t        j                  S r   )r  rH   r  rP  r   )rw  s   r5   pick_memory_formatz2meta_miopen_batch_norm.<locals>.pick_memory_format]	  sF    L)&&&%%E4K4K%L***&&&r7   r   r   )r   r   r  )rw  r+  r-  r  r  r  r  r  r   save_mean_shapesave_var_shaper  r   	save_meansave_vars   `              r5   meta_miopen_batch_normr  J	  s     ""I -9,Dl((&,,O*5*A[&&v||N' 
 
 
+
.
.=O=Q
.
RC **?;	)).9 **40	))$/	8##r7   c	           
            fd}	t         ||||||r|nd       }
d}d} j                  |      dk(  rd|
|<    j                  |
      }|j                   |	             }|S )Nc                  d   t               dk(  r&t               st              r+t        j                  S t               rt        j                  S  j	                  t        j
                        rt        j
                  S  j	                  t        j                        rt        j                  S y Nr   r   )r   r  rH   r  rP  r   preserve_format)rw  r+  s   r5   r  z%meta_conv.<locals>.pick_memory_format|	  s    |$.-1A&1I***-***%%E4K4K%L***''e6K6K'L((( Mr7   r   r   r   )r  r   r   r  )rw  r+  r-  r   r  rx  ry  r{  rz  r  	shape_outinput_channels_dimoutput_channels_dimr   s   ``            r5   	meta_convr  p	  s    
) *'T	I +,1)*	%&

 
 
+C
&&13&
4CJr7   mkldnnc
           
          t        | ||||d|g       }
| j                  |
      }t        j                  }| j	                         dk(  rt        j
                  }|j                  |      }|S )NFr  r   )r  r   rH   r  rw   channels_last_3dr  )rw  r+  r-  r  r   rx  rz  attrscalars	algorithmr  r   out_memory_formats                r5   meta_mkldnn_convolution_defaultr  	  sp     .&&'8UFB
	 $$Y/!//" % 6 6ff#4f5
r7   c                 b    | j                  g | j                  d d |j                  d         S Nr   r   r   r   )rw  r+  r-  r  r  r  s         r5   meta_linear_pointwise_defaultr  	  s5     %%&Q(:(:3B(?&Qa&QRRr7   mklc                 b    | j                  g | j                  d d |j                  d         S r  r  )rw  packed_weightorig_weightr-  r   s        r5   meta_mkl_linearr  	  s:    ))@,$$Sb)@;+<+<Q+?@ r7   onednnc           
         t        | ||||	d|
d       }|t        j                  t        j                  t        j                  t        j
                  fv sJ | j                  ||      }|j                  t        j                        }|S )NFr   r   )	r  rH   rL  rN  uint8r2  r   r  r  )rE   x_scalex_zpww_scalew_zpr-  r   r  rx  rz  output_scaleoutput_zero_pointoutput_dtyper  r  r  r  r   s                      r5   meta_qconv2d_pointwiser  	  s|    ( .	
	 u~~u{{EJJWWWWkk)<k8ff5#6#6f7
r7   c                     |dk(  sJ |S )Nsumr;   )rE   r  r  r  r  r  accumr-  r   r  rx  rz  r  r  r  accum_scaleaccum_zero_pointbinary_op_namer>  unary_op_nameunary_op_argsunary_op_algorithms                         r5   meta_qconv2d_pointwise_binaryr  	  s    2 &&&r7   c                     t        | j                        }|j                  d   |d<   |	t        j                  t        j                  t        j
                  t        j                  fv sJ | j                  ||	      }|S )Nr   r   r   )r   r   rH   rL  rN  r2  r  r   )rE   r  r  r  r  r  r-  r  r  r  post_op_namepost_op_argspost_op_algorithmrU  r   s                  r5   meta_qlinear_pointwiser  
  sa    " AGG}771:Ru~~uzz5;;WWWWkk,lk;
r7   c                    |dk(  r|S t        | j                        }|j                  d   |d<   |
t        j                  t        j                  t        j
                  t        j                  fv sJ | j                  ||
      }|S )Nr  r   r   r   )r   r   rH   rL  rN  r  r2  r   )rE   r  r  r  r  r  x_2r-  r  r  r  x2_scalex2_zpr  r>  r  r  r  rU  r   s                       r5   meta_qlinear_pointwise_binaryr  %
  sn    , U"JAGG}771:Ru~~u{{EJJWWWWkk,lk;
r7   c                 v    t        | j                        }|j                  d   |d<   | j                  |      }|S )Nr   r   )r   r   r   )rE   r  r-  rU  r   s        r5   meta_linear_dynamic_fp16r  D
  s6     AGG}771:Rkk,'
r7   	quantizedr   c                 .   t        | |||||      \  }}}| j                         dk(  r| j                  d      nd}	t        j                  }
| j                         dk(  r|||g}n|	|||g}t        j
                  || j                  | j                  |
      S Nr  r   r-   rl  )#max_pool2d_checks_and_compute_shaperw   r   rH   r  r|   rP   rt   r   r  r   r  rx  	ceil_modenInputPlaneoutputHeightoutputWidthr  r   r   s               r5   meta_quantized_max_pool2dr  U
  s     0;9
		
 $)99;!#3B++99;!{;DK{CD{{++<<'	
 	
r7   c                    t        j                  | j                         dk(  d| j                          d       t        j                  |j                         dk(  d|j                          d       t        j                  | j                  t         j                  t         j
                  t         j                  fv d| j                          t        j                  |j                  t         j                  k(  d|j                          t        j                  |j                  t         j                  k(  d|j                          t        j                  |j                  | j                  k(  d|j                          | j                  | j                  d	      |j                  d	      | j                  
      S )Nr   zx must be a 2D tensor, got Dzw must be a 2D tensor, got #expected x to be f32/f16/bf16, got expected w to be uint8, got z q_group_size must be int64, got z5q_scale_and_zeros must have the same dtype as x, got r   r   )rH   rY   rw   rP   rL  rM  rN  r  r   r   r   rE   r  q_group_sizeq_scale_and_zeross       r5   meta_int4mm_packed_weight_cpur  r
  s@   QUUW\%@	#KLQUUW\%@	#KLGGu}}ennEE1!'';	
 	QWW+/KAGG9-UV%++-.|/A/A.BC	
 	##qww.CDUD[D[C\]	
 {{166!9affQiqww{??r7   c                      t        j                   j                         k(  xr  j                     k(   fd       y )Nc                  j    d  d d ddj                          d dj                      z   S )NzExpected a tensor of dimension z and tensor.size[z] == rj   zbut got : dimension z] = rw   r   )rw   dim_sizer   r"  s   r5   rW   z check_dim_size.<locals>.<lambda>
  sP    1#6GzQVW[V\\^_ .?zfll[cNdMe
fg r7   )rH   rY   rw   r   )r"  rw   r  r   s   ````r5   check_dim_sizer  
  s6    	LL

>X 6$ >	gr7   c                     d } |d|      \  }}	t        j                  t        |      dv d        t        j                   j                  t         j                  t         j
                  t         j                  t         j                  fv fd       t        |      dk(  r||	}}
n%t        |      dk(  r|d   |d   }}
n |d|      \  }
} |d	|      \  }}t        j                  |d u xs |dk7  d
         j                         dk(  r j                  d      nd} j                  d      } j                  d      } j                  d      }t        ||||
d|      }t        ||	||d|      }t        j                         }t         ||	|
|||dd||||||        j                         dk(  r|||g}n||||g}t        j                  | j                   j                  |      S )Nc                      t        j                  t        |      dv  fd       |d   }t        |      dk(  r|n|d   }||fS )Nr   r   c                      d  dS )Nzavg_pool2d: 4 must either be a single int, or a tuple of two intsr;   r  s   r5   rW   z1meta_avg_pool2d.<locals>.unpack.<locals>.<lambda>
      l4&(\] r7   r   r   rH   rY   r   r  r  HWs   `   r5   unpackzmeta_avg_pool2d.<locals>.unpack
  G    H]	
 FSQACF!tr7   r  r   r   r   c                       yNzOavg_pool2d: stride must either be omitted, a single int, or a tuple of two intsr;   r;   r7   r5   rW   z!meta_avg_pool2d.<locals>.<lambda>
  r_   r7   c                  @    d j                   j                          dS )Nz""avg_pool2d" not implemented for ''rP   __str__r   s   r5   rW   z!meta_avg_pool2d.<locals>.<lambda>
      6u{{7J7J7L6MQQ r7   r   r   r   r  c                       yNzdivisor must be not zeror;   r;   r7   r5   rW   z!meta_avg_pool2d.<locals>.<lambda>
  r_   r7   r  r  r  r   r-   rl  )rH   rY   r   rP   r  uint16uint32uint64rw   r   pooling_output_shaper@   r  pool2d_shape_checkr|   rt   )r   r  r   r  r  count_include_paddivisor_overrider   kHkWdHdWpadHpadWr  r  inputHeight
inputWidthr  r  r   r   s   `                     r5   meta_avg_pool2dr  
  s    M;/FB	LLFy a 
LLEKKu||U\\RRQ 6{aRB	V	F1IB&)B	7+JD$	LLD 9$4$9*
  %yy{a/UZZ^QF**R.K**R.KBJ'Rr1iPL&z2tRINK//6M



		$ yy{a\;7\;?;;kk||#	 r7   c                     t        | ||||||dd|	|
||||       | j                         }|	}t        |||dz
  |       t        |||dz
  |       t        |||dz
  |       y )Nr   r-   r   )r  rw   r  )r   
gradOutputr  r  r  r  r  r  r  r  r  r  r  r  
mem_formatr   nOutputPlanes                    r5   avg_pool2d_backward_shape_checkr"  
  s    " 



		$ 99;DL:tTAX|<:tTAX|<:tTAX{;r7   c                    t        j                  t        |      dk(  xs t        |      dk(  d        |d   }t        |      dk(  r|n|d   }	t        j                  t        |      dk(  xs t        |      dk(  xs t        |      dk(  d        t        |      dk(  r|n|d   }
t        |      dk(  r|	nt        |      dk(  r|
n|d   }t        j                  t        |      dk(  xs t        |      dk(  d        |d   }t        |      dk(  r|n|d   }t        j                  |d u xs |dk7  d        |j                  }|j	                         dk(  r|d	   nd}|d
   }|d   }|d   }t        ||||
d|      }t        ||	||d|      }t        j                  |      }t        || |||	|
|||||||||       t        j                  ||j                  |j                  |      S )Nr   r   c                       y)NzKavg_pool2d: kernel_size must either be a single int, or a tuple of two intsr;   r;   r7   r5   rW   z*meta_avg_pool2d_backward.<locals>.<lambda>  r_   r7   r   c                       yr  r;   r;   r7   r5   rW   z*meta_avg_pool2d_backward.<locals>.<lambda>!  r_   r7   c                       y)NzGavg_pool2d: padding must either be a single int, or a tuple of two intsr;   r;   r7   r5   rW   z*meta_avg_pool2d_backward.<locals>.<lambda>'  r_   r7   c                       yr  r;   r;   r7   r5   rW   z*meta_avg_pool2d_backward.<locals>.<lambda>.  r_   r7   r  r  r  r  r   rl  )rH   rY   r   r   rw   r  r@   r  r"  r|   rP   rt   )gradOutput_r   r  r   r  r  r  r  r  r  r  r  r  r  
input_sizer  r  r  r  r  r  r   s                         r5   meta_avg_pool2d_backwardr*    s    
LLKA6[!1Q!6] 
QB;1$+a.B	LLFq@CK1,@Fq0@a 6{aVAYB6{a3v;!+;RB	LLG.S\Q.Y 1:Dw<1$4'!*D	LLD 9$4$9*
 J$yy{a/Z^QFR.KR.KBJ'Rr1iPL&z2tRINK,,U3J#



$ ;;kk|| 	 r7   c                     t        j                  t        |      dv d        |d   }t        |      dk(  r|n|d   }t        |      dk(  r|n|d   }	t        j                  | xs t        |      dv d        t        j                   j                  t         j                  t         j
                  t         j                  t         j                  fv fd       |s|n|d   }
|s|nt        |      dk(  r|
n|d   }|s|	nt        |      dk(  r|
n|d   }t        j                  t        |      dv d        |d   }t        |      dk(  r|n|d   }t        |      dk(  r|n|d   }t        j                   j                  d	v d
        t        j                  | xs |dk7  d         j                  d      } j                  d      } j                  d      } j                  d      } j                  d      }t        ||||
d|      }t        ||||d|      }t        ||	||d|      }t         ||||	|
|||||ddd||||||dd        j                  dk(  r j                  ||||f      S  j                  |||||f      S )Nr   r-   c                       yNzFavg_pool3d: kernel_size must be a single int, or a tuple of three intsr;   r;   r7   r5   rW   z!meta_avg_pool3d.<locals>.<lambda>c  r_   r7   r   r   r   c                       yNzJavg_pool3d: stride must be omitted, a single int, or a tuple of three intsr;   r;   r7   r5   rW   z!meta_avg_pool3d.<locals>.<lambda>k  r_   r7   c                  @    d j                   j                          dS )Nz""avg_pool3d" not implemented for 'r  r  r	  s   r5   rW   z!meta_avg_pool3d.<locals>.<lambda>o  r
  r7   c                       yNzBavg_pool3d: padding must be a single int, or a tuple of three intsr;   r;   r7   r5   rW   z!meta_avg_pool3d.<locals>.<lambda>w  r_   r7   r  r  c                       yNz9non-empty 4D or 5D (batch mode) tensor expected for inputr;   r;   r7   r5   rW   z!meta_avg_pool3d.<locals>.<lambda>  r_   r7   c                       yr  r;   r;   r7   r5   rW   z!meta_avg_pool3d.<locals>.<lambda>  r_   r7   r  r  r  r   zavg_pool3d()T)check_input_sizer  )rH   rY   r   rP   r  r  r  r  r   r   r  pool3d_shape_checkr   )r   r  r   r  r  r  r  kTr  r  dTr  r  padTr  r  r  nslicesitimeiheightiwidthotimeoheightowidths   `                       r5   meta_avg_pool3drD  V  s    
LLKF"X 
QB;1$+a.B;1$+a.B	LL
+c&kV+\ 
LLEKKu||U\\RRQ vayBc&kQ&6F1IBc&kQ&6F1IB	LLGT 1:Dw<1$4'!*Dw<1$4'!*D	LL

fK
 
LL5 0A 5*
 ZZ]FjjnGJJrNEjjnGZZ^F D"aCE"7Bb!YGG!&"dB9EF





			-2 zzQ@AAHIIr7   c                    t        j                  t        |      dv d        |d   }t        |      dk(  r|n|d   }	t        |      dk(  r|n|d   }
t        j                  | xs t        |      dv d        |s|n|d   }|s|	nt        |      dk(  r|n|d   }|s|
nt        |      dk(  r|n|d   }t        j                  t        |      dv d        |d   }t        |      dk(  r|n|d   }t        |      dk(  r|n|d   }t        j                  |j                  dv d	        t        j                  | xs |dk7  d
        |j	                  d      }|j	                  d      }|j	                  d      }|j	                  d      }t        ||||d|      }t        ||	||d|      }t        ||
||d|      }t        || |||	|
||||||||||||d       |j                  |j                        S )Nr,  c                       yr.  r;   r;   r7   r5   rW   z*meta_avg_pool3d_backward.<locals>.<lambda>  r_   r7   r   r   r   c                       yr0  r;   r;   r7   r5   rW   z*meta_avg_pool3d_backward.<locals>.<lambda>  r_   r7   c                       yr3  r;   r;   r7   r5   rW   z*meta_avg_pool3d_backward.<locals>.<lambda>  r_   r7   r4  c                       yr6  r;   r;   r7   r5   rW   z*meta_avg_pool3d_backward.<locals>.<lambda>  r_   r7   c                       yr  r;   r;   r7   r5   rW   z*meta_avg_pool3d_backward.<locals>.<lambda>  r_   r7   r  r  r  r   zavg_pool3d_backward())	rH   rY   r   r   r   r  avg_pool3d_backward_shape_checkr   r   )r  r   r  r   r  r  r  r  r:  r  r  r;  r  r  r<  r  r  r=  r>  r?  r@  otime_for_shape_checkoheight_for_shape_checkowidth_for_shape_checks                           r5   meta_avg_pool3d_backwardrO    s    
LLKF"X 
QB;1$+a.B;1$+a.B	LL
+c&kV+\ vayBc&kQ&6F1IBc&kQ&6F1IB	LLGT 1:Dw<1$4'!*Dw<1$4'!*D	LL

fK
 
LL5 0A 5*
 jjnGJJrNEjjnGZZ^F0D"aS27Bb!YW1&"dB9U#





', ??5;;''r7   c                 ,    t        j                   j                  dk(  xs  j                  dk(   fd        j                  d d t	        |      z   }t        j                         }t        j                  | j                   j                  |      S )Nr-   r  c                  "    d j                    S )Nz"Expected 3D or 4D tensor, but got r'  r   s   r5   rW   z*meta_adaptive_avg_pool2d.<locals>.<lambda>      4TZZLA r7   r  rl  )
rH   rY   r   r   rX   r@   r  r|   rP   rt   )r   output_sizerU  r   s   `   r5   meta_adaptive_avg_pool2drT     s|    	LL		Q($))q.A ::cr?U;%77L//5M ;;jj{{#	 r7   c                      t        j                   j                  dk(  xs  j                  dk(   fd        j                   j                  d d t        |      z         S )Nr  r  c                  "    d j                    S )Nz"Expected 4D or 5D tensor, but got r'  r   s   r5   rW   z*meta_adaptive_avg_pool3d.<locals>.<lambda>  rR  r7   r  )rH   rY   r   r   r   rX   )r   rS  s   ` r5   meta_adaptive_avg_pool3drW    sO    	LL		Q($))q.A >>$**Sb/E+,>>??r7   c                      j                   }t        d|      D ].  t        j                   j	                        dkD   fd       0 t        j                  |dk(  xs |dk(  fd       t        j                  j
                   j
                  k(   fd       t        j                  }t              rt        j                  }j                  j                        j                  |      S )	Nr   r   c                  *    d j                    d dS )Nz{adaptive_avg_pool2d_backward(): Expected grad_output to have non-zero                       size for non-batch dimensions,  with dimension  being emptyr'  )grad_outr   s   r5   rW   z4meta__adaptive_avg_pool2d_backward.<locals>.<lambda>!  s&     66>nn5EEUVWUXXdf r7   r-   r  c                  "    d j                    S )NzBadaptive_avg_pool2d_backward(): Expected 3D or 4D tensor, but got r'  r   s   r5   rW   z4meta__adaptive_avg_pool2d_backward.<locals>.<lambda>&  s    TUYU_U_T`a r7   c                  <    dj                    d j                    S Nexpected dtype z! for `grad_output` but got dtype r   )r\  r   s   r5   rW   z4meta__adaptive_avg_pool2d_backward.<locals>.<lambda>*  s    /$**-Nx~~N^_ r7   r   )r   r   rH   rY   r   rP   r   r  r  r   r   r  )r\  r   r   r   r   s   ``  @r5   "meta__adaptive_avg_pool2d_backwardra    s    ==D1d^ 
MM!q f	

 
LL	TQYa 
LL

hnn$_ ++M++>>$**%((}(EEr7   c                 d    t        | d       t        j                  |t        j                        S )Nadaptive_avg_pool3d_backwardr   )!_adaptive_pool_empty_output_checkrH   r   r  r  r   s     r5   "meta__adaptive_avg_pool3d_backwardrf  2  s(     &k3QRD0N0NOOr7   r  c                       j                   }t        d|      D ]/  t        j                   j	                        dkD   fd       1 y )Nr   r   c                  .      dj                    d dS )Nzc(): Expected grad_output to have non-zero size for non-batch dimensions, but grad_output has sizes rZ  r[  r'  )r  r  r   s   r5   rW   z3_adaptive_pool_empty_output_check.<locals>.<lambda>>  s/    * --8->->,??OPQsR^` r7   )r   r   rH   rY   r   )r  r  r   r   s   `` @r5   rd  rd  9  sG    D1d^ 
Q!#	

r7   c                      j                   }t        j                  |dv  fd       t        d|      D ].  t        j                   j	                        dkD   fd       0 t        j                  t        |      dk(  d        d}d}d} j                   dk(  r j	                  d      }|dz  } j	                  |dz
        }|\  }} j                   d	k(  r;|||f} j                  |      }	 j                  |t        j                  
      }
|	|
fS ||||f}t        j                         } j                  |      j                  |      }	 j                  |t        j                  
      j                  |      }
|	|
fS )Nr-   r  c                  "    d j                    S )Nz:adaptive_max_pool2d(): Expected 3D or 4D tensor, but got: r'  r	  s   r5   rW   z*meta_adaptive_max_pool2d.<locals>.<lambda>K      LU[[MZ r7   r   r   c                  *    dj                    d  dS )Nzjadaptive_max_pool2d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes rZ  r[  r'  r   r   s   r5   rW   z*meta_adaptive_max_pool2d.<locals>.<lambda>P  %    '',{{m3CA3lT r7   r   c                       y)NzCadaptive_max_pool2d(): internal error: output_size.size() must be 2r;   r;   r7   r5   rW   z*meta_adaptive_max_pool2d.<locals>.<lambda>X  r_   r7   r  r-   r   r   )r   rH   rY   r   r   r   r   r   r@   r  r  )r   rS  r   dimHsizeBsizeDosizeHosizeWr   r   r   r   r   s   `           @r5   meta_adaptive_max_pool2drv  E  s|    ::D	LLZ 1d^ 
JJqMA	

 
LLKAU
 DEEzzQ

1	JJtax E NFFzzQFF+	ooi(//)5;;/?G|E662	33E:ooi(++-+H//)5;;/?BB' C 
 G|r7   c                 N     j                   }t        j                  |dv  fd       t         d       t        j                  j                   j                  k(   fd       t        j                        }j                  j                        j                  |      S )Nrj  c                  "    d j                    S )NzKadaptive_max_pooling2d_backward(): Expected 3D or 4D grad_output, but got: r'  r  s   r5   rW   z3meta_adaptive_max_pool2d_backward.<locals>.<lambda>{  s    ]^i^o^o]pq r7   adaptive_max_pool2d_backwardc                  <    dj                    d j                    S r_  r   )r  r   s   r5   rW   z3meta_adaptive_max_pool2d_backward.<locals>.<lambda>  s!    /%++.OP[PaPaObc r7   r   )
r   rH   rY   rd  rP   r@   r  r   r   r  )r  r   r   r   r   s   ``   r5   !meta_adaptive_max_pool2d_backwardr|  u  s     D	LLq
 &k3QR	LL{(((c
 //6M??5;;'***GGr7   c                      j                   }t        j                  |dv  fd       t        d|      D ].  t        j                   j	                        dkD   fd       0 t        j                  t        |      dk(  d        d}d}d}|dk(  r j	                  d      }|dz  } j	                  |      }|\  }}}|d	k(  r||||f}	n|||||f}	 j                  |	      }
 j                  |	t        j                  
      }|
|fS )Nr4  c                  "    d j                    S )Nz:adaptive_max_pool3d(): Expected 4D or 5D tensor, but got: r'  r	  s   r5   rW   z*meta_adaptive_max_pool3d.<locals>.<lambda>  rl  r7   r   r   c                  *    dj                    d  dS )Nzjadaptive_max_pool3d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes rZ  r[  r'  rn  s   r5   rW   z*meta_adaptive_max_pool3d.<locals>.<lambda>  ro  r7   r-   c                       y)NzCadaptive_max_pool3d(): internal error: output_size.size() must be 3r;   r;   r7   r5   rW   z*meta_adaptive_max_pool3d.<locals>.<lambda>  r_   r7   r  r  r   )r   rH   rY   r   r   r   r   r   )r   rS  r   dimDrr  rs  osizeTrt  ru  r   r   r   r   s   `           @r5   meta_adaptive_max_pool3dr    s    ::D	LLZ 1d^ 
JJqMA	

 
LLKAU
 DEEqy

1	JJtE(FFFqyFFF3	E666:	
//)
$Cooiu{{o;G<r7   c                 P    t        | d       |j                  |j                        S )Nadaptive_max_pool3d_backward)rd  r   r   )r  r   r   s      r5   !meta_adaptive_max_pool3d_backwardr    s"     &k3QR??5;;''r7   c                 >    |t        d      | j                  |      S )Nz:cannot repeat_interleave a meta tensor without output_size)r  r   )repeatsrS  s     r5   meta_repeat_interleave_Tensorr    s%    WXX[))r7   c                     | j                   j                  sJ |j                   j                  sJ t        | j                  |j                        }| j	                  |t        | j                               S rw  )rP   r   r$   r   r   r   )realimagr   s      r5   meta_complexr    s[     ::''''::''''!$**djj9I>>)+Ftzz+R>SSr7   )
fill_valuer  c                d    | j                  || j                         ft        j                        S rw  )r   rw   rH   r   )r   r   r  s      r5   nonzero_staticr    s&     >>4,EJJ>??r7   c                 
   t        j                  t        j                  d        t        j                  | j                         | j                         fd| j                         ft         j                  | j                        S )Nc                       y)NaY  The register_meta function for torch.nonzero() raises unimplemented by default, as a correct data-independent implementation does not exist. This implementation returns a fake value, assuming all elements of the tensor are non-zero. To enable this registration, please set 'torch.fx.experimental._config.meta_nonzero_assume_all_nonzero' to True.r;   r;   r7   r5   rW   znonzero.<locals>.<lambda>  r_   r7   r   rP   rt   )	rH   _check_not_implemented
exp_configmeta_nonzero_assume_all_nonzeror  r   rw   r   rt   r   s    r5   nonzeror    sf     
  22	S 	txxz"	
DJJLjj{{	 r7   c           
          t        j                  t              d        g }t              D ]  \  ft        j                  j                  t         j
                  t         j                  t         j                  t         j                  fv d        j                  t         j                  t         j                  fv rȉj                         }t        |      t        j                  j                  z    j                  k   fd       t        j                        D ]`  t        j                  j                      j                  z      k(   fd       |j                  |j                  d             b ]|j                         p|j                          |t        j                  t               j                  k   fd       dd lm} t%         |j&                         t               j                  k  r*j                  d        t               j                  k  r*d}d}D ]  |dk(  rd}|dk(  rd	} n d
}|sg }g }t              D ]*  \  	|j                         |j                         , t              D ]*  \  	|j                         |j                         ,  j)                  |       |g g g t              D ]\  \  }	@rj                   j                  |	          )j                   j                  |	          Ht%        j                        ^ fd}
 j+                  z   z         }ddlm}  | j1                         dk(        r|S  |
       }t3        j4                  |      }t%        |      t%        t        t        |                  k7  r~t3        j6                  |j                  |      }t3        j8                  |      }t3        j6                  |t3        j:                  |            }|j=                  |j?                         |      }|S )Nc                       y)Nz#at least one index must be providedr;   r;   r7   r5   rW   z#meta_index_Tensor.<locals>.<lambda>  r_   r7   c                       y)Nz?tensors used as indices must be long, int, byte or bool tensorsr;   r;   r7   r5   rW   z#meta_index_Tensor.<locals>.<lambda>  r_   r7   c                  "    d j                    S )N)too many indices for tensor of dimension r  r   s   r5   rW   z#meta_index_Tensor.<locals>.<lambda>  s    G		{S r7   c            	      N    dj                    d  dj                    dz    S )NzThe shape of the mask z
 at index z0 does not match the shape of the indexed tensor r'  )r   r   jrS  r   s   r5   rW   z#meta_index_Tensor.<locals>.<lambda>  s<    "8ZPQs SJJN**U_`ade`e_f!h r7   r   c                  <    dj                    dt                dS )Nr  z (got rk   )r   r   )r   r   s   r5   rW   z#meta_index_Tensor.<locals>.<lambda>  s!    ;DII;fSQX\NZ[\ r7   r   Fr   Tc                     z   z   }t        | j                               }dgt              z  |t              t        | j                        t              z
   | j	                  ||      S )zI
        This follows restride_src in TensorAdvancedIndexing.cpp
        r   )r   r   r   r   r4  )r   r   r   after_shapebefore_shapereplacement_shapes      r5   _restride_srcz(meta_index_Tensor.<locals>._restride_srcG  so     00;>t{{}%KL#PSQ
 K
L!C

Oc+6F$FG ug..r7   guard_size_oblivious) rH   rY   rU  	enumeraterP   r   r   r2  r  r   r   r   r   r   r   selecttorch._refs_refsr   r%   r   r   r  r  r   r@   3compute_elementwise_output_logical_to_physical_perm
apply_permr   invert_permr4  r   )r   r   r  r  refsstatehas_contiguous_subspacerp  transposed_indicesrw   r  r   r  restrided_selfperm
perm_shaper%  r  r  r   r   r  rS  r  s   ``               @@@@@@@r5   meta_index_Tensorr    s   	LLg MN &(Fg& !5LL

EIIuzz5::NNY {{uzz5::66--/K""

Ndii/S uzz* 8A&&A$**QU*;;h
 MM'..A"678 e$MM% /!0 G	LLG		!\
 (4(('23G
g,
"t g,
" E# 'A: aZ} ' #'
 #!'* 	1HAu A"))%0	1 "'* 	1HAu}A"))%0	1 ||D!$ !LK#%( 2
U= ""4::c?3##DJJsO4 $U[[ 12	/ ..(99KG
HCJDJJLA-.

 #4(NDD^TD DzT%D	*++%%cii6
66zB
%%j%2C2CD2IJ
nnSXXZ4Jr7   c                     d }d }d }|
d   r| j                  |j                               }|
d   r| j                  |j                               }|
d   r| j                  |      }|||fS )Nr   r   r   r   r   )grad_output_input_weight_bias_sizes_optr   r  rx  
transposedr{  rz  output_maskbackend_grad_inputbackend_grad_weightbackend_grad_biass                 r5   meta_convolution_backwardr  h  sy      1~)33FKKMB1~*44W\\^D1~(22>B 35FGGr7   c                   j                  d      }j                  d      }| j                  ||f      } t        j                  j	                         dk(  d        t        j                  j	                         dk(  d        t        j                  j                  d      j                  d      k(  fd       t        j                  j                  d      j                  d      k(  fd       t        j                  | j                  d      |k(  xr | j                  d      |k(  d	        | j                  | j                               S )
Nr   r   r-   c                       yr4  r;   r;   r7   r5   rW   zmeta_addbmm.<locals>.<lambda>  r_   r7   c                       yr6  r;   r;   r7   r5   rW   zmeta_addbmm.<locals>.<lambda>  r_   r7   r   c                  P    d j                  d       dj                  d       S )Nz8batch1 and batch2 must have same number of batches, got r   r   r   r8  r9  s   r5   rW   zmeta_addbmm.<locals>.<lambda>  s.    J6;;WX>JZZ_`f`k`klm`n_op r7   c            
          d j                  d       d j                  d       dj                  d       dj                  d       d	S )Nz#Incompatible matrix sizes for bmm (r   rE   r   r   rk   r   r  s   r5   rW   zmeta_addbmm.<locals>.<lambda>  sQ    1&++a.1A6;;q>BR S;;q>"!FKKN#316 r7   c                       y)Nz.self tensor does not match matmul output shaper;   r;   r7   r5   rW   zmeta_addbmm.<locals>.<lambda>  r_   r7   )r   rq  rH   rY   rw   r   )r   r8  r9  r?  r>  rA  rB  s    ``    r5   meta_addbmmr    s    ;;q>D;;q>D;;d|$D	LL"$HI	LL"$HI	LLA&++a.(p 
LLA&++a.(	
 
LL		!51!5@ >>$))+&&r7   )
grad_scale	found_infc       	         n    | |||||fD ])  t        j                  t        t              fd       + y )Nc                       dt                S Nz'exponent must be a tensor list but got rl   ls   r5   rW   z#meta__fused_adam_.<locals>.<lambda>      =d1gYG r7   rH   rY   rb   r   )r   gradsexp_avgsexp_avg_sqsmax_exp_avg_sqsstate_stepslrbeta1beta2weight_decayepsamsgradmaximizer  r  r  s                  @r5   meta__fused_adam_r    s:    & E8[/;O 
q$G	

r7   c       	             | |||||fD ])  t        j                  t        t              fd       + d } ||        ||       ||       ||       ||      fS )Nc                       dt                S r  r  r  s   r5   rW   z"meta__fused_adam.<locals>.<lambda>  r  r7   c                 R    | D cg c]  }t        j                  |       c}S c c}w r2   rP  )tensor_listr  s     r5   empty_like_listz)meta__fused_adam.<locals>.empty_like_list  s!    -89  #999s   $r  )r   r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  s                   @r5   meta__fused_adamr    ss    & E8[/;O 
q$G	

: 	!$( r7   c                 j    t        j                   j                         dk(  d        t        j                  j                         dk(  d        t        j                   j                  t         j                  u  fd       t        j                  j                  t         j                  u fd       t        j                   j                  d      j                  d      k(   fd        j                   j                  d      j                  d      ft         j                  	      S )
Nr   c                       y)Nza must be a 2D tensorr;   r;   r7   r5   rW   zmeta__int_mm.<locals>.<lambda>  r_   r7   c                       y)Nzb must be a 2D tensorr;   r;   r7   r5   rW   zmeta__int_mm.<locals>.<lambda>  r_   r7   c                  "    d j                    S )Nzexpected self to be int8, got r   )rk  s   r5   rW   zmeta__int_mm.<locals>.<lambda>      0	: r7   c                  "    d j                    S )Nzexpected mat2 to be int8, got r   )rl  s   r5   rW   zmeta__int_mm.<locals>.<lambda>  r  r7   r   r   c            
          d j                  d       d j                  d       dj                  d       dj                  d       d	S )Nz'Incompatible matrix sizes for _int_mm (r   rE   r   r   rk   r   rk  rl  s   r5   rW   zmeta__int_mm.<locals>.<lambda>  sM    5affQi[!&&) M66!9+Qqvvayk, r7   r   )rH   rY   rw   rP   r2  r   r   r3  r  s   ``r5   meta__int_mmr    s     
LLA>?	LLA>?	LL	5::: 
LL	5::: 
LL	q	QVVAY	
 ;;q	166!9-U[[;AAr7   c                 f    t        j                   j                         dk(  d        t        j                   j                  t         j                  u  fd        j                  d      } j                  d      dz  } j                  |dz  ||dz  z  d|dz  ft         j                  	      S )
Nr   c                       yNzw must be a 2D tensorr;   r;   r7   r5   rW   z2meta__convert_weight_to_int4pack.<locals>.<lambda>  r_   r7   c                  "    d j                    S Nr  r   r  s   r5   rW   z2meta__convert_weight_to_int4pack.<locals>.<lambda>      .qwwi8 r7   r   r      rK      r   )rH   rY   rw   rP   r  r   r   r3  r  inner_k_tilesr   rS  s   `   r5    meta__convert_weight_to_int4packr    s    	LLA>?	LL	5;;8 	
q	A	q	AA;;F-"$%Q		
 kk   r7   c                 J    t        j                   j                         dk(  d        t        j                   j                  t         j                  u  fd        j                  d      } j                  d      } j                  ||dz  ft         j                        S )Nr   c                       yr  r;   r;   r7   r5   rW   z:meta__convert_weight_to_int4pack_for_cpu.<locals>.<lambda>  r_   r7   c                  "    d j                    S Nzexpected w to be int32, got r   r  s   r5   rW   z:meta__convert_weight_to_int4pack_for_cpu.<locals>.<lambda>  r  r7   r   r   r   )rH   rY   rw   rP   r3  r   r   r  r  s   `   r5   (meta__convert_weight_to_int4pack_for_cpur   
  s    	LLA>?	LL	5;;8 	
q	A	q	A;;	
AFkk   r7   c                 .    t        j                   j                         dk(  d        t        j                  j                         dk(  d        t        j                   j                  t         j                  t         j
                  t         j                  fv  fd       t        j                  j                  t         j                  u fd        j                   j                  d      j                  d      dz   j                  	      S )
Nr   c                       yNzx must be a 2D tensorr;   r;   r7   r5   rW   z*meta__weight_int4pack_mm.<locals>.<lambda>  r_   r7   r  c                       y)Nzw must be a 4D tensorr;   r;   r7   r5   rW   z*meta__weight_int4pack_mm.<locals>.<lambda>  r_   r7   c                  "    d j                    S Nr  r   rE   s   r5   rW   z*meta__weight_int4pack_mm.<locals>.<lambda>      5aggY? r7   c                  "    d j                    S r  r   r  s   r5   rW   z*meta__weight_int4pack_mm.<locals>.<lambda>#  r  r7   r   r  r   )
rH   rY   rw   rP   rL  rM  rN  r3  r   r   r  s   ``  r5   meta__weight_int4pack_mmr
    s    	LLA>?	LLA>?	LL	EMM5==%..AA? 
LL	5;;8 ;;qvvay!&&)a-qww;??r7   c                 (    t        j                   j                         dk(  d        t        j                  j                         dk(  d        t        j                   j                  t         j                  t         j
                  t         j                  fv  fd       t        j                  j                  t         j                  u fd        j                   j                  d      j                  d       j                        S )Nr   c                       yr  r;   r;   r7   r5   rW   z2meta__weight_int4pack_mm_for_cpu.<locals>.<lambda>*  r_   r7   c                       yr  r;   r;   r7   r5   rW   z2meta__weight_int4pack_mm_for_cpu.<locals>.<lambda>+  r_   r7   c                  "    d j                    S r  r   r  s   r5   rW   z2meta__weight_int4pack_mm_for_cpu.<locals>.<lambda>.  r  r7   c                  "    d j                    S r  r   r  s   r5   rW   z2meta__weight_int4pack_mm_for_cpu.<locals>.<lambda>2  r  r7   r   r   )
rH   rY   rw   rP   rL  rM  rN  r  r   r   r  s   ``  r5    meta__weight_int4pack_mm_for_cpur  (  s    	LLA>?	LLA>?	LL	EMM5==%..AA? 
LL	5;;8 ;;qvvay!&&)177;;;r7   rk  rl  c                     | |z   dz
  |z  |z  S r   r;   r  s     r5   kai_roundupr  7  s    UQY1!!r7   c                   	
 | dk(  ry||k(  r(d}d}d}dddd fdfd} ||||||      S |dz  d	k(  rC||z  d	k(  r:d}d}d}dddd		fd
}	
fdd 
	fd	fd |||||||      S y y y )Nr  r  rK  r   c                 8    t        ||z  d      }t        | |      S )Nr  r  )rS  krsrkr_sr_roundedup4s       r5   kai_k_roundedupz3get_kai_packed_weight_size.<locals>.kai_k_roundedupF  s#     $/rBw#: "1&677r7   c                 X     | ||      }|dz  dk(  sJ d       ||dz  z   z   z   z  S )Nr   r   zk_internal must be evenr;   )	rS  nrr  r  
k_internalr  kai_num_bytes_biaskai_num_bytes_multiplier_rhskai_num_bytes_sum_rhss	        r5   9kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0z]get_kai_packed_weight_size.<locals>.kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0L  sY     -QB7
"Q1,G.GG,1_23+, )) r7   c                 >    t        | |      |z  }| ||||      z  S r2   r  )r   rS  r  r  r  num_rowsr   s         r5   7kai_get_rhs_packed_size_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0z[get_kai_packed_weight_size.<locals>.kai_get_rhs_packed_size_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0Z  s6     'q"-3 O2r2r7   r  r   c                 |    ||z  dk(  sJ |	z  dk(  sJ |z  dk(  sJ t        | |      |z  }| |||||      z  S ra  r  )
r   rS  r  r  r  blr"  kai_bl_multiple_of;kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0kai_nr_multiple_ofs
          r5   9kai_get_rhs_packed_size_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0z]get_kai_packed_weight_size.<locals>.kai_get_rhs_packed_size_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0r  sp     RA~%~//A555//A555&q"-3 Q2r2rr7   c                     ||z  dk(  sJ |
z  dk(  sJ |z  dk(  sJ  	       } | |      } ||      }|||z  z   z   z  S ra  r;   )rS  r  r  r  r%  num_bytes_multiplier_rhsnum_blocks_per_rownum_bytes_per_blockr&  #kai_get_bf16_datatype_size_in_bytesr(  kai_num_blocks_per_rowr  kai_num_bytes_per_blockr  s           r5   r'  z_get_kai_packed_weight_size.<locals>.kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0  s     RA~%~//A555//A555 ,O+P(%;Ar%B"&=0'# (+==+,() r7   c                       y)Nr   r;   r;   r7   r5   r.  zGget_kai_packed_weight_size.<locals>.kai_get_bf16_datatype_size_in_bytes  s    r7   c                 6    |z  dk(  sJ t        | |      |z  S ra  r  )rS  r%  r&  s     r5   r/  z:get_kai_packed_weight_size.<locals>.kai_num_blocks_per_row  s)    //A555"1b)R//r7   c                 (    | z  dk(  sJ | dz  |z   S )Nr   r   r;   )r%  r+  r&  s     r5   r0  z;get_kai_packed_weight_size.<locals>.kai_num_bytes_per_block  s'    //A555a#;;;r7   r;   )n_bitsrj  K	groupsizekai_nrkai_krkai_srr#  r)  r&  r.  r'  r   r  r(  r/  r  r  r0  r  s            @@@@@@@@@@@r5   get_kai_packed_weight_sizer:  ;  s    {>FFF$%!+,(!"8
 K1fff  ^q Q]a%7FFF$%!!"!"!#  ,0< M1fffi u &8 [ r7   c                 V    t        j                   j                  t         j                  u  fd       t         j                  j
                  j                         r||k(  r|j                  t         j                  k(  s2||k  re|dz  dk(  r]||z  dk(  rU|j                  t         j                  k(  r8t        d|||      } j                  t        |      t         j                        S  j                         |j                         z   } j                  |t         j                        S )Nc                  "    d j                    S r  r   )weightss   r5   rW   z2meta__dyn_quant_pack_4bit_weight.<locals>.<lambda>  s    .w}}o> r7   r  r   r  r   )rH   rY   rP   r  backendskleidiaiis_availablerL   rN  r:  r   r   r   )r=  scales_zerosr-  
block_sizein_featuresout_featurespacked_weight_sizes   `      r5    meta__dyn_quant_pack_4bit_weightrF    s     
LL$> ~~++-	{	"|'9'9U[['H$R1$j(A-""enn4 8|[*
   %7!8 LL <+=+=+??/u{{CCr7   c                     t        j                   j                         dk(  d        t        j                   j                  t         j                  fv  fd        j                  d      } j                  || j                        S )Nr   c                       y)Nzinput must be a 2D tensorr;   r;   r7   r5   rW   z-meta__dyn_quant_matmul_4bit.<locals>.<lambda>  r_   r7   c                  "    d j                    S )Nzexpected input to be f32, got r   )inps   r5   rW   z-meta__dyn_quant_matmul_4bit.<locals>.<lambda>  s    0< r7   r   r   )rH   rY   rw   rP   rL  r   r   )rJ  packed_weightsrB  rC  rD  r  s   `     r5   meta__dyn_quant_matmul_4bitrL    sg     
LLa!DE	LL		emm_$< 	A==L		=::r7   c                 (    t        j                   j                         dk(  d        t        j                   j                  t         j                  t         j
                  t         j                  fv  fd       t        j                  j                         dk(  d        t        j                  j                  t         j                  u fd        j                   j                  d      j                  d       j                        S )Nr   c                       yr  r;   r;   r7   r5   rW   z*meta__weight_int8pack_mm.<locals>.<lambda>  r_   r7   c                  "    d j                    S r  r   r  s   r5   rW   z*meta__weight_int8pack_mm.<locals>.<lambda>  r  r7   c                       yr  r;   r;   r7   r5   rW   z*meta__weight_int8pack_mm.<locals>.<lambda>  r_   r7   c                  "    d j                    S )Nzexpected w to be int8, got r   r  s   r5   rW   z*meta__weight_int8pack_mm.<locals>.<lambda>  s    -aggY7 r7   r   r   )
rH   rY   rw   rP   rL  rM  rN  r2  r   r   )rE   r  q_scaless   `` r5   meta__weight_int8pack_mmrS    s    	LLA>?	LL	EMM5==%..AA? 
LLA>?	LL	5::7 ;;qvvay!&&)177;;;r7   c                 f    t        j                   j                         dk\   fd       t        j                  j                         dk\  fd       t        j                   j                  d      j                  d      k(   fd       t        j                  t	        j
                   j                        d        t        j                  t	        j
                  j                        d        t        j                  |dk\  d	        t        j                  d
v fd        j                  d      }j                  d      } j                  d d }j                  d d }t        t        j                  ||            }|j                  ||g        j                  |      S )Nr   c                  ,    d j                          dS )Nz1cdist only supports at least 2D tensors, X1 got: r  r   )x1s   r5   rW   z$meta_cdist_forward.<locals>.<lambda>      CBFFH:QO r7   c                  ,    d j                          dS )Nz1cdist only supports at least 2D tensors, X2 got: r  r   )x2s   r5   rW   z$meta_cdist_forward.<locals>.<lambda>  rW  r7   r   c                  P    d j                  d       dj                  d       S )Nz4X1 and X2 must have the same number of columns. X1: r   z X2: r   )rV  rY  s   r5   rW   z$meta_cdist_forward.<locals>.<lambda>  s,    Frwwr{mSXY[Y`Y`acYdXef r7   c                       y)Nz=cdist only supports floating-point dtypes, X1 got: {x1.dtype}r;   r;   r7   r5   rW   z$meta_cdist_forward.<locals>.<lambda>  r_   r7   c                       y)Nz=cdist only supports floating-point dtypes, X2 got: {x2.dtype}r;   r;   r7   r5   rW   z$meta_cdist_forward.<locals>.<lambda>  r_   r7   r   c                       y)Nz)cdist only supports non-negative p valuesr;   r;   r7   r5   rW   z$meta_cdist_forward.<locals>.<lambda>  r_   r7   Nr   r   c                      d  S )Nz%possible modes: None, 1, 2, but was: r;   )compute_modes   r5   rW   z$meta_cdist_forward.<locals>.<lambda>  s    7~F r7   r  )rH   rY   rw   r   r@   is_float_dtyperP   r   r   broadcast_shapesextendr   )	rV  rY  r(  r`  r1r2batch_tensor1batch_tensor2rU  s	   `` `     r5   meta_cdist_forwardrh    sJ   	LL
AO 
LL
AO 
LL
rwwr{"f 
LLRXX&O 
LLRXX&O 
LLaLM	LL$F 
B	BHHSbMMHHSbMM..}mLMLR!<<%%r7   c                 4   |j                   d   }|j                   d   }|j                   d   }|j                   d d }|j                   d d }	t        t        j                  ||	            }
|
j	                         }|j                  ||g       t        j                  |
      }|dk(  s|dk(  s
|dk(  s|dk(  rt        j                  |      S |t        |j                         k7  r|j                  |      }t        j                  |t        j                        S )Nr   r  r   r   )r   r   rH   rb  copyrc  mathprod
zeros_likerq  r   r   )r,  rV  rY  r(  cdistc1rd  re  rf  rg  rm  tensor1_expand_sizebatch_products                r5   meta_cdist_backwardrr    s     
"B	"B	"BHHSbMMHHSbMM 6 6}m TU.335Bx(II23M	Qw"'R1W(:##d288n,YY*+Be.E.EFFr7   c	                     t        j                  j                  t         j                  t         j                  fv fd       t        j                  j                  t         j                  t         j                  fv fd       t        j                  t        j                   j                         fd       j                  d      }	|rt        j                  |	dk\  d        |	dz  }	 j                  |	 j                  d            }
}t        j                  |t        k(  d        t        j                  j                  dk(  fd       t        j                  j                         j                         k(  fd	       fd
d fd}t              dk7  r|j                  j                  d            }j                  j                               }|t        k(  r"j                  |	 j                  d            }nj                  d      }n | |
|      }|t        t        fv s|s!j                  j                  d            }nj                  d      }j                  |	      }j                  d   }|t        k(  rA|rt        j                  |dk\  d        |dz  }j                  | j                  d         }nj                  |j                               }|
|||fS )Nc                  "    d j                    S )Nz(expected indices to be long or int, got r   )r   s   r5   rW   z$meta_embedding_bag.<locals>.<lambda>*      :7==/J r7   c                  "    d j                    S )Nz(expected offsets to be long or int, got r   )re  s   r5   rW   z$meta_embedding_bag.<locals>.<lambda>.  ru  r7   c                  "    d j                    S )Nz/expected weight to be floating point type, got r   )r+  s   r5   rW   z$meta_embedding_bag.<locals>.<lambda>2  s    A&,,P r7   r   r   c                       yNz1include_last_offset: numBags should be at least 1r;   r;   r7   r5   rW   z$meta_embedding_bag.<locals>.<lambda>9  r_   r7   c                       y)Nz@embedding_bag: per_sample_weights only supported with mode='sum'r;   r;   r7   r5   rW   z$meta_embedding_bag.<locals>.<lambda>B  r_   r7   c                  $    d j                    dS )Nz1expected per_sample_weights to be 1D tensor, got r  r  )per_sample_weightss   r5   rW   z$meta_embedding_bag.<locals>.<lambda>F  s    GHZH_H_G``ab r7   c                  N    dj                          d j                          dS )Nz%expected per_sample_weights.numel() (z$ to be the same as indices.numel() (rk   r   )r   r|  s   r5   rW   z$meta_embedding_bag.<locals>.<lambda>J  s4    78J8P8P8R7S T66=mmo5FaI r7   c                 D     | ||      xr |j                  d      dk(  S Nr   r   r   )r  rZ  r   padding_idxis_fast_path_index_selects       r5   is_fast_path_index_select_scalez;meta_embedding_bag.<locals>.is_fast_path_index_select_scaleP  s(    %c6;?XELLQROWXDX	
r7   c                     | j                   t        j                  k(  xs | j                   t        j                  k(  xr1 | j	                  d      dk(  xr |j	                  d      dk(  xr |dk  S Nr   r   )rP   rH   rL   rJ   r   )r  r   r  s      r5   r  z5meta_embedding_bag.<locals>.is_fast_path_index_selectU  sb    YY%++%@ejj)@  

1" a A%  a		
r7   c                 2    | | |||      S  | ||      S r2   r;   )r  rZ  r   r  r  r  s       r5   is_fast_pathz(meta_embedding_bag.<locals>.is_fast_path]  s)    23v{SS,S&+FFr7   cpuc                       yry  r;   r;   r7   r5   rW   z$meta_embedding_bag.<locals>.<lambda>w  r_   r7   )rH   rY   rP   r   r   r@   ra  r   r   MODE_SUMr   r   r   MODE_MAX	MODE_MEANr   )r+  r   re  scale_grad_by_freqrJ  sparser|  include_last_offsetr  num_bagsr   r  
offset2bagbag_sizemax_indicesfast_path_sumnumBagsr  r  s   ```   `          @@r5   meta_embedding_bagr    s}    
LL%**eii00J 
LL%**eii00J 
LLV\\*P
 ||AHMG	
 	AhA7F%HV	
 	##q(b	
 	$$&'--/9	



G 7u$&&w||A7
$$W\\^48!++Hfkk!nEK!++A.K$V-?UIx(( **7<<?;J **1-J$$X.--"8"qLO 1!++GV\\!_EK!++HMMO<K:x44r7   c                     t        | ||g| \  }}}}t        |      dk(  r|j                  |j                               }||||fS )Nr  )r  r   r   r   )r+  r   re  rB   r   r  r  r  s           r5   meta_embedding_bag_forward_onlyr    sX    0B1#'1-FJ+ 7u$$$W\\^4:x44r7   c                     |r|S | j                   j                  s| j                   j                  r| j                   S |rt        j                  S | j                   S r2   )rP   r   r   rH   r   )r   rP   promote_int_to_longs      r5   _get_reduction_dtyper    sD    {{$$(>(>{{	zz;;r7   r   c                    t        | |d      }t        j                  | j                  |      }t	        | ||      }| j                  ||      S )NT)r  r   )r  r@   ry  r   rz  r   )r   rp  r|  rP   r  rU  s         r5   meta_nansumr    sI     (u$OLT2D+E4AL??<|?<<r7   c           	          t        j                  | j                  t        t	        | j                                           }| j                  |      S r2   )r@   rq  r   rX   r   rw   r   )r   rU  s     r5   meta_medianr    s<    77U5-.L ??<((r7   c                    t        |       dk(  rt        j                  d       t        j                  | j                  |f      }t        | ||      }| j                  |      | j                  |t        j                        fS )Nr   zmedian CUDA with indices outputr   )	r   r@   alert_not_deterministicry  r   rz  r   rH   r   )r   rw   r|  rU  s       r5   meta_median_mode_dimr    sp     5V#%%&GH


u{{SF
3C+E3@L%EJJ7 r7   c                     | S r2   r;   r   s    r5   meta_logical_not_r    rL  r7   c                 R   t        j                  t        |      | j                         k\  d        t        |      | j                         z
  }d|z  t	        | j
                        z   }t        t        |            D cg c]  }||   ||   z   }}| j                  |      S c c}w )Nc                       y)NzZNumber of dimensions of repeat dims can not be smaller than number of dimensions of tensorr;   r;   r7   r5   rW   zmeta_repeat.<locals>.<lambda>  r_   r7   r  )rH   rY   r   rw   rX   r   r   r   )r   r  num_new_dimensionspadded_sizer   target_sizes         r5   meta_repeatr    s    	LLG
"l W
2++eDJJ.??K8=c'l8KL1;q>GAJ.LKL>>+&& Ms   ?B$c                     | S r2   r;   r   s    r5   
meta_zero_r    rL  r7   c                 z    t        |t        j                        r t        | j                  |j                         | S r2   )rb   rH   r   r[   r   r   r   s     r5   meta_binop_inplacer    s)     %&

EKK8Kr7   c                     d }d }d } ||       r ||      rt        d       ||       r ||      st        d      t        |t        j                        r t	        | j
                  |j
                         | S )a*  
    Some checks for inplace ops.
    Checks for promotion rules for some dtypes.
    int.add/sub_(float) and bool.add/sub_(others) are rejected.
    Promoting in these in-place operations would require reallocating
    and copying over elements, hence not allowed.
    Checks for alpha param.
    c                     t        | t              rt        j                  | j                        S t        | t
              S r2   )rb   r   r@   r,  rP   r   re   s    r5   is_integericz.meta_binop_inplace_alpha.<locals>.is_integeric  s.    c:&))#))44c7++r7   c                     t        | t              rt        j                  | j                        S t        | t
              S r2   )rb   r   r@   ra  rP   r   r  s    r5   
is_floaticz,meta_binop_inplace_alpha.<locals>.is_floatic   s.    c:&''		22c9--r7   c                     t        | t              rt        j                  | j                        S t        | t
              S r2   )rb   r   r@   is_boolean_dtyperP   r   r  s    r5   is_booleanicz.meta_binop_inplace_alpha.<locals>.is_booleanic  s.    c:&))#))44c8,,r7   z]Promotion of int.add/sub_(float) in in-place ops are not possible due to element size change.z_Promotion of book.add/sub_(others) in in-place ops are not possible due to element size change.)r  rb   rH   r   r[   r   )r   r   r>  r  r  r  s         r5   meta_binop_inplace_alphar    sz    $,.- Dj/k
 	

 D,u"5m
 	
 %&

EKK8Kr7   c                 8    t        | t        j                        S Nr=   rF   r   rA   )r   kwargss     r5   
meta_roundr    s    <DD r7   c                 l    t        j                  t        j                  j                         fd       t        t         j                        r8t        j                  t        j                  j                         fd       y t        j                  t        t               fd       y )Nc                  &      dj                    S )Nz7: Expected input tensor to have an integral dtype. Got r   )r  r   s   r5   rW   z#shift_dtype_check.<locals>.<lambda>'  s    7)RSWS]S]R^_ r7   c                  &      dj                    S )Nz6: Expected shift value to have an integral dtype. Got r   r  r  s   r5   rW   z#shift_dtype_check.<locals>.<lambda>,  s    wiUVYV_V_U`a r7   c                        d S )Nz): Expected shift value to be an int. Got r;   r  s   r5   rW   z#shift_dtype_check.<locals>.<lambda>1  s    wiHN r7   )rH   rY   r@   r,  rP   rb   r   r   )r  r   r  s   ```r5   shift_dtype_checkr  $  sp    	LLtzz*_ #u||$""399-a	

 	sG$N	
r7   c                 T    t        d| |       t        | |t        j                        S )Nrshiftr  r  rF   r   rA   r  s     r5   meta_rshiftsr  5  )    he,e$C$K$K r7   c                 T    t        d| |       t        | |t        j                        S )Nlshiftr  r  r  s     r5   meta_lshiftsr  =  r  r7   c                 8    | j                  | j                        S r2   r  r   s    r5   	meta_zeror  E  s    >>$**%%r7   c                     | S r2   r;   r   r  s     r5   
meta_fill_r  J  rL  r7   c                 ,    t        j                  |       S r2   rP  r  s     r5   	meta_fillr  O      D!!r7   c                     | S r2   r;   r   s    r5   
meta_relu_r  T  rL  r7   c                 :    t        | |t        j                        S r  r  )r   r   r>  s      r5   meta__add_relur  Y  s     e$C$K$K r7   c                 ,    t        j                  |       S r2   rP  r   noiselowerr  r  r   s         r5   meta_rrelu_with_noiser  a  s    
 D!!r7   c                 V    t        j                  |       t        j                  |      fS r2   rP  r  s         r5    meta_rrelu_with_noise_functionalr  i  s%     D!5#3#3E#:::r7   c                     | S r2   r;   )r   r  r  r  r   s        r5   meta_rrelu_with_noise_r  p  s	     Kr7   c                 ,    t        j                  |       S r2   rP  r   r   r   
accumulates       r5   meta_index_putr  w  r  r7   c                 F    t        | j                  |j                         | S r2   r[   r   )r   rb  values      r5   meta_masked_fill_r  |  s    DJJ

3Kr7   c                     | j                  | j                               j                  t        j                  |             }|S r   )r   r   r  r@   r  )r   rb  rZ  masked_scales       r5   meta__masked_scaler    s<    >>$))+.1111$7 2 L r7   c                      t        j                  |j                  t         j                  t         j                  fv d        t        j                   j                  j                  k(   fd        S )Nc                       y)NzMask must be bool or uint8r;   r;   r7   r5   rW   z&meta_masked_scatter_.<locals>.<lambda>  r_   r7   c                  <    d j                    dj                    S )NzEmasked_scatter: expected self and source to have same dtypes but got r   r   )r   rX  s   r5   rW   z&meta_masked_scatter_.<locals>.<lambda>  s"     **U6<<.: r7   )rH   rY   rP   rU  r  )r   rb  rX  s   ` `r5   meta_masked_scatter_r    sU    	LL

uzz5;;//1U 
LL

fll"	:
 Kr7   c                     t        | |      \  } }t        j                  | t        j                        }t	        |||      S r   )r%   rH   r   r   r  )r   rb  rX  r   s       r5   meta_masked_scatterr    s;     "$-JD$d%2I2IJFf55r7   c                 $    | j                  |      S r2   rs  )r   rb  r5  s      r5   meta_masked_scatter_backwardr    s    >>%  r7   c                     | S r2   r;   r  s       r5   meta_index_put_r    rL  r7   c                 8    | j                  | j                        S r2   )viewr   r   s    r5   
meta_aliasr    s    99TZZ  r7   c                 H  	
 t        j                  | j                         dk(  d        t        j                  |j                         dk(  d        | j                         }|j                         |d   	|d   
|d   }d   }	||ft        j                  d   	k(  xr d   
k(  	
fd       |j	                        }|sUSt        j                  j                         dk(  d        t        j                  j                         k(  fd	       |S )
Nr-   c                       yr4  r;   r;   r7   r5   rW   z)common_meta_baddbmm_bmm.<locals>.<lambda>  r_   r7   c                       yr6  r;   r;   r7   r5   rW   z)common_meta_baddbmm_bmm.<locals>.<lambda>  r_   r7   r   r   r   c            	      .    d d d d    d d    d	S r;  r;   r=  s   r5   rW   z)common_meta_baddbmm_bmm.<locals>.<lambda>  s5    RSURV
l<?*;2l1o=NbR r7   c                       y)Nzself must be a 3D tensorr;   r;   r7   r5   rW   z)common_meta_baddbmm_bmm.<locals>.<lambda>  r_   r7   c                  0    d  dj                          S )Nz*Expected an input tensor shape with shape z but got shape: r   )rS  self_baddbmms   r5   rW   z)common_meta_baddbmm_bmm.<locals>.<lambda>  s!    @M]^j^o^o^q]rs r7   )rH   rY   rw   r   r   )r8  r9  is_bmmr  rD  res_rowsres_colsr   r>  r?  r@  rS  s      `    @@@@r5   common_meta_baddbmm_bmmr    s	   	LL"$HI	LL"$HI;;=L;;=L	aB#AAHAHx*K	LLQ2E,q/5E"E	R k*Fl.\%%'1,.PQ;.s	

 Mr7   c                     t        | |d      S )NT)r  )r   r:  s     r5   meta_bmmr    s    "4t44r7   c                 h    | |z  }| |z  }|dk7  r"t        |dk        t        |dk        k7  r|dz  }|S r  )rU  )rE   yqrN  s       r5   div_rtnr    sB    	QA	AA 	Av4A;$q1u+-	QHr7   c                     t        | |z   |z   ||dz
  z  z
  dz
  |r|dz
  ndz   |      dz   }|r|dz
  |z  | |z   k\  r|dz  }|S r  )r  )	inputSize
kernelSizer  r  r   rx  r  
outputSizes           r5   pooling_output_shape_pad_lrr    s     	 *q.)* 	
 'vzA/ 	
 		  Nf$	E(99!OJr7   c           	          t        j                  |dk7  d        t        j                  dk\  fd       t        j                  dz
  z  dz   dz  k  fd       t        | ||      S )Nr   c                       y)Nzstride should not be zeror;   r;   r7   r5   rW   z&pooling_output_shape.<locals>.<lambda>  r_   r7   c                      d  S )Nz'pad must be non-negative, but got pad: r;   )pads   r5   rW   z&pooling_output_shape.<locals>.<lambda>  s    %LSE#R r7   r   r   c                      d d d  S )NzApad should be at most half of effective kernel size, but got pad=z, kernel_size=z and dilation=r;   )rx  r  r  s   r5   rW   z&pooling_output_shape.<locals>.<lambda>  s'    OPSu U%,nXJ@ r7   )rH   rY   r  )r  r  r  r   rx  r  s    `` ` r5   r  r    ss    	LL1AB	LLRS	LLa8+a/A55	
 ':sC9 r7   c           	      >   	
  j                         }	t        j                  dkD  xr dkD  d        t        j                  |dkD  xr |dkD  d        t        j                  |dkD  xr |dkD  d         j                  d      dk7  xr  j                  d      dk7  }|t        j                  k(  r5t        j                  |dk(  xr |xr  j                  d      dk7  d	        nWt        j                  |dk(  xr  j                  d      dk7  xr |xs |dk(  xr |xr  j                  d      dk7   fd
       t        j                  dz  k\  xr dz  k\  fd       t        j                  dk\  xr dk\  
	fd       y )Nr   c                       y)NzCkernel size should be greater than zero, but got kH: {kH}, kW: {kW}r;   r;   r7   r5   rW   z$pool2d_shape_check.<locals>.<lambda>  r_   r7   c                       y)Nz>stride should be greater than zero, but got dH: {dH}, dW: {dW}r;   r;   r7   r5   rW   z$pool2d_shape_check.<locals>.<lambda>"  r_   r7   c                       y)Nz\dilation should be greater than zero, but got dilationH: {dilationH}, dilationW: {dilationW}r;   r;   r7   r5   rW   z$pool2d_shape_check.<locals>.<lambda>&  r_   r7   r   r   r  r-   c                       y)NzExpected 4D (batch mode) tensor expected for input with channels_last layout with optional 0 dim batch size for input, but got: {input.size()}r;   r;   r7   r5   rW   z$pool2d_shape_check.<locals>.<lambda>.  r_   r7   c                  *    d j                          S )NzYExpected 3D or 4D (batch mode) tensor with optional 0 dim batch size for input, but got: r   r	  s   r5   rW   z$pool2d_shape_check.<locals>.<lambda>5  s    opupzpzp|o}~ r7   c                       d d d d  S )NzKpad should be smaller than or equal to half of kernel size, but got padW = z	, padH = z, kW = z, kH = r;   )r  r  r  r  s   r5   rW   z$pool2d_shape_check.<locals>.<lambda>:  s&     ygbT> r7   c                  .    d d  d d d d dS NzGiven input size: (rE   z). Calculated output size: (z). Output size is too smallr;   )r  r  r  r!  r  r  s   r5   rW   z$pool2d_shape_check.<locals>.<lambda>@  s:    %k]!K=* N$$0><.+ O## r7   )rw   rH   rY   r   r  )r   r  r  r  r  r  r  	dilationH	dilationWr  r  r  r  r  r   r   
valid_dimsr!  s   ```  ``  `````   @r5   r  r    s   " 99;DL	LL
Q26U 
LL
Q26P 
LLA')a-n
 A!#:

1(:J+++AI;*;A!);Q	
 	QY<5::a=A-<* A	?j?UZZ]a-?~	
 
LL
a4+B!GtO	> 
LLq.\Q.	# 	#r7   r=  r:  r  r  r;  r  r  pTpHpW	dilationTr   r!  r>  r?  r@  rA  rB  rC  r8  c           
      J   	
  j                   }t        j                  dkD  xr dkD  xr dkD  fd       t        j                  dkD  xr dkD  xr dkD  fd       t        j                  dkD  xr dkD  xr dkD  fd       t        j                  |dv  fd       t        |      D ]:  |dk(  rdk(  rt        j                   j	                        dkD   fd       < |r/t        j                  k\  xr k\  xr k\  fd	       t        j                  d
z  k\  xr d
z  
k\  xr d
z  	k\  	
fd       t        j                  dk\  xr dk\  xr dk\  fd       y )Nr   c                      d d  d S )Nz5kernel size should be greater than zero, but got kT: z, kH: z, kW: r;   )r  r:  r  s   r5   rW   z$pool3d_shape_check.<locals>.<lambda>b  s#    $fRDrd, r7   c                      d d  d S )Nz0stride should be greater than zero, but got dT: z, dH: z, dW: r;   )r  r;  r  s   r5   rW   z$pool3d_shape_check.<locals>.<lambda>i  s     >rd&FSURVW r7   c                      d d  d S )Nz9dilation should be greater than zero, but got dilationT: z, dilationH: z, dilationW: r;   )r   r&  r!  s   r5   rW   z$pool3d_shape_check.<locals>.<lambda>o  s$    #M)M)V r7   r4  c                  &      dj                    S )Nz/: Expected 4D or 5D tensor for input, but got: r'  )r  r   s   r5   rW   z$pool3d_shape_check.<locals>.<lambda>w  s    7)J5;;-X r7   r  c                  L      dj                    dj                         dS )NzZ: Expected input's non-batch dimensions to have positive length, but input has a shape of z and non-batch dimension z has length zero!)r   r   )r  r   r   s   r5   rW   z$pool3d_shape_check.<locals>.<lambda>  s3    ) --2[[M+EJJqM?:KM r7   c                  .    d d  d d d d dS )Nzinput image (T: r  r  z ) smaller than kernel size (kT:  kH:  kW: rk   r;   )r?  r>  r@  r  r:  r  s   r5   rW   z$pool3d_shape_check.<locals>.<lambda>  s9    "5'gYd6( C$$&4uRDbT< r7   r   c                  ,    d d d  d d d S )NzHpad should be smaller than or equal to half of kernel size, but got kT: r/  r.  z padT: z padW: z padH: r;   )r  r:  r  r$  r#  r%  s   r5   rW   z$pool3d_shape_check.<locals>.<lambda>  s6    $eB4uRDt72$gbTK r7   r   c                  :    d d d  d d d d d dS r  r;   )r?  r>  r@  r=  rB  rA  rC  s   r5   rW   z$pool3d_shape_check.<locals>.<lambda>  sI    !'!E7!G9AfX F((/y%'!F8 L'( r7   )r   rH   rY   r   r   )r   r=  r:  r  r  r;  r  r  r#  r$  r%  r&  r   r!  r>  r?  r@  rA  rB  rC  r  r8  r   r   s   `````````````````````  @r5   r9  r9  F  s   0 ::D	LL
Q$26$b1f	
 
LL
Q$26$b1f	
 
LLA9)a-9IM	
 
LLX
 4[ 
19aJJqMA	
	
 RK:GrM:fl 	
 
LL
Q"6a26"q&B,	
 	
 
LL
3v{3w!|	
 	
r7   c                 j   | j                   }t        | |||||||	|
||||||||||||       t        |||dz
  |       t        |||dz
  |       t        |||dz
  |       t        |||dz
  |       t        |||dz
  |       t        |||dz
  |       t        |||dz
  |       t        |||dz
  |       y )Nr  r-   r   r   r   r9  r  )r   r  r   r=  r:  r  r  r;  r  r  r#  r$  r%  r&  r   r!  r>  r?  r@  rA  rB  rC  r  r   s                           r5   max_pool3d_backward_shape_checkr4    s    2 ::D








+0 ;dQh8;dQh6;dQh8;dQh77D$(G47D$(E27D$(G47D$(F3r7   c                     | j                   }t        | ||||||||	|
|ddd|||||||d       t        |||dz
  |       t        |||dz
  |       t        |||dz
  |       t        |||dz
  |       y )Nr   Tr  r-   r   r3  )r   r  r=  r:  r  r  r;  r  r  r#  r$  r%  r>  r?  r@  rA  rB  rC  r  r   s                       r5   rK  rK    s    * ::D








			-2 ;dQh8;dQh6;dQh8;dQh7r7   c                    d } |d|      \  }}t        j                  t        |      dv d        t        |      dk(  r||}
}	n |d|      \  }	}
 |d|      \  }} |d|      \  }}| j                  d	      }| j                  d
      }| j                  d      }t	        j
                  |       }|t         j                  k(  r)t        j                  | j                         dk(  d        nR|t         j                  k(  r(t        j                  | j                         dv d        nt        j                  dd        t        ||||	||      }t        ||||
||      }t        | |||	|
||||||||||       |||fS )Nc                      t        j                  t        |      dv  fd       |d   }t        |      dk(  r|n|d   }||fS )Nr  c                      d  dS )Nzmax_pool2d: r  r;   r  s   r5   rW   zEmax_pool2d_checks_and_compute_shape.<locals>.unpack.<locals>.<lambda>"  r  r7   r   r   r  r  s   `   r5   r   z3max_pool2d_checks_and_compute_shape.<locals>.unpack  r  r7   r  r  c                       y)NzOmax_pool2d: stride must either be omitted, a single int, or a tuple of two intsr;   r;   r7   r5   rW   z5max_pool2d_checks_and_compute_shape.<locals>.<lambda>,  r_   r7   r   r   r  rx  r  r  r   r  c                       y)NzMnon-empty 4D (batch mode) tensor expected for input with channels_last layoutr;   r;   r7   r5   rW   z5max_pool2d_checks_and_compute_shape.<locals>.<lambda>=  r_   r7   rj  c                       y)Nz9non-empty 3D or 4D (batch mode) tensor expected for inputr;   r;   r7   r5   rW   z5max_pool2d_checks_and_compute_shape.<locals>.<lambda>B  r_   r7   Fc                       y)Nz?Unsupport memory format. Supports only ChannelsLast, Contiguousr;   r;   r7   r5   rW   z5max_pool2d_checks_and_compute_shape.<locals>.<lambda>G  r_   r7   )rH   rY   r   r   r@   r  r  rw   r   r  r  )r   r  r   r  rx  r  r   r  r  r  r  r  r  r   r!  r  r  r  r   r  r  s                        r5   r  r    s    M;/FB	LLFy a 6{aRB&)B	7+JD$!*h7Iy**R.K**R.KBJ//6M+++IIK1c	
 
%11	1IIK6!O	

 	U	

 (Rr9iXL&z2tRIVK



$ k11r7   c                 |    t        |||||      \  }t        j                  j                   j                  k(   fd       |j                  fd}	 |	         |	|       t        j                        }
t        j                  j                  j                  j                  |
      S )Nc                  <    dj                    d j                    S )NzExpected dtype z  for `gradOutput` but got dtype r   re  s   r5   rW   z7meta_max_pool2d_with_indices_backward.<locals>.<lambda>w  s     /$**-MkN_N_M`a r7   c                 l    t        | dz
         t        | dz
         t        | dz
         y )Nr-   r   r   )r  )r  r!  r   r  r  s    r5   _check_dim_sizez>meta_max_pool2d_with_indices_backward.<locals>._check_dim_size}  s9    q$q,7q$q,7q$q+6r7   rl  )
r  rH   rY   rP   r   r@   r  r|   r   rt   )r  r   r  r   r  rx  r  r   r  r@  r   r!  r   r  r  s   ``         @@@@r5   %meta_max_pool2d_with_indices_backwardrA  b  s     	,k67Hi		
 
LL

k'''a
 L99D7
 K G//5M;;

jj{{#	 r7   c                    t        | |||||      \  }}}| j                         dk(  r| j                  d      nd}	t        j                  |       }
| j                         dk(  r|||g}n|	|||g}t        j                  || j                  | j                  |
      t        j                  |t
        j                  | j                  |
      fS r  )
r  rw   r   r@   r  rH   r|   rP   rt   r   r  s               r5   meta_max_pool2d_with_indicesrC    s     	,{FGXy		
  %yy{a/UZZ^QF//6Myy{a\;7\;?++<<'		
 	++<<'		
 r7   c           	         
 t        j                   j                  dv  fd        j                  }t        |dz
  |      D ]?  
t        j                   j	                  
      dkD  d j	                          d
 d       A t        j                  t              dk(  d	        t        j                  t        |      dk(  d
         j	                  d      } j	                  d       j	                  d      |dk(  r j	                  d      }nd}t        j                   j                  j                  k(  d        t        j                  j                  dk(  fd       j	                  d      }j	                  d      }j	                  d      
t        j                  ||k\  d       t        j                  ||k(  d        t        j                  
dk(  
fd       t        j                  |d   d   z   dz
  k  fd       t        j                  |d   d   z   dz
  k  fd        j                         dk(  r|||d   |d   g}	n||d   |d   g}	t        j                  |	 j                   j                        t        j                  |	t         j                   j                        fS )Nrj  c                  "    d j                    S )Nz:fractional_max_pool2d: Expected 3D or 4D tensor, but got: r  r   s   r5   rW   z,meta_fractional_max_pool2d.<locals>.<lambda>  s    LTYYKX r7   r-   r   z^fractional_max_pool2d: Expected input to have non-zero  size for non-batch dimenions, but got rZ  z emptyr   c                       y)NzNfractional_max_pool2d: kernel_size musteither be a single int or tuple of Intsr;   r;   r7   r5   rW   z,meta_fractional_max_pool2d.<locals>.<lambda>  r_   r7   c                       y)NzOfractional_max_pool2d: output_size must either be a single int or tuple of Intsr;   r;   r7   r5   rW   z,meta_fractional_max_pool2d.<locals>.<lambda>  r_   r7   r  r  r   r  r   c                       y)Nz6Expect _random_samples to have the same dtype as inputr;   r;   r7   r5   rW   z,meta_fractional_max_pool2d.<locals>.<lambda>  r_   r7   c                  "    d j                    S )Nz1Expect _random samples to have 3 dimensions got, r  )random_sampless   r5   rW   z,meta_fractional_max_pool2d.<locals>.<lambda>  s    CNDWDWCXY r7   z=Expect _random_samples.size(0) no less then input batch size.c                       y)Nz<Expect _random_samples.size(1) equals to input channel size.r;   r;   r7   r5   rW   z,meta_fractional_max_pool2d.<locals>.<lambda>  r_   r7   c                      d  dS )Nz/Expect _random_samples.size(2) equals to 2 got .r;   )r   s   r5   rW   z,meta_fractional_max_pool2d.<locals>.<lambda>  s    #RSTRUUV!W r7   c                      dd    d  S )Nz%fractional_max_pool2d: kernel height r   z' is too large relative to input height r;   )input_heightr  s   r5   rW   z,meta_fractional_max_pool2d.<locals>.<lambda>  s    7A7GGno{n|} r7   c                      dd    d  S )Nz$fractional_max_pool2d: kernel width r   z& is too large relative to input width r;   )input_widthr  s   r5   rW   z,meta_fractional_max_pool2d.<locals>.<lambda>  s    6{1~6FFlmxlyz r7   r  )rH   rY   r   r   r   r   rP   rw   r|   rt   r   )r   r  rS  rJ  r   input_channelsinput_batchr   cr   r   rO  rQ  s   `` `      @@@r5   meta_fractional_max_pool2drU    s   	LL		VX 99D4!8T" 
IIaL166:iik]BRSTRUU[]	

 
LLKA	2
 
LLKA	2 YYr]N99R=L))B-Kqyiil	LL

n***H 
LLq Y
 	AAAAAA	LL	[G 
LL	^N 
LLaWX	LLAQ'!+|;} 
LLAQ'!+{:z
 xxzQ^[^[^LAA? 	**;;	

 	++;;	
 r7   c                 |   t        j                  t        |      dv d        |d   }t        |      dk(  r|n|d   }t        |      dk(  r|n|d   }t        j                  | xs t        |      dv d        |s|n|d   }	|s|nt        |      dk(  r|	n|d   }
|s|nt        |      dk(  r|	n|d   }t        j                  t        |      dv d        |d   }t        |      dk(  r|n|d   }t        |      dk(  r|n|d   }t        j                  t        |      dv d        |d   }t        |      dk(  r|n|d   }t        |      dk(  r|n|d   }t        j                  | j                  d	v d
        | j                  dk(  r| j	                  d      nd}| j	                  d      }| j	                  d      }| j	                  d      }| j	                  d      }t        ||||	||      }t        ||||
||      }t        ||||||      }t        | |||||	|
|||||||||||||d       | j                  dk(  xr& t        j                  |       t         j                  k(  }| j                  dk(  rK| j                  d      }|j                          xr  |j                  t         j                        }||||f}n|||||f}| j                  |      }| j                  |t         j                        }|r@|j                  t         j                        }|j                  t         j                        }||fS )Nr,  c                       yNzMmax_pool3d: kernel_size must either be a single int, or a tuple of three intsr;   r;   r7   r5   rW   z.meta_max_pool3d_with_indices.<locals>.<lambda>  r_   r7   r   r   r   c                       yNzQmax_pool3d: stride must either be omitted, a single int, or a tuple of three intsr;   r;   r7   r5   rW   z.meta_max_pool3d_with_indices.<locals>.<lambda>  r_   r7   c                       yNzImax_pool3d: padding must either be a single int, or a tuple of three intsr;   r;   r7   r5   rW   z.meta_max_pool3d_with_indices.<locals>.<lambda>&  r_   r7   c                       yNzJmax_pool3d: dilation must be either a single int, or a tuple of three intsr;   r;   r7   r5   rW   z.meta_max_pool3d_with_indices.<locals>.<lambda>.  r_   r7   r4  c                       yr6  r;   r;   r7   r5   rW   z.meta_max_pool3d_with_indices.<locals>.<lambda>6  r_   r7   r  r  r  r  r   zmax_pool3d_with_indices()r  r   r   )rH   rY   r   r   r   r  r9  r@   r  r  r  rP  r   r   r  )r   r  r   r  rx  r  r:  r  r  r;  r  r  r#  r$  r%  r&  r   r!  r  r=  r>  r?  r@  rA  rB  rC  r  input_channels_last_checkr   r   r   s                                  r5   meta_max_pool3d_with_indicesrb  
  sJ    
LLKF"_ 
QB;1$+a.B;1$+a.B	LL
+c&kV+c vayBc&kQ&6F1IBc&kQ&6F1IB	LLG[ 
B7|q gajB7|q gajB	LLH\ I ]a/	Xa[I ]a/	Xa[I	LL

fK
  %zzQUZZ^AFjjnGJJrNEjjnGZZ^F BIyIE"7BB	9MG!&"b"iKF








#+2 	

aXE77>%BXBXX  zzQ$)OOA$6!)7799
'5500 6 
 	
 eWf5	WeWf=	
//)
$Cooiu{{o;Gff5#9#9f:**5+A+A*B<r7   c                    t        j                  t        |      dv d        |d   }t        |      dk(  r|n|d   }	t        |      dk(  r|n|d   }
t        j                  | xs t        |      dv d        |s|n|d   }|s|	nt        |      dk(  r|n|d   }|s|
nt        |      dk(  r|n|d   }t        j                  t        |      dv d        |d   }t        |      dk(  r|n|d   }t        |      dk(  r|n|d   }t        j                  t        |      dv d        |d   }t        |      dk(  r|n|d   }t        |      dk(  r|n|d   }t        j                  |j                  d	v d
        |j	                  d      }|j	                  d      }|j	                  d      }|j	                  d      }| j	                  d      }| j	                  d      }| j	                  d      }t        || ||||	|
|||||||||||||||d       |j                  dk(  xr& t        j                  |      t         j                  k(  }|j                  dk(  rD|j                  d      }|j                          xr  |j                  t         j                        }|j                  |j                        }|r |j                  t         j                        }|S )Nr,  c                       yrX  r;   r;   r7   r5   rW   z7meta_max_pool3d_with_indices_backward.<locals>.<lambda>  r_   r7   r   r   r   c                       yrZ  r;   r;   r7   r5   rW   z7meta_max_pool3d_with_indices_backward.<locals>.<lambda>  r_   r7   c                       yr\  r;   r;   r7   r5   rW   z7meta_max_pool3d_with_indices_backward.<locals>.<lambda>  r_   r7   c                       yr^  r;   r;   r7   r5   rW   z7meta_max_pool3d_with_indices_backward.<locals>.<lambda>  r_   r7   r4  c                       yr6  r;   r;   r7   r5   rW   z7meta_max_pool3d_with_indices_backward.<locals>.<lambda>  r_   r7   r  r  r  r   z"max_pool3d_with_indices_backward()r  r  r   )rH   rY   r   r   r   r4  r@   r  r  r  rP  r   r   r  )r  r   r  r   r  rx  r  r   r:  r  r  r;  r  r  r#  r$  r%  r&  r   r!  r=  r>  r?  r@  rA  rB  rC  r  ra  r  s                                 r5   %meta_max_pool3d_with_indices_backwardri  s  s    
LLKF"_ 
QB;1$+a.B;1$+a.B	LL
+c&kV+c vayBc&kQ&6F1IBc&kQ&6F1IB	LLG[ 
B7|q gajB7|q gajB	LLH\ I ]a/	Xa[I ]a/	Xa[I	LL

fK
 jjnGJJrNEjjnGZZ^FR Er"Gb!F#








,/6 	

aXE77>%BXBXX  zzQ$)OOA$6!)7799
'5500 6 
 	 -J]]1G1G]H
r7   gridc                 z    t        j                   j                  j                  k(   fd       t        j                   j                  t         j                  k(  xr j                  t         j                  k(   fd       t        j                   j
                  d   j
                  d   k(   fd       t        j                  j
                  d    j                  dz
  k(   fd       t        d j                        D ],  t        j                   j
                     dkD   fd       . y )	Nc                  <    dj                    d j                    S )NzNgrid_sampler(): expected input and grid to be on same device, but input is on z and grid is on r  rj  r   s   r5   rW   z+check_grid_sampler_common.<locals>.<lambda>  s'    \\N"24;;-A r7   c                  <    dj                    d j                    S )NzTgrid_sampler(): expected input and grid to have torch.strided layout, but input has z and grid has )rs   rm  s   r5   rW   z+check_grid_sampler_common.<locals>.<lambda>  s&    nT[[MC r7   r   c                  <    dj                    d j                    S )NzZgrid_sampler(): expected grid and input to have same batch size, but got input with sizes  and grid with sizes r'  rm  s   r5   rW   z+check_grid_sampler_common.<locals>.<lambda>  s'      %},A$**O r7   r   r   c                  B    dj                   dz
   d j                   S )Nz+grid_sampler(): expected grid to have size r   z, in last dimension, but got grid with sizes )r   r   rm  s   r5   rW   z+check_grid_sampler_common.<locals>.<lambda>  s,    9%**q.9I J226**? r7   c                  *    dj                    d  dS )NzYgrid_sampler(): expected input to have non-empty spatial dimensions, but input has sizes rZ  r[  r'  rn  s   r5   rW   z+check_grid_sampler_common.<locals>.<lambda>  ro  r7   )rH   rY   rt   rs   r  r   r   r   )r   rj  r   s   ``@r5   check_grid_sampler_commonrs    s    	LL#	
 
LL%F$++*F	
 
LLA$**Q-'	
 
LL

2%**q.(	
 1ejj! 
KKNQ	

r7   c                       e Zd ZdZdZdZy)GridSamplerInterpolationr   r   r   N)rm   
__module____qualname__BILINEARNEARESTBICUBICr;   r7   r5   ru  ru    s    HGGr7   ru  interpolation_modec                     t        j                   j                  dk(  xr  j                  j                  k(   fd       t        j                   j                  dk(  xr |t        j                  j
                  k(   d        y )Nr  c                  <    dj                    d j                    S )Nzdgrid_sampler(): expected 5D input and grid with same number of dimensions, but got input with sizes rp  r'  rm  s   r5   rW   z'check_grid_sampler_3d.<locals>.<lambda>  s&    449KK=#DJJ<1 r7   c                       y)Nz<grid_sampler(): bicubic interpolation only supports 4D inputr;   r;   r7   r5   rW   z'check_grid_sampler_3d.<locals>.<lambda>  r_   r7   )rH   rY   r   ru  rz  r  )r   rj  r{  s   `` r5   check_grid_sampler_3dr    sp    	LL

a3EJJ$))3	
 
LLJJ!O M"&>&F&F&L&LL	
 	Or7   c                     |d   }|r&t        j                  |t         j                        }nd }t        j                  |t         j                        }	||	fS Nr   r   )rH   rm  r   r   
r  r   rj  r{  padding_modealign_cornersr  input_requires_gradr  	grad_grids
             r5   grid_sampler_2d_backward_metar    sQ     &a.%%e5;R;RS

  U5L5LMI	""r7   c                     t        | |       t        | ||       | j                  d   }| j                  d   }|j                  d   }|j                  d   }|j                  d   }	| j                  |||||	f      S )Nr   r   r   r-   )rs  r  r   r   )
r   rj  r{  r  r  rj  Cout_Dout_Hout_Ws
             r5   grid_sampler_3dr  -  sv     eT*%'9:AAAAJJqMEJJqMEJJqME??Aq%677r7   r  c                     t        ||       t        |||       |d   }|r&t        j                  |t        j                        }nd }t        j
                  |t        j                        }	||	fS r  )rs  r  rH   rm  r  r   r  s
             r5   grid_sampler_3d_backwardr  @  sm     eT*%'9:%a.%%!?!?

 
  U5S5STIy  r7   c                     |j                  dd       }|st        j                  |      }||d<   t        j                  | g|i |S )NrP   )rO   r@   	get_dtyperH   r|   )r   r  rB   r  rP   s        r5   fullr  X  sE    JJw%E
+F7O;;t-d-f--r7   c                 N   |t         j                  k(  rt        j                  |d u d        t        j                  d|| j                  n|||| j
                  n||      }| j                  r>|j                  | j                         | j                         | j                                n/|j                  | j                         | j                         d       |j                  d       |S t        j                  j                  | |||||      }|j!                  d       |S )Nc                       y)Nz9memory format option is only supported by strided tensorsr;   r;   r7   r5   rW   zzeros_like.<locals>.<lambda>n  r_   r7   r   r   Tr  )rH   
sparse_coorY   r|   rP   rt   	is_sparsesparse_resize_and_clear_r   
sparse_dim	dense_dimrw   _coalesced_r*   r   r  fill_)r   rP   rs   rt   ru   r   r  s          r5   rm  rm  b  s     !!!T!O	

 kk %$**5"(.4;;f!
 >>((		T__.0@ ((dhhj!D

//
!
!# " C IIaLJr7   c                     ddl m}  j                         }t        j                  |dk7  d        dk\  rn|z    j                        }t        j                   | |kD        xs  ||k\          fd       dk\  rn|z   t         j                               }t         j                               } j                         |   z  z   }|= |=  j                  |||      S )Nr   r  c                       y)Nz-select() cannot be applied to a 0-dim tensor.r;   r;   r7   r5   rW   zmeta_select.<locals>.<lambda>  r_   r7   c                  6    d dj                          d  S )Nzselect(): index z! out of range for tensor of size z at dimension r   rw   r   r   s   r5   rW   zmeta_select.<locals>.<lambda>  s%    "5')J99;-~cU, r7   )
r  r  rw   rH   r   r   r   r   r   r4  )	r   rw   r   r  r   r   new_sizer%  new_storage_offsets	   ```      r5   meta_selectr    s    J88:D		?
 #sTzC99S>D	 %$/V3GQU3V	
	,	 aZEUT\EDIIK Hdkkm$J,,.C1HH3??8Z1CDDr7   c                 ,    t        j                  |       S r2   r@   clone_preserve_strides)r   r  rw   r   s       r5   meta_select_scatterr        ''--r7   c                 ,    t        j                  |       S r2   r  )r   r  rw   ro   rn   steps         r5   meta_slice_scatterr    r  r7   dim_post_exprwrap_scalarc                 v    |dk  r|sJ d}| }|dz
  }| |k  s| |kD  rJ d|  d| d| d       | dk  r| |z  } | S )Nr   r   zdim z out of bounds (rj   rk   r;   )rw   r  r  r   r  s        r5   r   r     sm    {.C
!
Cc	S3YR4u4DSEC5PQ)RR'
Qw}Jr7   c                 J    | j                         dk(  rdS | j                  |   S r  r  )r  rw   s     r5   ensure_nonempty_sizer    s!    11.!''#,.r7   c                 :    t         j                         d      }t        j                         d      }t        j                  ||k(  d        t	        |      D ];  k7  s	t        j                  t              t               k   fd       = y )Nr   c                       y)NzDIndex tensor must have the same number of dimensions as input tensorr;   r;   r7   r5   rW   z$gather_shape_check.<locals>.<lambda>  r_   r7   c                  N    d dj                    dj                    d  z   S )Nz!Size does not match at dimension z expected index  to be no larger than self  apart from dimension r'  )rw   r   r   r   s   r5   rW   z$gather_shape_check.<locals>.<lambda>  s7    ;A3>Nu{{m\/

|;QRUQVWX r7   )r  rw   rH   rY   r   r  )r   rw   r   	self_dims
index_dimsr   s   ```  @r5   gather_shape_checkr    s    DHHJ"IUYY[!$J	LLZV 9 8LL$UA.2FtQ2OOXr7   c                 2   ddl m} t        || j                               } |j	                         dk(        }|sAt        j                  j                  t
        j                  k(  fd       t        | |       | j                  j                        S )Nr   r  c                  "    d j                    S )Nz2gather(): Expected dtype int64 for index, but got r   r   s   r5   rW   zmeta_gather.<locals>.<lambda>  s    HV r7   )r  r  r   rw   r   rH   rY   rP   r   r  r   r   )r   rw   r   sparse_gradr  wrapped_dimis_index_emptys     `    r5   meta_gatherr    sp    J dhhj1K)%++-1*<=NKK5::%V	
 	4e4>>%++&&r7   c                     |r6| dk(  ry| dk(  ry| dk(  ry| dk(  ry| d	k(  ry
t        j                  dd        y | dk(  ry| dk(  ryt        j                  dd        y )Nr  
REDUCE_ADDrl  REDUCE_MULTIPLYmeanREDUCE_MEANamaxREDUCE_MAXIMUMaminREDUCE_MINIMUMFc                       y)Nz=reduce argument must be either sum, prod, mean, amax or amin.r;   r;   r7   r5   rW   z#get_operator_enum.<locals>.<lambda>  r_   r7   addmultiplyc                       y)Nz/reduce argument must be either add or multiply.r;   r;   r7   r5   rW   z#get_operator_enum.<locals>.<lambda>  r_   r7   rP  )reduce_use_new_optionss     r5   get_operator_enumr    s{    e$ ##S	
 	e
"$UUVr7   c                     ddl m}  ||j                         dk7        r4t        j                  |j
                  t        j                  k(   fd       |1t        j                  |j
                  |j
                  k(   fd       y y )Nr   r  c                        dS )Nz"(): Expected dtype int64 for indexr;   method_names   r5   rW   z,scatter_gather_dtype_check.<locals>.<lambda>  s    {m#EF r7   c                        dS )Nz0(): Expected self.dtype to be equal to src.dtyper;   r  s   r5   rW   z,scatter_gather_dtype_check.<locals>.<lambda>  s    {m#ST r7   )r  r  r   rH   rY   rP   r   )r  r   r   src_optr  s   `    r5   scatter_gather_dtype_checkr    sd    JEKKMQ./KK5::%F	

 JJ'--'T	
 r7   c                     t        | d      S r   )r  r   s    r5   ensure_nonempty_dimr    s    sA;r7   c                     ddl m}  |j                         dk(        ry t        j                  t         j                               t        j                               k(  d        d}t         j                               }t        |      D ]'  }t        |      }|k(  r|t         |      kD  s%d} n |s1/t        |      D ]!  }t        |      }|t        |      kD  sd} n ft        j                  t         j                               t        j                               k(  d        t        j                  |  fd       y t        j                  |  fd       y )	Nr   r  c                       yNzCIndex tensor must have the same number of dimensions as self tensorr;   r;   r7   r5   rW   z%scatter_shape_check.<locals>.<lambda>#  r_   r7   FTc                       yr  r;   r;   r7   r5   rW   z%scatter_shape_check.<locals>.<lambda>=  r_   r7   c                  b    dj                    dj                    d  dj                    z   S )NExpected index r  r  z and to be no larger than src r'  )rw   r   r   r  s   r5   rW   z%scatter_shape_check.<locals>.<lambda>A  s8    oekk]2Mdjj\Z&se+I'--YZ r7   c                  H    dj                    dj                    d  z   S )Nr  r  r  r'  r  s   r5   rW   z%scatter_shape_check.<locals>.<lambda>G  s,    oekk]2Mdjj\Z&se,- r7   )	r  r  r   rH   rY   r  rw   r   r  )	r   rw   r   r  r  is_wrong_shaper  r   index_d_sizes	   ````     r5   scatter_shape_checkr    sP   JEKKMQ./	LLDHHJ'+>uyy{+KKU
 N#DHHJ/I 9 +E158.tQ77!N g1y! 	A/q9L27A>>!%		 
+/B599;/OOY	
 	Z	
 	-	
r7   c                     t        || j                               }t        d| ||       t        | |||       |t	        ||       y y )Nscatter)r   rw   r  r  r  )r   rw   r   r  r  r  r  s          r5   scatter_meta_implr  M  sE     dhhj1Ky$s;k5#6'?3 r7   c                 V    t        | |||d       | j                  | j                        S Nr  r  r   r   r   rw   r   r  s       r5   meta_scatter_addr  V  s%    dCU3>>$**%%r7   c                 $    t        | |||d       | S r  r  r  s       r5   meta_scatter_add_r  \  s    dCU3Kr7   c                     t        |t        j                        r|nd }t        | ||||       | j	                  | j
                        S r2   )rb   rH   r   r  r   r   r   rw   r   src_or_valuerY  r  s         r5   meta_scatterr  b  s;     %\5<<@,dCdCV4>>$**%%r7   c                 `    t        |t        j                        r|nd }t        | ||||       | S r2   )rb   rH   r   r  r  s         r5   meta_scatter_r  q  s-     %\5<<@,dCdCV4Kr7   queryr   r  	dropout_p	is_causalreturn_debug_maskrZ  c           	      V   | j                  d      }| j                  d      }| j                  d      }	| j                  d      }
|j                  d      }| j                  dd      }t        j                  |      j                  dd      }t        j                  |||	ft        j
                  | j                        }|ra|
dkD  rdnd}t        j                  |	|z        }|dk  rd}n|dk  rd}t        j                  |||	|f| j                  | j                        }n,t        j                  d| j                  | j                        }t        j                  j                  rkt        j                  j                         rMt        j                  d	t        j                  d
      }t        j                  d	t        j                  d
      }nLt        j                  dt        j                  d
      }t        j                  d	t        j                  d
      }||d d |	||||f	S )Nr   r   r   r-   r  @         r;   rr   )r   r  rH   r   r|   rL   rt   rk  ceilrP   versionhipr   r@  r   r  )r  r   r  r  r  r  rZ  r   	num_headsmax_seqlen_batch_qhead_dimmax_seqlen_batch_kquery_t	attention	logsumexpblocksize_cmax_seqlen_k
debug_maskseedoffsets                       r5   (meta__scaled_dot_product_flash_attentionr	    s    AJ

1IAzz!}H!ooa#G  )33Aq9I	Y 23kk||I %]cyy!3k!AB$L3&L[[$6E++<<

 [[%++ellK
 }}UZZ446{{2UZZ?Ruzz&A{{Aell6BRu||FC 	
 
r7   	attn_biascompute_log_sumexpc	           	          | j                  d      }	| j                  d      }
| j                  d      }|j                  d      }|j                  d      }t        j                  |	|
||f| j                  | j                        }t        j                  |	|
|ft        j
                  | j                        }t        j                  dt        j                  d      }t        j                  dt        j                  d      }||d d ||||d f	S Nr   r   r   r   r  r;   rr   r   rH   r|   rP   rt   rL   r   )r  r   r  r
  r  r  r  r  rZ  r  r  S_QS_KVD_Vr  
logsum_expr  r  s                     r5   (meta__scaled_dot_product_cudnn_attentionr    s     	

1A

1A
**Q-C88A;D
**R.C
++q!S#&ekk%,,
OC	
Askk||J ;;rF;D[[5::f=F 	
 
r7   c           	          | j                  d      }| j                  d      }	| j                  d      }
|j                  d      }|j                  d      }t        j                  ||	|
|f| j                  | j                        }t        j                  ||	|
ft        j
                  | j                        }t        j                  dt        j                  d      }t        j                  dt        j                  d      }||d d |
|||d f	S r  r  )r  r   r  r
  r  r  r  rZ  r  r  r  r  r  r  r  r  r  s                    r5   5meta__scaled_dot_product_fused_attention_overrideabler    s     	

1A

1A
**Q-C88A;D
**R.C
++q!S#&ekk%,,
OC	
Askk||J ;;rF;D[[5::f=F 	
 
r7   r\  r  	cum_seq_q	cum_seq_kmax_qmax_kphilox_seedphilox_offsetc                 J   t        j                  |j                  dd            j                  dd      }t        j                  |j                  dd            j                  dd      }t        j                  |j                  dd            j                  dd      }|||fS r^  )rH   r   r  )r\  r  r   r  r   r  r  r  r  r  r  r  r  r  rZ  grad_qgrad_kgrad_vs                     r5   'meta__scaled_dot_product_flash_backwardr     s    , eooa34>>q!DFcmmAq12<<QBFeooa34>>q!DF66!!r7   	attn_maskc                     | j                  d      }| j                  d      }| j                  d      }	t        j                  |       }
t        j                  ||	|ft        j                  | j
                        j                  dd      }|
|fS )Nr   r   r   r  )r   rH   r   r|   rL   rt   r  )r  r   r  r  r  r!  rZ  r   r  r  r  r  s               r5   0meta__scaled_dot_product_flash_attention_for_cpur#  /  s     AJ

1IA  'I	

 kk|| i1o  	 r7   c
                    |j                  d      }
|j                  d      }|j                  d      }|j                  d      }|j                  d      }t        j                  |
|||fd|j                  |j                        }t        j                  |
|||fd|j                  |j                        }t        j                  |
|||fd|j                  |j                        }|||fS )Nr   r   r-   r   r   r   r   r-   r  )r   rH   empty_permutedrP   rt   )r\  r  r   r  r   r  r  r  r!  rZ  r   r  r  len_qlen_kr  r  r  s                     r5   9meta__scaled_dot_product_flash_attention_for_cpu_backwardr)  Q  s    & AJ

1Izz!}HJJqMEHHQKE!!	Yx0kk||	F !!	Yx0iizz	F !!	Yx0kk||	F 66!!r7   c                 @   | j                  dd      } |j                  dd      }|j                  dd      }| j                  d      }| j                  d      }	| j                  d      }
|j                  d      }t        j                  ||	|
|| j                  | j
                        }t        j                  j                  r&t        j                  j                         r	 |r|	nd}n|rt        j                  |	dz        dz  nd}t        j                  ||
|ft        j                  | j
                        }|j                  dd      }t        j                  dt        j                  d	      }t        j                  dt        j                  d	      }||||fS )
Nr   r   r   r  r   r  r  r;   rr   )r  r   rH   r|   rP   rt   r  r  r   r@  rk  r  rL   r   )r  r   r  r
  r  r  r  rZ  r  r  r  Kvr  logsumexp_dimr  r  r  s                    r5   ,meta__scaled_dot_product_efficient_attentionr-    sD    OOAq!E
--1
COOAq!E

1A

1A

2I	BB
++aIrU\\
RC}}UZZ446	 0Q2D		!b&)B.!	
I}%kk||J --1
C ;;rF;D[[5::f=F
D&((r7   grad_input_maskc                    |j                  d      }|j                  d      }|j                  d      }|j                  d      }|j                  d      }|j                  d      }t        j                  ||||fd|j                  |j                        }t        j                  ||||fd|j                  |j                        }t        j                  ||||fd|j                  |j                        }d }|~|
d   ry|j                  d      }|dz  dk(  r|n
|dz   |dz  z
  }t        |j                               }||d<   t        j                  ||j                  |j                        }|d	d |f   }||||fS )
Nr   r   r   r-   r%  r  r   rK  .)r   rH   r&  rP   rt   r   r|   )r\  r  r   r  r
  r   r  r  r  r  r.  r  rZ  r   r  r  r  
head_dim_vr  r  r  r  	grad_biaslastDimlastDimAligned	new_sizess                             r5   +meta__scaled_dot_product_efficient_backwardr5    s{   ( AJ

1IJJqMEzz!}HAJHHQKE!!	Yx0kk||	F !!	Yx0iizz	F !!	Yz2kk||	F I!3..$$+bLA$57R<'TV,;V)*	&	"KKY__Y5E5E
	 c8G8m,	669,,r7   c                     t        j                  |      }t        j                  |      }t        j                  |      }|||fS r2   rP  )r\  r  r   r  r   r  r  r  r
  r  r  r  r  r  r  rZ  r  r  r  s                      r5   'meta__scaled_dot_product_cudnn_backwardr7    sA    . e$Fc"Fe$F66!!r7   window_size_leftwindow_size_right	seqused_kalibi_slopesc                    || j                  d      n|j                         dz
  }|| j                  d      n|}||j                  d      n|}| j                  d      }| j                  d      }t        j                  |       }|4t        j                  |||ft        j
                  | j                        }nC| j                  d      }t        j                  ||ft        j
                  | j                        }|	ra|dkD  rdnd}t        j                  ||z        }|dk  rd}n|dk  rd}t        j                  ||||f| j                  | j                        }n,t        j                  d| j                  | j                        }d	\  }}t        j                  j                  rkt        j                  j                         rMt        j                  d
t        j                  d      }t        j                  d
t        j                  d      }nLt        j                  dt        j                  d      }t        j                  d
t        j                  d      }|||||fS )Nr   r   r  r   r  r  r  r  NNr;   rr   r   )r   r   rH   r   r|   rL   rt   rk  r  rP   r  r  r   r@  r   r  )r  r   r  r  r  r  r  r  r  r  rZ  r8  r9  r:  r;  r   r  r   r  r  r  r  total_qr  r  r  r  r  s                               r5   meta__flash_attention_forwardr?    s   4 #,"3A9JQ9NJ*3*;A(1(9!u

2Izz"~H   'IKK$67++<<
	 **Q-KK ELL
	 %]cyy!3k!AB$L3&L[[$6E++<<

 [[%++ellK
 LD&}}UZZ446{{2UZZ?Ruzz&A{{Aell6BRu||FC r7   c                     t        j                  |      }t        j                  |      }t        j                  |      }|||fS r2   rP  )r\  r  r   r  r   r  r  r  r  r  r  r  r  r  rZ  r8  r9  
grad_querygrad_key
grad_values                       r5   meta__flash_attention_backwardrD  V  sA    0 !!%(J$H!!%(Jx++r7   cu_seqlens_qcu_seqlens_kmax_seqlen_qr  custom_mask_typecausal_diagonalseqlen_kwindow_sizec                    | j                  d      }| j                  d      }|j                  d      }| j                  d      }|j                  d      }t        j                  ||||| j                  | j                        }||j                  d      dz
  n|}|}||J |}||n|}|
rt        j                  |dz        dz  nd}t        j                  |||ft        j                  | j                        }t        j                  dt        j                  d      }t        j                  dt        j                  d      }||||||fS )	Nr   r   r  r   r  r  r;   rr   )	r   rH   r|   rP   rt   rk  r  rL   r   )r  r   r  r-  rE  rF  rG  r  r  rH  r  rZ  rI  rJ  rK  r  r  rj  r  r+  r  logsumexp_batch_dimactual_max_seqlen_qactual_max_seqlen_kr,  r  r  r  s                               r5   !meta__efficient_attention_forwardrP  u  s9   , 	

1A

1AA

2I	BB
++aIrU\\
RC7C7O,++A.2VW'''**6*B,4F		%*+b0A  	i7kk||J ;;rF;D[[5::f=F
D&*=?RRRr7   bias_requires_gradnum_splits_keyshared_storage_dqdkdvc                    |rt        j                  |j                  d   |j                  d   k(  d        t        j                  |j                  d   |j                  d   k(  d        t        j                  g |j                  dd d|j                  d   |j                  d   |j                  |j
                        }|j                  d	d      }|j                  d	d      }|j                  d	d
      }n?t        j                  |      }t        j                  |      }t        j                  |      }|z|j                  d      }|dz  dk(  r|n
|dz   |dz  z
  }t        |j                               }||d<   t        j                  ||j                  |j
                        }|dd |f   }n!t        j                  d|j
                        }||||fS )Nr   c                       y)Nz,seqlen must match for `shared_storage_dqdkdvr;   r;   r7   r5   rW   z4meta__efficient_attention_backward.<locals>.<lambda>  r_   r7   r-   c                       y)Nz3embedding dim must match for `shared_storage_dqdkdvr;   r;   r7   r5   rW   z4meta__efficient_attention_backward.<locals>.<lambda>  r_   r7   r   r  r   r  r  r   rK  .r;   r  )
rH   rY   r   r|   rP   rt   r  r   r   r   )r\  r  r   r  r-  rE  rF  rG  r  r  r  r  r  rH  rQ  rZ  rR  rS  chunkrA  rB  rC  r2  r3  r4  r1  s                             r5   "meta__efficient_attention_backwardrX    s   2 KKNciil*B	
 	KKNciil*I	
 Eekk!BEEEKKOEU[[_E++<<

 \\"a(
<<A&\\"a(
%%e,
##C(%%e,
))B-$+bLA$57R<'TV,;V%	&	"KK	DKKP	c8G8m,	KK5<<8	xY66r7   scale_ascale_bscale_resultuse_fast_accumc                     d }t        j                   j                         dk(  xr j                         dk(   fd       t        j                   | j                        xr  |j                         fd       t	               dk(  rhd }	d }
d }t        j                   |	 j                               xs  |        fd	       t        j                   |
j                               xs  |      fd
       t        j                   j                  d      dz  dk(   fd       t        j                  j                  d      dz  dk(  xr j                  d      dz  dk(  fd       t        j                  j                  t         j                  k(  xr j                  t         j                  k(  d         j                  \  }j                  d      j                         dk(  rj                         dk(  rnt        j                  j                         dk(  xr j                         dk(  fd       j                  d      k(  rtj                  d      dk(  r`j                  d      dk(  rLj                  d      k(  r8t        j                  j                         xr j                         d        nt        j                  dfd       ||n j                  }t        j                   j                  d      j                  d      | j                        S )Nc                     | t         j                  t         j                  t         j                  t         j                  fv S r2   )rH   rO  float8_e5m2float8_e4m3fnuzfloat8_e5m2fnuzr   s    r5   is_fp8_typez#meta_scaled_mm.<locals>.is_fp8_type  s8    !!!!	
 
 	
r7   r   c                  L    dj                          d j                          S )Nz%Inputs must be 2D but got self.dim()=z and mat2.dim()=r   r:  r   s   r5   rW   z meta_scaled_mm.<locals>.<lambda>  s'    7
|CSTXT\T\T^S_` r7   c                  <    dj                    d j                    S )Nz8Expected both inputs to be fp8 types but got self.dtype=z and mat2.dtype=r   rd  s   r5   rW   z meta_scaled_mm.<locals>.<lambda>  s"    J4::,Vfgkgqgqfrs r7   r   c                 ,    | d   | d   kD  xr | d   dk(  S r  r;   r  s    r5   is_row_majorz$meta_scaled_mm.<locals>.is_row_major  s"    !9vay(;VAY!^;r7   c                 &    | d   dk(  xr | d   dkD  S r  r;   r  s    r5   is_col_majorz$meta_scaled_mm.<locals>.is_col_major  s    !9>3fQi!m3r7   c                 V    | j                  d      dk(  xs | j                  d      dk(  S r  r   )	tensor_2ds    r5   has_zero_dimz$meta_scaled_mm.<locals>.has_zero_dim	  s)    >>!$)CY^^A->!-CCr7   c                  *    d j                          S )Nz#self must be row_major, got stride r  r   s   r5   rW   z meta_scaled_mm.<locals>.<lambda>      9$++-I r7   c                  *    d j                          S )Nz#mat2 must be col_major, got stride r  r:  s   r5   rW   z meta_scaled_mm.<locals>.<lambda>  rn  r7   r   rK  r   c                  ,    d j                  d       S )NzBExpected self.size(1) to be divisible by 16, but got self.size(1)=r   r   r   s   r5   rW   z meta_scaled_mm.<locals>.<lambda>  s    XY]YbYbcdYeXfg r7   c                  "    d j                    S )Nz>Expected both dimensions of mat2 to be divisble by 16 but got r'  rp  s   r5   rW   z meta_scaled_mm.<locals>.<lambda>  s    TUYU_U_T`a r7   c                       y)Nz6Both scale_a and scale_b must be float (fp32) tensors.r;   r;   r7   r5   rW   z meta_scaled_mm.<locals>.<lambda>   r_   r7   c                  L    d j                         dj                         S )NzLFor non-tensorwise scaling, scale tensors must be 2D, but got scale_a.dim()=z and scale_b.dim()=r   )rY  rZ  s   r5   rW   z meta_scaled_mm.<locals>.<lambda>+  s,    gY`YdYdYfXhh|nunynyn{m}~ r7   c                       y)Nz@Both scale_a and scale_b must be contiguous for rowwise scaling.r;   r;   r7   r5   rW   z meta_scaled_mm.<locals>.<lambda>7  r_   r7   Fc                      d  d dj                  d       dj                  d       dj                  d       dj                  d       dS )	Nz}Invalid scaling configuration. For tensorwise scaling, both scales should be scalar. For rowwise scaling, scale_a should be (z, 1), scale_b should be (1, z). Got scale_a.size()=(r   rj   r   z) and scale_b.size()=(rk   r   )rT  r   rY  rZ  s   r5   rW   z meta_scaled_mm.<locals>.<lambda>=  sk    CCD#Eabcad e//6||A.?r',,q/AR S//6||A.?r',,q/ARRS	U r7   r  )rH   rY   rw   rP   r   r   r   rL  r   r   rP  r|   rt   )r   r:  rY  rZ  r-  r[  r/  r\  rb  rg  ri  rl  _k
_out_dtyperT  r   s   ````          @@r5   meta_scaled_mmry    s   
 
LL
a+DHHJ!O` 
LLDJJ;K

$;s
 4F"	<	4	D 	'=<+=I	
 	'=<+=I	
 	IIaL2"g	
 	IIaL2"=tyy|b'8A'=a	
 	MMU]]*Mw}}/ML	
 

2IIaL==?aGMMOq$8 LL"9w{{}'9~ Q1$LLOq(LLOq(LLOq( ))+G0E0E0G^ 	 (3J;;tyy|TYYq\DKKXXr7   c                 Z    t        | ||||d       | j                  | j                        S NT)r  r  r   rw   r   r  rY  rW  s         r5   meta_scatter_reduce_twor}  J  s)     dCVTJ>>$**%%r7   c                 (    t        | ||||d       | S r{  r  r|  s         r5   meta_scatter_reduce__twor  Q  s    dCVTJKr7   c                t    t        j                  d j                         cxk  xr dk  nc  fd        j                         dk(  r0t        j                  |t         j                   j
                        S t        j                   j                  d      |t         j                   j
                        S )Nr   r   c                  *    d j                          S )Nz@The probabilty distributions dimensions must be 1 or 2, but got r   r	  s   r5   rW   z"meta_multinomial.<locals>.<lambda>\  s    RSXS\S\S^R_` r7   r   r  )rH   rY   rw   r|   r   rt   r   )r   num_samplesreplacementr   s   `   r5   meta_multinomialr  W  s|     
LL	EIIK1` yy{a{{;ejjNN;;

1{%**U\\ r7   c                 "    d}| D ]  }||z  }	 |S r   r;   )vsrN  vs      r5   multiply_integersr  e  s$    	A 	QHr7   c                 L    t        j                  t              k(  fd       dz   t        j                  t               k(   fd       t        j                  t        d  dd  D              xr t        d D               fd        d d \  }}||gS )Nc                  &    d  dt               S )Nz%It is expected output_size equals to , but got size r  )num_spatial_dimsrS  s   r5   rW   z'upsample_common_check.<locals>.<lambda>o  s    78H7IY\]hYiXjk r7   r   c                  &    d  dt               S )Nz$It is expected input_size equals to r  r  )expected_input_dimsr)  s   r5   rW   z'upsample_common_check.<locals>.<lambda>t  s    67J6K?[^_i[jZkl r7   c              3   &   K   | ]	  }|d kD    ywr  r;   rd   r  s     r5   rf   z(upsample_common_check.<locals>.<genexpr>x  s     *aAE*r  c              3   &   K   | ]	  }|d kD    ywr  r;   r  s     r5   rf   z(upsample_common_check.<locals>.<genexpr>x  s     2NQ1q52Nr  c                      d  d S )NzDInput and output sizes should be greater than 0, but got input size z and output size r;   )r)  rS  s   r5   rW   z'upsample_common_check.<locals>.<lambda>y  s      \!2;-A r7   )rH   rY   r   r  )r)  rS  r  r  channelsr  s   ```  @r5   upsample_common_checkr  l  s    	LLK,,k +Q.	LLJ..l
 
LL*:ab>**Ns2N+2N/N	A ""1~FHH+{++r7   c                 4    t        j                   j                         dk7  xs t         j	                         dd         fd       t         j	                         |d      } j                  |      j                  t        j                               S )Nr   r   c                  *    d j                          S )Nz>Non-empty 3D data tensor expected but got a tensor with sizes r   r	  s   r5   rW   z$upsample_nearest1d.<locals>.<lambda>      PQVQ[Q[Q]P^_ r7   r  r   
rH   rY   r   r  r   r  r   r  r@   r  )r   rS  scalesfull_output_sizes   `   r5   upsample_nearest1dr         
LLA/

QR0@A_ -

kA ??+,//11%8 0  r7   c                     t        j                   j                         dk7  xs t         j	                         dd         fd       t         j	                         |d      } j                  |      }t        j                         } j                  \  }}}} j                  j                  dk(  r|dk  rt         j                  }|j                  |      }|S )	Nr   r   c                  *    d j                          S Nz>Non-empty 4D data tensor expected but got a tensor with sizes r   r	  s   r5   rW   z$upsample_nearest2d.<locals>.<lambda>  r  r7   r   r  r   r  r   )rH   rY   r   r  r   r  r   r@   r  r   rt   rl   r   
contiguous)	r   rS  scales_hscales_wr  r   r   rC   
n_channelss	   `        r5   upsample_nearest2dr    s     
LLA/

QR0@A_ -

kA __-.F //6M  ++Az1a||F"zA~//];FMr7   rS  r)  r  r  c                 X    t        ||d      t        j                   j                  dk(   fd       t	        d      D ]2  t        j                   j                           k(   fd       4  j                  |      j                  t        j                               S )Nr   r  r  c                  "    d j                    S NzFExpected grad_output to be a tensor of dimension 4 but got: dimension r  ry  s   r5   rW   z-upsample_nearest2d_backward.<locals>.<lambda>      XYdYiYiXjk r7   c            
      D    d d     d dj                         S )NzCExpected grad_output to have the same shape as output; output.size() = z but got grad_output.size(r   r  r  r   s   r5   rW   z-upsample_nearest2d_backward.<locals>.<lambda>  s>      !s$'7':&;,QCtK4D4DQ4G3HJ r7   r   )
r  rH   rY   r   r   r   r   r  r@   r  )r  rS  r)  r  r  r  r   s   `    @@r5   upsample_nearest2d_backwardr    s     -K! 
LLAk 1X 
Q#3A#66	

   ,//11+> 0  r7   c                 4    t        j                   j                         dk7  xs t         j	                         dd         fd       t         j	                         |d      } j                  |      j                  t        j                               S )Nr   r   c                  *    d j                          S )Nz>Non-empty 5D data tensor expected but got a tensor with sizes r   r	  s   r5   rW   z$upsample_nearest3d.<locals>.<lambda>  r  r7   r-   r  r   r  )r   rS  scales_dr  r  r  s   `     r5   upsample_nearest3dr    r  r7   c                    t        j                  |       t        j                  | t         j                        }}||t        |t              sJ t        |t              sJ |j
                  }|j                         }	t        ||      }t        ||      }|j                  ||	       |j                  ||	       t        ||       t        ||       ||fS ||fS )Nr   )r}  r~  )
rH   r   r   rb   r   r   r   r    r   r"   )
r   stablerw   
descendingr   r   r  r   r   
out_strides
             r5   	meta_sortr    s     D!5#3#3D#LqAg1&*---':... GG	XXZ
"695#GY79j1Iz2F3G4wa4Kr7   c           	          t        j                   j                  dk(   fd       t        j                   j                  j                  k(   fd        j	                  d      t        j                  j                  dk(  fd       t        j                  j                         k(  fd       t        j                  j                  j                  k(  fd       t        j                  j                  dk(  fd        j	                  d	      z  z  t        j                  j                         k(   fd
       t        j                  t         fdfD              d        y )Nr   c                  "     j                    dS Nz != 2r  )input_gatess   r5   rW   z%rnn_cell_checkSizes.<locals>.<lambda>      ;3C3C2DE0J r7   c                  :    j                    d j                    S N != r'  )hidden_gatesr  s   r5   rW   z%rnn_cell_checkSizes.<locals>.<lambda>  s     ;$$%T,*<*<)=> r7   r   c                  "     j                    dS )Nz != 1r  )
input_biass   r5   rW   z%rnn_cell_checkSizes.<locals>.<lambda>  s    joo5Fe3L r7   c                  .    j                          d  S r  r~  )
gates_sizer  s   r5   rW   z%rnn_cell_checkSizes.<locals>.<lambda>  s    z'')*$zl; r7   c                  :    j                    d j                    S r  r'  )hidden_biasr  s   r5   rW   z%rnn_cell_checkSizes.<locals>.<lambda>  s     z''([->->,?@ r7   c                  "     j                    dS r  r  )prev_hiddens   r5   rW   z%rnn_cell_checkSizes.<locals>.<lambda>  r  r7   r   c            
      `    j                          dj                  d       d d d  d
S )Nr  r   z * z // z (aka rk   )r   r   )expected_prev_hidden_numelfactorr  r  r  s   r5   rW   z%rnn_cell_checkSizes.<locals>.<lambda>  sB    ;$$&'tK,<,<Q,?+@J<tTZS[[ab|a}}~ r7   c              3   P   K   | ]  }|j                   j                   k(    y wr2   r  )rd   rE   r  s     r5   rf   z&rnn_cell_checkSizes.<locals>.<genexpr>  s(      
 HH***
s   #&c                       y)Nz%expected all inputs to be same devicer;   r;   r7   r5   rW   z%rnn_cell_checkSizes.<locals>.<lambda>  r_   r7   )rH   rY   r   r   r   r   r  )r  r  r  r  r  r  r  r  s   ``````@@r5   rnn_cell_checkSizesr    s@    
LL!!Q&(JK	LL\///> !!!$JZ__)+LM*,;	
 	 1 11@	
 
LL!!Q&(JK!,!1!1!!4z!AV!K	LL99 
LL 
"J[I
 	
 	8r7   c                 
   t        | |||d|       t        j                  | t        j                        }t        j                  |t        j                        }t        j                  |t        j                        }|||fS )Nr  r   )r  rH   r   r   )r  r  cxr  r  	workspacehycys           r5   _thnn_fused_lstm_cell_metar     sk     \:{ArR  E<S<STI			"E,C,C	DB			"E,C,C	DBIr7   c                 b   t        |      dk7  }|r t        |      }|d   }| j                  d   }nB|
r| j                  d   n| j                  d   }|
r| j                  d   n| j                  d   }d}|rdnd}|dk7  r|n|}|r|||z  g}n|
r||||z  gn||||z  g}| j                  |      }|	|z  ||g}|"t        j                  d| j
                        }n|j                  |      }|j                  |	|z  ||g      }|rdnd}| j                  |t        j                        }|||||fS )Nr   r   r   r   r  r   )r   r   r   rH   r|   rt   r  )r   r+  weight_stride0
weight_bufhxr  rJ  hidden_size	proj_size
num_layersbatch_firstdropouttrainbidirectionalbatch_sizesdropout_stateis_input_packed
seq_length
mini_batchbatch_sizes_sumnum_directionsout_sizer   r   
cell_shaper  r  reserve_shapereserves                                r5   
_cudnn_rnnr  /  sS   & +&!+O%
 ^
++a.'2U[[^A
'2U[[^A
'QQN%NyH$h&?@	  X%>?j(^*CD 	
 __Y'F~-z;GJ	z[[5<<0\\*%	zN2JI	JB AAMoom5;;o?G2r7J..r7   c                 (   |r| j                   d   n| j                   d   }|r| j                   d   n| j                   d   }|
}|r|||gn|||g}| j                  |      }|"t        j                  d| j                        }n|j                  |j                         }|"t        j                  d| j                        }n|j                  |j                         }t        j                  d| j                  t        j
                        }||||fS )Nr   r   r  r   )r   r   rH   r|   rt   r  )r   w0w1w2w3hx_cx_r   r  rJ  r  r  
has_biasesr  r  r  r  r  output_chanelsr   r   r  r  r  s                           r5   mkldnn_rnn_layerr  g  s    & $/QEKKNJ#.QEKKNJ N  
Z0*n5 
 __Y'F
{[[5<<0]]399%
{[[5<<0]]399%Aell%++FI2r9$$r7   c                     | j                   dk(  r%t        j                  dk(  xs dk(  fd       y t        j                  | j                        dk7  fd       y )Nr   r   c                       d  S )Nz4: Expected reduction dim -1 or 0 for scalar but got r;   rw   r  s   r5   rW   z'zero_numel_check_dims.<locals>.<lambda>  s    wiSTWSXY r7   c                       d  dS )Nz: Expected reduction dim z to have non-zero size.r;   r  s   r5   rW   z'zero_numel_check_dims.<locals>.<lambda>  s    wi8=TU r7   )r   rH   r   r   )r   rw   r  s    ``r5   zero_numel_check_dimsr    sR    yyA~1H!r	Y	

 	IIcNaU	
r7   c                      |(t        ||j                               }t        ||        y t        j                  |j                         dk7   fd       y )Nr   c                        dS )Nz@: Expected reduction dim to be specified for input.numel() == 0.r;   r  s   r5   rW   z%check_argmax_argmin.<locals>.<lambda>  s    tf\] r7   )r   rw   r  rH   rY   r   )r  r   rw   s   `  r5   check_argmax_argminr    sC    
S$((*-dC.JJLA]	
r7   c                     t        d| |       t        j                  | j                  ||fnd       }t	        | ||      }| j                  |t        j                        S )Nargmaxr   )r  r@   ry  r   rz  r   rH   r   )r   rw   r|  rp  r   s        r5   argmax_argmin_metar    sQ    $,

coSF4PD$T49E>>%u{{>33r7   c                 |    |t         j                  k(  rt         j                  }t        j                  d||||      S )Nr;   r   )rH   jaggedr  r|   )r  rP   rs   rt   ru   s        r5   scalar_tensorr    s5    
 ;;
%v* r7   c                    t        || j                         d      }| j                         dk(  rdn| j                  |      }t        j                  |       t        j
                  ||k  d        t        | j                        }t        |      dkD  r|||<   | j                  |      | j                  |t        j                        fS )NTr  r   r   c                       y)Nzk not in range for dimensionr;   r;   r7   r5   rW   ztopk_meta.<locals>.<lambda>  r_   r7   r   )r   rw   r   rH   r  rY   r   r   r   r   r   )r   rS  rw   largestsorted	sliceSizetopKSizes          r5   	topk_metar
    s     dhhjd
;CXXZ1_$))C.I		LLi!GHDJJH
8}q>>(#T^^HEKK^%PPPr7   c                     |	|J d       |j                         }| j                         }	t        j                  ||	j                  |	j                  |	j
                        S )Nz;segment_reduce(): Either lengths or offsets must be defined)rP   rt   rs   )r  rH   r   rP   rt   rs   )
r,  r   ri  rY  rd  re  rf  rh  data_contiggrad_contigs
             r5   meta__segment_reduce_backwardr    sj    
 '"5 E5 //#K//#K!!!!	 r7   c                    t        | j                         d      | j                         dkD  r| j                        nd}t        j                  |dk\  xr ||k  fd       t        | j                  d  | j                  dz   d  z         }|r%| j                         dkD  r|j                  d       | j                  |      | j                  |t        j                        fS )NTr  r   r   c                      d  S )Nz9kthvalue(): selected number k out of range for dimension r;   r   s   r5   rW   zkthvalue_meta.<locals>.<lambda>  s    KC5Q r7   r   )
r   rw   r   rH   rY   r   r   r!  r   r   )r   rS  rw   r|  dimSizer   s     `   r5   kthvalue_metar    s     dhhjd
;C $
QdiinAG	LL	Q1<Q
 DS!DJJsQwy$99:E488:>S!>>% $..ekk."JJJr7   c                    | | n|}t        j                  |j                         dk(  d        |j                         }| (t        j                  | j                         |k(  d        |(t        j                  |j                         |k(  d        t        j                  |j                         |k(  d        t        j                  |j                         |k(  d        t        j                  |j                         dk(  d        t        j                  |j	                         |d   |d	   z  d
z  k(  d        y )Nr   c                       yN r;   r;   r7   r5   rW   z(checkLSTMBackwardSizes.<locals>.<lambda>  r_   r7   c                       yr  r;   r;   r7   r5   rW   z(checkLSTMBackwardSizes.<locals>.<lambda>  r_   r7   c                       yr  r;   r;   r7   r5   rW   z(checkLSTMBackwardSizes.<locals>.<lambda>  r_   r7   c                       yr  r;   r;   r7   r5   rW   z(checkLSTMBackwardSizes.<locals>.<lambda>  r_   r7   c                       yr  r;   r;   r7   r5   rW   z(checkLSTMBackwardSizes.<locals>.<lambda>  r_   r7   c                       yr  r;   r;   r7   r5   rW   z(checkLSTMBackwardSizes.<locals>.<lambda>  r_   r7   r   r   r  c                       yr  r;   r;   r7   r5   rW   z(checkLSTMBackwardSizes.<locals>.<lambda>  r_   r7   )rH   rY   rw   r   r   )grad_hygrad_cyr  r  r  defined_gradexp_sizes          r5   checkLSTMBackwardSizesr!    s    %17wL	LL!!#q(*5  "HW\\^x/<W\\^x/<	LLh&
3	LLh&
3	LLA%z2	LL"hqkHQK&?!&CCZPr7   c                     | |yt        | ||||       t        j                  |t              }t        j                  |t              }|r|j	                  dd      nd }|||fS )NNNNr   r   F)r|  )r!  rH   r   legacy_contiguous_memory_formatr  )	r  r  r  r  r  has_bias
grad_gatesgrad_cxr1  s	            r5   #_thnn_fused_lstm_cell_backward_implr(    sl    7?7GRY?!!!@J r1PQG4<
q%0$Iw	))r7   c                    d }d }d }|d   r|j                  | j                               }|d   s|d   rQ|j                  |j                  d      | j                  d      f      }|j                  |j                  d            }|||fS )Nr   r   r   r   r  )r  r  r  r  r  grad_weightr1  s          r5   linear_backwardr+    s    JKI1~!++FKKM:
1~Q",,l.?.?.CV[[QS_-UV **<+<+<R+@A	Y//r7   c                     t         j                        dkD  r j                  d   ||z  z  dk(  sJ d j                   d|        d  fd} j                  d   ||z  z  } j                  d   |z  } j                  d	   |z  }g  j                  d d |||} j                  |      }|j                   |       
      }|S )Nr   r  r   z'Invalid input shape for pixel_shuffle: z with upscale_factor = c                 b    t         j                  j                  |       t         j                  k(  S r2   r  r  s    r5   r  z,meta_pixel_shuffle.<locals>.is_channels_last"  s$    ""88=ATATTTr7   c                  2           r.t              dk(  rt        j                  S t        j                  S j	                  t        j                        rt        j                  S j	                  t        j
                        rt        j
                  S y r  )r   rH   r   r  rP  r  )r  r   s   r5   r  z.meta_pixel_shuffle.<locals>.pick_memory_format%  s|    D!4 F*...***e.E.EF***e.C.CD((( Er7   r  r   r   )r   r   r   r  )	r   upscale_factorr  r  HrWrr   r   r  s	   `       @r5   meta_pixel_shuffler2    s     	DJJ!

2.>2Q RVW W 2$**=TUcTdeW
U	) 	

2>N:;A	B.	(B	B.	(B-$**Sb/-1-b-"-I
..
#C
&&13&
4CJr7   c                 X   | j                  | j                        }|j                  |j                        }|j                  |j                        }|j                  |j                        }|j                  |j                        }|j                  |j                        }|||||||fS r2   r  )r   weight0weight1weight2weight3r  cx_tmpr   hy_cy_grad_output_r_optgrad_hy_r_optgrad_cy_r_optr   rJ  r  r  r  r  r  r  r  r  diff_xdiff_hxdiff_cxdiff_w1diff_w2diff_bs                                r5   mkldnn_rnn_layer_backwardrD  :  s    4 __U[[)FmmCII&Gv||,G.G.Gw}}-F7GVVWgEEr7   )	out_int32r   c                    t        j                  | |rt         j                  nt         j                  t         j                        S )NrP   r   )rH   r   r3  r   r   )r   
boundariesrE  r   s       r5   meta_bucketizerI  ]  s2     &ekkEKK-- r7   c                 0    dt               dk(  r't        j                   j                          fd       t        j                  t	        t
              fd       t        j                  dkD  fd       t        j                  t	        t              fd       t        j                  t	        t              fd       t        j                  k\  d	        t        j                   j                   j                  
      S )Nzhistc()r  c                  $    d j                    dS )Nz%"histogram_cpu" not implemented for 'r  r   r	  s   r5   rW   zmeta_histc.<locals>.<lambda>n  s    =ekk]!L r7   c                  $     dt                S )Nz#: argument 'bins' must be int, not r  binsr  s   r5   rW   zmeta_histc.<locals>.<lambda>r  s    7)>tDzlK r7   r   c                       d  S )Nz: bins must be > 0, but got r;   rM  s   r5   rW   zmeta_histc.<locals>.<lambda>t  s    gY.J4&#Q r7   c                  $      dt               S )Nz%: argument 'min' must be Number, not r  )r  r   s   r5   rW   zmeta_histc.<locals>.<lambda>w      7)@cL r7   c                  $      dt               S )Nz%: argument 'max' must be Number, not r  )r  r  s   r5   rW   zmeta_histc.<locals>.<lambda>{  rQ  r7   c                       y)Nz&{fn_name}: max must be larger than minr;   r;   r7   r5   rW   zmeta_histc.<locals>.<lambda>}  r_   r7   r   )
r   rH   rY   r   rb   r   r   r|   rt   rP   )r   rN  r   r  r  s   ````@r5   
meta_histcrT  g  s     G5U"##%L	
 
LL4!K 
LLQR	LL3L 
LL3L 
LLMN;;tELLDDr7   c                 B    t         j                         |d      }t        j                   j	                         dk7  xs# t        d  j                         dd  D               fd        j                  |      j                  t        j                               S )Nr   r  r   c              3   &   K   | ]	  }|d kD    ywr  r;   )rd   r   s     r5   rf   z,meta_upsample_bimode2d_aa.<locals>.<genexpr>  s     !Ht$(!Hr  r   c                  *    d j                          S r  r   r	  s   r5   rW   z+meta_upsample_bimode2d_aa.<locals>.<lambda>  r  r7   r   )
r  r   rH   rY   r   r  r   r  r@   r  )r   rS  r  r  r  r  s   `     r5   meta_upsample_bimode2d_aarX    s     -

kA 
LLHc!Huzz|AB7G!HH_ ??+,//11%8 0  r7   c                 T    t        ||d      t        j                   j                  dk(   fd       t	        d      D ]0  t        j                   j
                        k(   fd       2  j                  |      j                  t        j                               S )Nr   r  r  c                  "    d j                    S r  r  ry  s   r5   rW   z4meta_upsample_bimode2d_aa_backward.<locals>.<lambda>  r  r7   c            
      D    d d     d dj                         S )NzD
Expected grad_output to have the same shape as output; output.size(r  z
but got grad_output_size(r   r  s   r5   rW   z4meta_upsample_bimode2d_aa_backward.<locals>.<lambda>  s@     DDE3dK[\]K^J_ `D!1!1!!4 59 r7   r   )
r  rH   rY   r   r   r   r   r  r@   r  )r  rS  r)  r  r  r  r  r   s   `     @@r5   "meta_upsample_bimode2d_aa_backwardr\    s     -K! 
LLAk 1X 
a $4Q$779	

   ,//11+> 0  r7   c                 P   t        j                  |j                         dk(  d        t        j                  |j                         dk(  d        t        j                  |j                  j                  d        t        j                  |j                  j                  d        y )Nr   c                       y)Nz%found_inf must be a 1-element tensor.r;   r;   r7   r5   rW   z<_amp_foreach_non_finite_check_and_unscale_.<locals>.<lambda>  r_   r7   c                       y)Nz%inv_scale must be a 1-element tensor.r;   r;   r7   r5   rW   z<_amp_foreach_non_finite_check_and_unscale_.<locals>.<lambda>  r_   r7   c                       y)Nz!found_inf must be a float tensor.r;   r;   r7   r5   rW   z<_amp_foreach_non_finite_check_and_unscale_.<locals>.<lambda>  r_   r7   c                       y)Nz!inv_scale must be a float tensor.r;   r;   r7   r5   rW   z<_amp_foreach_non_finite_check_and_unscale_.<locals>.<lambda>  r_   r7   )rH   rY   r   rP   r   )r   r  	inv_scales      r5   *_amp_foreach_non_finite_check_and_unscale_rc    s|    	LLQ O 
LLQ O 
LL))3 
LL))3r7   c                 V    t        | j                               }| j                  |      S r2   )r   r   r   )r   nanposinfneginfrb  s        r5   
nan_to_numrh    s#     tyy{#K>>+&&r7   c                    | j                   t        j                  t        j                  t        j                  t        j
                  hvsJ d| j                    d       | j                  }t        ||      }t        ||      }||k(  r| S t        | j                               }t        | j                               }||   ||   c||<   ||<   ||   ||   c||<   ||<   | j                  ||       | S )Nz>torch.transpose_: in-place transposition is not supported for z layout)rs   rH   r  
sparse_cscr  
sparse_bscr   r   r   r   r   r   )r   dim0rA  ndimsr   r   s         r5   r  r    s    ;;	   IU\]  IIE$&D$&Dt|		D$++- F!'vd|F4L&,!$ZdDJT
T6"Kr7   c                    | j                   }| j                  r8| j                         }| j                         }|dk  r|dk(  s,J d| d| d       | j	                         dk  sJ d| d       t        | d|dk  rd      S d      S )	Nr   r   zEt_ expects a tensor with <= 2 sparse and 0 dense dimensions, but got z sparse and z dense dimensionsz6t_ expects a tensor with <= 2 dimensions, but self is r  r   )r   r  r  r  rw   r  )r   rm  r  r  s       r5   t_ro    s    IIE~~__&
NN$	Q9> 	
!l,yk9JL	
1
 xxzQ 	
DUG1M	
 dAEAIq55155r7   )rE  r   sidesorterc                \    t        j                  t         j                        dk  xs  j                  d d j                  d d k(   fd       t        j                  d u xs  j                  j                  k(   fd       t        j                  |dk7  xs | d       |rt         j                  nt         j
                  }t        t         j                        r&t        j                  |t         j                        S t        j                  d| j                  	      S )
Nr   r   c                  `    dt        j                         dt         j                         S )Nztorch.searchsorted(): boundaries tensor should be 1 dimension or the first N-1 dimensions of boundaries tensor and input value tensor must match, but we got boundaries tensor z and input value tensor r   r   )r   sorted_sequences   r5   rW   z#meta_searchsorted.<locals>.<lambda>  s8    3378M8M3N2O P""&tzz"2!35 r7   c                  l    dt         j                         dt        j                         S g  S )Nz[torch.searchsorted(): boundary and sorter must have the same size, but got boundary tensor z and got sorter tensor rt  )ru  rq  s   r5   rW   z#meta_searchsorted.<locals>.<lambda>  sO    ##'(=(=#>"??V%+%7tFLL!@B  >@@B r7   r   zetorch.searchsorted(): side and right can't be set to opposites, got side of left while right was TruerG  r;   r  )rH   rY   r   r   r3  r   rb   r   r   r   r|   rt   )ru  r   rE  r   rp  rq  rP   s   ``   ` r5   meta_searchsortedrw    s     
LLO!!"a' 	9  "%CR8	
	 
LL$?///6<<?	
 
LL#e)	$ %EKK%++E$%U-D-D
 	
 {{2U?3I3IJJr7   c                      t        j                   t         j                  t         j                  t         j                  fv fd       y )Nc                      d  S )Nz/Unsupported input type encountered for isin(): r;   r   s   r5   rW   z3_check_for_unsupported_isin_dtype.<locals>.<lambda>3  s    A%I r7   )rH   rY   rU  
complex128	complex64r   s   `r5   !_check_for_unsupported_isin_dtyper|  0  s/    	LLejj%"2"2EOODDIr7   c                 j    |	rt         j                  | ||||||||
|
      S t        | ||||||||
|
      S r2   )r*   _embedding_bag_sparse_backward!meta_embedding_bag_dense_backward)r,  r   re  r  r  maximum_indicesnum_weightsr  rJ  r  r|  r  s               r5   meta_embedding_bag_backwardr  7  se     22
 	
 1
 	
r7   c
                 N    t        j                   j                  t         j                  t         j                  t         j
                  t         j                  fv  fd       |t        k(  rt        j                  |d u        j                  | j                  d      f      }
|
S )Nc                  "    d j                    S )Nz$Unsupported input type encountered: r   )r,  s   r5   rW   z3meta_embedding_bag_dense_backward.<locals>.<lambda>q  s    6tzzlC r7   r   )
rH   rY   rP   rM  rN  rL  float64r  r   r   )r,  r   r  r  r  r  r  rJ  r|  r  index_grad_weights   `          r5   r  r  b  sv     
LL

u}}ennemmU]]SSC x_D01TYYq\'BCr7   c                    | j                  d      }t        j                  |t        k(  d       t        j                  | j	                         dk(         t        j                  |j	                         dk(         |j                  d      }t        j                  |j	                         dk(         t        j                  |j                  d      |k(         | j                  |f      }	|	S )Nr   zHembedding_bag_backward: per_sample_weights only supported for mode='sum'r   r   )r   rH   rY   r  rw   r   )
r,  r+  r   re  r  rJ  r  embedding_featuresr  r   s
             r5   .meta_embedding_bag_per_sample_weights_backwardr  y  s     1	LLR 
LLq!	LL!#$,,q/K	LL"#	LLQ#556^^[N+FMr7   )assume_uniqueinvertc                   t        j                  t        | t              xs t        |t              d        t        | t              s!t        j                  | |j
                        } t        |t              s!t        j                  || j
                        }t        | j                         t        |j                         t        j                  | t         j                        S )Nc                       y)Nz<At least one of elements and test_elements must be a Tensor.r;   r;   r7   r5   rW   zmeta_isin.<locals>.<lambda>  r_   r7   r  r   )
rH   rY   rb   r   r"  rt   r|  rP   r   rU  )elementstest_elementsr  r  s       r5   	meta_isinr    s     
LL8V$I
=&(IN h'<<1E1EFmV,]8??K%hnn5%m&9&9:HEJJ77r7   r   c                     t        j                  | dk\  d        t        |t        j                        \  }}t        j
                  ||      S )Nr   c                       y)Nz,polygamma(n, x) does not support negative n.r;   r;   r7   r5   rW   z meta_polygamma.<locals>.<lambda>  r_   r7   r  r   )rH   rY   r   r   r  r   )r   r   rC   rD   s       r5   meta_polygammar    sF     
LLaOP(;HHOA| D55r7   c                     t        d      )Nz.Tensor.item() cannot be called on meta tensors)r  r   s    r5   meta_local_scalar_denser    s    
G
HHr7   c                 ,    t        j                  |       S r2   rP  r   s    r5   silur    rR  r7   c                 l    t        | t        j                        \  }}t        j                  | |      S r  )r   r   r  rH   r   )r   rC   rD   s      r5   sigmoidr    s3     );HHOA| D55r7   rE   half_to_floatc                    |r| j                   t        j                  k(  sJ t        j                  | t        j
                  j                        \  }}|s|n|}t        j                  | |t        j                        }|S )Nr  rG  )	rP   rH   rJ   r@   r   r   rA   r   r   )rE   rw   r  computation_dtyperD   r  s         r5   softmaxr    so     ww%**$$$&+&>&>	uDDLL'#| (5<:KL


1L@W@W
XCJr7   r  r  r  c                    | j                         dk(  sJ d       | j                  }|j                  }|j                  dk(  r|d   f}n$|j                  dk(  r|d   |d   f}n
g ||d   }| j                  }| j	                  ||      S )Nr   z'weight' must be 2-Dr   r   r   )rw   r   r   rP   r   )	r+  r   r  r  r  weight_shapeindices_shaper   r/  s	            r5   	embeddingr    s     ::<1444<<LMMM||q&21o%7			"1%|A7	5m5\!_5	IIY77r7   max_lengthspadding_valuec                     t        |      dk(  sJ t        |      dk(  sJ |d   j                  d   dz
  }|d   }||g| j                  dd  }| j                  |      S r  )r   r   r   )r   re  r  r  r  rf  rU  s          r5   $meta__jagged_to_padded_dense_forwardr    sv     w<1{q   
aAAAq,6<<+,LL))r7   c                 B    t        |       t               d               }|S )Nc                 8    t        | t        j                        S r  rF   r   r  r  s    r5   _fz)_create_unary_float_meta_func.<locals>._f  s      =JJ
 	
r7   r<   r#   funcr  s     r5   _create_unary_float_meta_funcr    *    4]
  

 Ir7   c                 B    t        |       t               d               }|S )Nc                 :    t        | |t        j                        S r  r  )rE   r
  s     r5   r  z*_create_binary_float_meta_func.<locals>._f  s      q!@!M!M
 	
r7   r  r  s     r5   _create_binary_float_meta_funcr    r  r7   c                      t                fd       } j                   d}||_         t        t        t        |            |      }|S )Nc                 `     | g|i |}t        | j                  |j                         | S r2   r  )r   rB   r  r   r4   s       r5   _fnz#_register_inplace_meta.<locals>._fn*  s.    '''

CII6r7   rC   )r   rm   r<   getattrr*   )r4   r  inplace_names   `  r5   _register_inplace_metar  )  sO    
2Y 
 kk]!$LCL
4-l3
4S
9CJr7   c                 f    t        j                   j                  j                  k(   fd        g}t        t              rQj
                  dk7  r1t        j                   j                  j                  k(   fd       |j                         t        |dt        j                  iS )Nc                  <    dj                    d j                    S )Nr`  z for `end`, but got dtype r   )rn   ro   s   r5   rW   zlerp.<locals>.<lambda><  s    /%++.HT r7   r   c                  <    d j                    dj                    S )Nr`  z for `weight`, but got dtype r   )ro   r+  s   r5   rW   zlerp.<locals>.<lambda>C  s!    /%++6STZT`T`Sab r7   r=   )
rH   rY   rP   rb   r   r   r   rF   r   rA   )ro   rn   r+  rB   s   ``` r5   lerpr  7  s     
LLsyy T 3<D&*%;;!LLv||+b 	F	=EE r7   )r  c                <    t        | ||t        j                        S r  r  r   tensor1tensor2r  s       r5   addcmulr  K  s!     w0O0W0W r7   c                    t        j                  t        j                  |j                        xr t        j                  |j                         d        t        | ||t        j                        S )Nc                       y)N)zFInteger division with addcdiv is no longer supported, and in a future zErelease addcdiv will perform a true division of tensor1 and tensor2. z4The historic addcdiv behavior can be implemented as zA(input + value * torch.trunc(tensor1 / tensor2)).to(input.dtype) zfor integer inputs and as z6(input + value * tensor1 / tensor2) for float inputs. z?The future addcdiv behavior is just the latter implementation: z4(input + value * tensor1 / tensor2), for all dtypes.r;   r;   r7   r5   rW   zaddcdiv.<locals>.<lambda>[  r_   r7   r  )rH   rY   r@   r,  rP   rF   r   rA   r  s       r5   addcdivr  S  sb     
LL""7==1 6&&w}}5	
		
  w0O0W0W r7   c                     i } dD ]  }t         |   }|D ]  }|| vs||   | |<    ! | j                         D ]  \  }}t        |t        j                  j
                        r,t        |t              sJ  |j                  t        j                  j                  j                        |       t        j                  j                  |j                         d      r|t         d   v st        | d      |j                  r|j                         dv rd|j                         v rt        j!                  ||       	d|j                         v rt"        j!                  ||       3d|j                         v rt$        j!                  ||       ]d	|j                         v rt&        j!                  ||       t(        j!                  ||        y )
N)rr   post_autogradpre_autogradCompositeImplicitAutogradrr   z is a CompositeImplicitAutograd op, we shouldn't register meta function for it. Instead, we should let the decomposition run and write meta kernels for the base operators.>   aten::cloneaten::copy_aten::rot90aten::_to_copyaten::empty_stridedaten::constant_pad_ndaten::as_strided_scatterzmkldnn::zmkl::zonednn::zquantized::)r   itemsrb   rH   _opsHigherOrderOperatorr   py_impl_CDispatchKeyr,   %_dispatch_has_kernel_for_dispatch_keyr  r  is_view2_meta_lib_dont_use_me_use_register_meta_for_mkldnnimpl/_meta_lib_dont_use_me_use_register_meta_for_mkl2_meta_lib_dont_use_me_use_register_meta_for_onednn5_meta_lib_dont_use_me_use_register_meta_for_quantized'_meta_lib_dont_use_me_use_register_meta)activate_meta_tablerl   registryopoop_overloadr4   s         r5   activate_metar  w  s    : 9-d3 	9C--+3C=#C(	99 /446 6NR
 k5::#A#AB+z2226EHH00556r:8899 ;
 8@@""m $; ; 
    	 [--//BGGUWXK,,..?DD[RTU{//11BGGUWX+"2"2"44EJJ 8<<["Mm6Nr7   )Fr#  r2   )NNNFr   r   Fr=  )Tr  )r  )r  T)FF)TT)rL  )FTN)TFF)TF)r   )g      ?N)r.   str)r;   r  r  F)r;   r  FTN)Fr   FNFr   )NF)r   F)g      ?gUUUUUU?FN)NNNNN)r   NNr   )NNF)        FFN)Nr  FFN)r  FNN)r  FN)FN)FNNNN)NNNF)Nr   FNN)NNNN)r   TT)NNr   N)d   r   r   )r   )r   FF)r  (  rk  collections.abcr   enumr   	functoolsr   typingr   r   r   r	   typing_extensionsr
   rH   torch._prims_commonr  r@   r   r   r   torch._decompr   r   r   r   
torch._opsr   torch._primsr   r   r   r   r   r   r   r   r   r   r   r   torch._prims_common.wrappersr   r    r!   r"   r#   r  r$   r%   torch.fx.experimentalr&   r  torch.utilsr'   r8   r(   r)   opsr*   libraryLibraryr  r   r  r  r  r<   rF   rR   r[   linspacelogspacer  r~   taker  r   r   r   r   cummaxcumminr   r   r   r   r   rU  r   _fft_c2cr   r   r   _fft_r2cr   randpermgenerator_outr   r   r   randintr  r  low_outr  randr  _fft_c2rr  r  r  r&  
unsqueeze_r*  _sparse_semi_structured_linearr  rP   r7  _sparse_semi_structured_mmr=  _sparse_semi_structured_addmmrA  _cslt_sparse_mmrV  index_reducer]  index_reduce_r_  index_selectrc  segment_reducerp  r  	unary_outrt  rw   r}  r   r  r  r  r  r  _assert_asyncr  msgr  _printr  _make_dep_tokenr  r  _functional_sym_constrain_ranger  r  (_functional_sym_constrain_range_for_sizer  _functional_assert_asyncr  r   r  r   r  r  r  r  _linalg_eighr  r  _linalg_eigvalslinalg_eigvalsr  
linalg_eigr  r  r  r  r  r   r  r  linalg_inv_exr  linalg_ldl_factor_exrX   r!  linalg_ldl_solver0  	linalg_lur6  linalg_lu_factor_exr:  linalg_lu_solverC  	lu_unpackrI  rR  	linalg_qrrY  r]  rZ  _linalg_svdrg  r-  r  rw  r  linalg_solve_triangularr  r  r  _linalg_detr  r  r  r  reflection_pad1dr  replication_pad1dr  r  reflection_pad1d_backwardr  replication_pad1d_backwardr  r  reflection_pad2dr  replication_pad2dr  reflection_pad2d_backwardr  replication_pad2d_backwardr  r  reflection_pad3dr  replication_pad3dr   reflection_pad3d_backwardreplication_pad3d_backwardr'  _pdist_forwardrL   r+  _pdist_backwardr1  baddbmmrE  	bernoullirH  
bernoulli_rK  r(  rN  poissonrQ  _fused_moving_avg_obs_fq_helperrc  mmrm  rz  r   r  r  miopen_batch_normr  convolutionr  r  _has_mkldnnr  r  _convolution_pointwiser  _linear_pointwiser  has_mklr  r  _mkl_linearr  r  r  qconv2d_pointwiser  binaryr  qlinear_pointwiser"  r  binary_tensorr  linear_dynamic_fp16linear_relu_dynamic_fp16r  r  r  
max_pool2dr  int4mm_packed_weight_cpur  r  
avg_pool2dr  r"  avg_pool2d_backwardr*  
avg_pool3drD  avg_pool3d_backwardrO  _adaptive_avg_pool2drT  _adaptive_avg_pool3drW  _adaptive_avg_pool2d_backwardra  _adaptive_avg_pool3d_backwardrf  rd  adaptive_max_pool2drv  rz  r|  adaptive_max_pool3dr  r  r  repeat_interleaver  rc   r  r  r  r   _unsafe_indexr  convolution_backwardr  addbmmr  _fused_adam__fused_adamw_r  _fused_adamr  _int_mmr  _convert_weight_to_int4packr  #_convert_weight_to_int4pack_for_cpur   _weight_int4pack_mmr
  _weight_int4pack_mm_for_cpur  r  r:  _dyn_quant_pack_4bit_weightrF  _dyn_quant_matmul_4bitrL  _weight_int8pack_mmrS  _cdist_forwardrh  _cdist_backwardrr  _embedding_bagr  _embedding_bag_forward_onlyr  r  nansumr  median	nanmedianr  
dim_valuesrJ  r   r  logical_not_r  repeatr  zero_r  mul_Scalardiv_logical_and_logical_or_logical_xor_r  add_sub_r  rounddecimalsr  r  
__rshift__r  
__lshift__r  zeror  r  r  fillr  relu_r  	_add_relur  rrelu_with_noiser  rrelu_with_noise_functionalr  rrelu_with_noise_r  	index_put_unsafe_index_putr  masked_fill_r  _masked_scaler  masked_scatter_r  masked_scatterr  masked_scatter_backwardr  
index_put_r  aliasr  r  bmmr  r  r  r  r  r9  r4  rK  r   max_pool2d_with_indices_backwardrA  max_pool2d_with_indicesrC  fractional_max_pool2drU  max_pool3d_with_indicesrb   max_pool3d_with_indices_backwardri  rs  ru  r  grid_sampler_2d_backwardr  r  r  r  rm  r  r  select_scatterr  slice_scatterr  r   r  r  gatherr  r  r  r  r  r  scatter_addr  scatter_add_r  r  r  r  rY  value_reducer  scatter_r  #_scaled_dot_product_flash_attentionr	  #_scaled_dot_product_cudnn_attentionr  0_scaled_dot_product_fused_attention_overrideabler  ,_scaled_dot_product_flash_attention_backwardr   +_scaled_dot_product_flash_attention_for_cpur#  4_scaled_dot_product_flash_attention_for_cpu_backwardr)  '_scaled_dot_product_efficient_attentionr-  0_scaled_dot_product_efficient_attention_backwardr5  ,_scaled_dot_product_cudnn_attention_backwardr7  _flash_attention_forwardr?  _flash_attention_backwardrD  _efficient_attention_forwardrP  _efficient_attention_backwardSymIntrX  
_scaled_mmry  scatter_reducetwotwo_outr}  scatter_reduce_r  multinomialr  r  r  r  _upsample_nearest_exact1dr  _upsample_nearest_exact2dr  "_upsample_nearest_exact2d_backwardr  _upsample_nearest_exact3dr   r  values_stabler  r  _thnn_fused_lstm_cellr  r  r  r  r  r  argminr  r  topkr
  _segment_reduce_backwardr  kthvaluer  r   r$  r!  r(  r+  pixel_shuffler2  rD  	bucketize
Tensor_outrI  histcrT  _upsample_bilinear2d_aa_upsample_bicubic2d_aarX   _upsample_bilinear2d_aa_backwardr\  rc  rh  r  ro  searchsortedrw  r|  _embedding_bag_backwardr  _embedding_bag_dense_backwardr  *_embedding_bag_per_sample_weights_backwardr  isinr  	polygammar  _local_scalar_denser  r  r  _softmaxr  r  _jagged_to_padded_dense_forwardr  r  r  special_airy_aispecial_bessel_y0special_bessel_y1special_modified_bessel_i0special_modified_bessel_i1special_modified_bessel_k0special_modified_bessel_k1!special_scaled_modified_bessel_k0!special_scaled_modified_bessel_k1special_chebyshev_polynomial_tspecial_chebyshev_polynomial_uspecial_chebyshev_polynomial_vspecial_chebyshev_polynomial_w&special_shifted_chebyshev_polynomial_t&special_shifted_chebyshev_polynomial_u&special_shifted_chebyshev_polynomial_v&special_shifted_chebyshev_polynomial_wspecial_hermite_polynomial_hspecial_hermite_polynomial_hespecial_laguerre_polynomial_lspecial_legendre_polynomial_pr  r  r  r  lerp_addcmul_addcdiv_torch._refs.nn.functionaltorch._refs.specialr  r;   r7   r5   <module>r	     snD    $   5 5 '  # + +  " U    < 7 ) T]t_yy~~*/--*?*?PV*W ' %a )X
8BF#3"4hr2v6F"FG 
3(* t}}-.
 

==5  /5p 		!!499==12'  3' !!))4+<+<+@+@AB%' %  C%$ t%%&I  'I 	[[$++//4;;+>+>P Xy! " !!))4+<+<+@+@ABI  CI3lV $s) 4  %%t}}'8'89:K  ;K $s)  %%t}}'8'89:=
  ;=
@ t}}**+"& 3 ,3 t}}$$% **
 &
 $$dll&6&678
 **  9&   $,,"6"678 **  9& 		!!499==12%)$tPT   3 %%t}}'8'89:$Dv $DDI $Dc $DC $D  ;$DN tzz!!" #0( t&&' ( t223
 "%)'+  6
	
 c] $ 4B t../
 (,	
  $	 00 t112 	
'+
  	 $ 3D t##$ ""'+"$-<,,-<\\-< 6
-< F	-<
 $-< -< -< -< -< %-<` t  (() 	I
	I		I 	I LL		I
 	I 	I 	I *	I t!!))* 	
			 	 LL		
 	 	 	 +	 t  (()' * ' t""**+
 !% $ $ W
 W W f	 W
 f W f W  W  W  W , WF   $(("4"456  7 txx||    $(("4"456  7 txx||  tzz!!"6 #6 tzz~~( (
 t!!))* + t!!%%& ' t{{""# $ t##++, ) -) t''//0, 1, t33;;< =
 t008896 :6& t<<DDE F
 t,,001 2
F C    F  #  N (,


 !%
$V S C 
 
F 
$ 
 
"  	  	C  !!))4+<+<+H+HIJ]N+ s T  , K" $$,,d.A.A.E.EFGB B6 B  HB  !]N+	6 	 , "	Q QF Q t**+) )F )4 )F )  ,) t""#J JF J4 JF J  $J t}})6 )$ )6 )  ) t$$%)6 )$ )6 )  &) t&&../&  T  0" 	$$,,d.M.M.Q.QR .f .6 .f . .d t!!))*&   + ))1143L3L3P3PQRT8V$ 	
  	
 666!" % S& %%--t/D/D/H/HIJ ''' '
 ' '  K'T &&(:(:;<S#s/3 f  fff>T8U   =4 ((00$2J2J2N2NOPT8V$ 	  	
 666!" % QD $$,,d.B.B.F.FGH 444 4
 4 4 4  I4n t~~S#s 	$$$ $ 	$
 666!"$  $P tTz!2 * &&(:(:;<S#f C ffn8M   =4 $$,,d.B.B.G.GHIV[$1'v '%(F"G ' 2 J'$ t''(   	""" " SM	" )"J.
.
. 49d3i .".
.
. 3-. 66>	.(f V   t$$%
 ##!777 	7
 7 V7 	7 V7 6
7 6666)*7 &7t ,,44d6R6R6V6VWX   	
   
&	  Y2 t$$%S#4( +(
+(+( +( 	+(
 +( 66>+( ) &+(^ t''(
 )
 tzz
 WW	W W 	W
 W W  Wt>#;L t$$%=  &= t%%&>  '>(< t--.\S  /S t../\T  0T2Ej t$$%=  &= t%%&>  '> &&..&&11''//''22	 \& &:<G~ t$$%=  &= t%%&>  '> &&..&&11''//''22	 \$( $(N t""#

f 

 

v 

  $

 t##$Pv PV P Pf PQW P  %P $$dll&6&678/0 '  9'2 &&(:(:;<&* I  =I
 t$$% & t~~ I !I
 $$dll&6&678"  9" t33;;< * =*. tww	  	B
* 7;h,,hLLh $s)S.!h 49c>"	h
 DIsN#h h h U49c>23hVQ t%%--."$,,"$LL"$ 5<<
 "$ 5<<(	"$
 %,,'"$ "$ !&"$ "$ /"$J t''(),,)LL) ,,) I	)
 #Y) 3i) ) I) ) ))X 	889>9N9N&&:6 599##::BBC D, 599##55==>S ?S
 xx:?--:O:O66;
7 
uyy}}00	1	 
2	
 :?9N9N&&:6 599##55==>  ? D 599##55<<= >6 599##55==>599##55<<= > ?, 599##55<<=599##55CCD E >: 599##77??@599##<<DDE	 F A	 =BMM<Q<QVV=9 599&&112 
 3
8 599&&??@@ A@( t&&' M (Mb(<X t''//0E 1EP t UJ   UJp t''(\K(  )K(\ t((001 2" t((001@ 2@ t1199:F ;F, t112\P  3P
	
6 	
S 	
 t''(UI+  )+\ t001\H  2H$ t''(UI'  )'T t001\(  2(
 t%%,,-* .* $$dll&6&678T  9T ##++T-@-@-D-DEF46 @c @  G@ 		&&..		0F0F0J0JKL  M" 

!!4#5#5#<#<=> ?D ))1123H 4H: ##T[[__56./q '  7'0 !!))4+=+=+E+EFG  !
 H
2   (()*  !! +!H ~B  B* 0012 3& 889: ; (()*@ +@ 0012< 3<"3 "3 "3 "jZ 0012D!)&!1D 3D0 ++,-; .;  (()*< +< t""**+ & , &F t##$G  %G* t""**+
 	
`5 ,`5F t//7785 95
 ##T[[__56=$ =  7= ##T^^%;%;<=) >) !!					 Xy! "	 t  (() * t{{""#' $' tzz!!" # 								!!  !!

 									**Z 

""DJJ$7$789 :
" &&(>(>?@ A &&(>(>?@ A tyy  !& "& 

!!4::#4#456 7 		  $))"2"234" 5" tzz!!" # t~~$$%F   & %%&'RV"  (" 0012RV; 3; &&'(KO ) &&(>(>(F(FGH" I" t  ''( )
 t!!))* + t##$	 %	 t""#6  $6 t++,! -! t&&' ( tzz!!"! #!B txx 5 !56;h #-YYY 	Y 		Y
 	Y 	Y 	Y 	Y 	Y 	Y 	Y Y Y Y Y  !Y" #Y$ %Y& 'Y( )Y* +Y, -Yx;4|383838 38 		38
 	38 	38 	38 	38 	38 	38 	38 	38 38 38 38  !38" #38$ %38& '38lI2X t44<<=( >(V t++334 # 5#L t))112Q 3Qh t++,UI d  -dN t445\b  6bJ%
V %
6 %
Pt  v 3 $ t,,445# 6#$ t##$8  %8" t,,-\;'! ( .!, 		!!"#. $. t&&' ) ()X t{{E  E@ t""**+. ,. t!!))*. +.
	 	C 	d 	/
  t{{""#' $'6
 
-
b4 t''(& )&
 t  ! "
 !!	 & & ""	 889:
 #!==	= = 	=
 = = E?= ;=@ 889: #!&&	& & 	&
 & & & & E?& ;&R EEFG
 #'#!%%	% % 	%
 % % % E?% H%P 99( """" 
" 	"
 
" " " " " " " " " " E?"
". 88 "&!	  	
   E?
: AA #'!'"'"'" 
'" 	'"
 
'" '" '" '" '" E?'"
'"T <<=> !))))	)) )) 	))
 )) )) E?)) ?))X ==" !4-4-4- 
4- 	4-
 4- 
4- 4- 4- 4- 4- $Z4- 4- E?4-
4-n 99* "!""" 
" 	"
 
" " " " " " " " " " "  E?!"
"0 %%  "&*'+"&%)HH	H H 	H
 H H H H H H E?H smH  }H H 6"H
HV &&( "&*'+#,,, 
, 	,
 
, , , , , , , , , , E?,  sm!,"  }#,
,4 ))   %!(,!%!%,S,S	,S ,S 6
	,S
 6",S 6",S 3-,S 3-,S ,S ,S ,S E?,S f%,S v,S #,S
,S^ *** "$("'%474747 
47 	47
 6
47 6"47 6"47 ,,47 ,,47 47 47 47 47 47 47  E?!47" SM#47$  %47
47n ''() $(+/'+ aY
,,aY
,,aY \\aY \\	aY
 5<<
 aY 5<<(aY $aY aY *aYH ##'')<)<)D)DEF&  G&
 t##''( )
   (($*:*:*>*>?@	 	  A	,* 	$$d&D&D&L&LM

 	$$d&D&D&L&LM. ((00//77 !% $%U\\ 123 sELL012 uo	
 uo: 	$$d&D&D&L&LM

 									&$N t))112
  3 t&&'4/ (4/n t$$,,-$% .$%N


 ##T[[%8%89:4 ;4 t!!))* + tyy  !
Q "
Q t,,-LP  .  %%t}}';';<=Xy!K " >K #("9"9 Q t77??@	* A	* t##++,	0 -	0 t!!))* +> t--556F 7FD %%t~~'@'@AB27u   C 

|E  E0 	!!))4+F+F+N+NO & 55==>?  @8 t>>FFG H$ '')<)<=>'  ?'
 uyy~~(() *8 uyy~~  !6 "6$ t  !
 
	-K  "-K` t++, '
 -'
T t112  3, t>>?  @. tyy8=e 8  8  t~~6c 6 6F 6  6 t''(I& I )I tyy"v "& "  " t||6& 6V 6  6 t}}	v 	C 	 	 	  	 t~~ $888 8 	8
 8 8  8, t33;;<
 	**&\* c* 	* =*  d22 3 d44 5 d44 5 d== > d== > d== > d== > dDD E dDD E tBB C tBB C tBB C tBB C tJJ K tJJ K tJJ K tJJ K t@@ A tAA B tAA B tAA B tyy  $ t||./    t||./   , 	tyy)!$,,/!$,,/
    BNJ r7   