8i -pddlZddlmZmZddlmZddlmZddlm Z ddl m Z ddl Z ddl mZddl mZmZmZddlmZmZmZmZdd lmZdd lmZmZdd l mZmZmZm Z m!Z!m"Z"m#Z#m$Z$m%Z%m&Z&m'Z'dd l(m)Z)m*Z*m+Z+m,Z,m-Z-dd l.m/Z/m0Z0ddl1m2Z3ddl4m5Z5m6Z6ddl7m8Z9e dZ:e dZ;e j<j=Z=e j>?dddZ@eAd\ZBZCZDdZEdZFdeee;e:fgee;e:fffdZGde!fdZHdZIdZJeGe=jKe=jLge-ddde jMddfd ZNeGe=jOjPe=jOjQge-d!ZReGe=jSjPe=jSjQge-d"d#d$ZSeGe=jTe-d%ZTeGe=jUjPe=jUjQe=jVjPe=jVjQge-d&d'd(ZWeGe=jXjPe=jXjQge-d)ZXd*ZYd;d+ed,eZe[d-e\fd.Z]eGe=j^jPe=j^jQge-d/Z_dZ`d,eZe[fd0ZaeGe=jbjPe=jbjQge-d1ZceGe=jdjedd2d3ZfeGe=jdjPe jgdddd4d5ZheGe=jijPe=jijQge-e jgdddd4d6ZjeGe=jijke=jijlge-e jgdddd4d7ZmeGe=jnjPe=jnjQge-ddddd4d8ZoeGe=jpjPe=jpjQge-d+ed,eZe[d9e[d:e[fd;ZqeGe=jrjPd;d<Zsd=ZteGe=jujPd>ZveGe=jw ddMe jdNe jdBedzdOedzdDe jydzdPe\dQe[dRe[dSe[fdTZeGe=jjPdUdVd+ed,e[dWedXe jdYexdZe\defd[ZeGe=jjPdUdVd+ed,e[dWedXe jdYexdZe\defd\Ze-eGe=jjPd]ZeGe=jjPddddddd^d_edYexd`edzd'edzdaedzdbe[dce\defddZeGe=jjPe=jjge-deZeGe=jjd;dfZeGe=jjPe=jjge-dgZeGe=jjd;dhZeGe=jjPdiZeGe=jjQdjZeGe=jjPdkZeGe=jjdlZeGe=jjPdmZeGe=jjPddddddndoZeGe=jjPd?dpZeGe=jjPd?dddZeGe j<jjjPdZeGe j<jjjPdZe jjrEe j>?dddZeGe j<jjdZ e j>?d ddZ!eGe j<j"j#jPeGe j<j"j$jPeGe j<j"j$j%d Z&eGe j<j"j#j'eGe j<j"j#j(d Z)eGe j<j"j*jPeGe j<j"j*j%d Z+eGe j<j"j*j'eGe j<j"j*j(d Z,eGe j<j"j-jPeGe j<j"j.jPdZ/e j>?dddZ0eGe j<j1j2 dMdZ3eGe j<j1j4dZ5dZ6eGe=j7jP dNdZ8dZ9eGe=j:jPdZ;eGe=j<e- dNdZ=eGe=j>e-d֦dZ?eGe=j@jPdZAeGe=jBjPdZCeGe=jDjPdZEeGe=jFe-d֦dZGded}exfd ZHeGe=jIe-dd'd!ZJeGe=jKe-d֦d"ZLeGe=jMe-dd'd#ZNeGe=jOe-d֦d$ZPeGe=jQjd=d%ZReGe=jSjPe=jSjQge-d&ZTeGe=jUjPe=jUjQge-d"d'd(e[fd)ZUeGe j<j=jVjPe j<j=jVjQge-d*ZVeGe=jWje=jXjgd+ZYeGe=jZjPgd,Z[eGe=j\jPe=j\jQge-dU̦dJdJdd-Z]eGe=j^jgd.Z_eGe=j`jPe=jajPgddd/d0ZbeGe=jcjPgddd/d1ZdeGe=jege-d2ZfeGe=jggd3ZheGe=jigd4ZjeGe=jkgd5ZleGe=jmgd6ZneGe=jogd7Zod8e[d9e[de[fd:Zpd;ZqeGe=jrgdBedzfd<ZseGe=jtgd=ZueGe=jvgd>ZweGe=jxjPd?ZyeGe=jze-d@Z{eGe=j|jP dOdAZ}eGe=j~jPdBZd@dCZeGe=jjPe=jjQge-dPddDdEZeGe=jjPe=jjPgdFZeGe=jje=jje=jje=jje=jjPe=jjge-d&d'dQdGZeGe=jjPdHZeGe=jjPdIZeGe=jjPdJZeGe=jje=jje=jje=jje=jjPe=jjPe=jjPgdKZeGe=jje=jje=jje=jjgddLZeGe=jje=jjgddMZeGe=jjPe=jjgdNZdOZeGe=jje=jjgdPZeGe=jje=jjgdQZeGe=jjPdRZeGe=jje=jjgdSZeGe=jje=jjgdTZeGe=jjPdUZeGe=jje-ddefdVZeGe=jge- dRdYZeGe=jg dRdZZeGe=jg dRd[ZeGe=jjPe=jjPgd;d\ZeGe=jjd]ZeGe=jjPd^ZeGe=jd_ZeGe=je-d`ZeGe=jdaZeGe=jjPd;dbZÐd?dcZeGe=jjPddZeGe=jjydeZǐdfZȐdgZɐdhZʐdiZ d;d?edje[dke[dle[dme[dne[doe[dpe[dqe[dre[dse[dte[due[dve[dwe[dxe[dye[dze[d{e[d|e[dexd}e\f,d~Z̐dZd?ededje[dke[dle[dme[dne[doe[dpe[dqe[dre[dse[dwe[dxe[dye[dze[d{e[d|e[dexf&dZΐdZeGe=jjPdZeGe=jjP dMdZeGe=jjPdZeGe=je-dd' dMdZeGe=je-d֦dZd?edefdZGddeZd?edede[fdZeGe=jjPdZeGe=je-dZeGe=je-d֐ddZeGe=jjPgdZeGe=jjP dSdZeGe=jjPe=jjQge-dddddddZeGe=jjPe=jjQge-dddddddZeGe=jjPdZeGe=jjPdTdZd@d,e[de[de\fdZdZdZeGe=jjPd;dZd;dZd=dZdZd=dZdUdZeGe=jjPdZeGe=jdZeGe=jje=jje=jje=jjge-d=dZeGe=jje=jje=jje=jjgd=dZeGe=jg dVdedededede\de\dedzfdZdedee[dffdZeGe=jg dVdededededzde\dede\de\dedzfdZeGe=jg dWdededededzdede\de\dedzfdZeGe=jg d=dedededededededede[de[dede\dedededzfdZ eGe=j g dXdedededede\dedzdedzfdZ eGe=j g d?dededededededede\dedzdedzfd„Z eGe=jg dYdededededzdede\dedzdedzdeeeffdĄZeGe=jg dZdededededzde\de\dedzfdńZeGe=jg d[dedededededzdedededededeZe\de\dedzfdDŽZeGe=jg d=dededededededededededede[de[dede\dedzf dȄZeGe=jg dSdededededzdedzde[de[dede\de\dedzde[dzde[dzdedzdedzfd̈́ZeGe=jg d<dedededededededede[de[dede\dedededzde[dzde[dzf"d΄ZeGe=jg d\dedededBedzdedzdedzde[dzde[dzdede[de\dedzdedzdedzde[dzfdׄZeGe=jg dUdededededBedzdedzdedzde jde jdedededede[de\dedzde[dzde\f$dۄZ d]d+e jdHe jde jde jdBe jdzde jdzdDe jydzde\fdZeGe=j jPg d]d+e jdHe jde jde jdBe jdzde jdzdDe jydzde\fdZ! d^d+e jdHe jdeZe jdeZe5deZe jdeZe5dBe jdzdDe jydzdeZe6dzdeZe6dzde\fdZ"eGe=j#jPg d]d+e jdHe jdeZe jdeZe5deZe6deZe jdeZe5deZe6dBe jdzde jydzdeZe[dzde\fdZ$eGe=j%j&e=j%j'ge-d@dZ(eGe=j)j&d@dZ*eGe=j+jPe=j+jQge-d;dd2dZ,dZ-dZ.eGe=j/jPe=j0jPgd=dZ/eGe=j1jPe=j2jPgd?dZ1eGe=j3jPe=j4jPg d?dedee[e jzdee[e jzdedzdedzf dZ3eGe=j5jPe=j6jPgd<dZ5eGe=j7jPe=j7j8e=j7je=j7j9gd_dZ:dZ;eGe=j<jP d?dZ=eGe=j>jPdZ>eGe=j?jPdZ?dZ@dZAeGe=jBjPe=jCjPgdPdZDeGe=jEjPd`dZEeGe=jFjPdadZGeGe=jHe- dbdZIeGe=jJjPe=jJjge-d&d'dQdZKe jLZMdZNeGe=jOjPdZOeGe=jPjPdZPeGe=jQjPdZReGe=jSjPdZSeGe=jTje=jTjUge-dddd ZVeGe=jWge-dcd ZXeGe=jYjPe=jZjPg d?d Z[eGe=j\jPg d?d Z]eGe=j^jPdZ^eGe=j_jPe=j_jQge-d<dZ_eGe j<j=j`dZ`eGe j<j=jadZaeGe=jbe-ddddddZcdZdeGe=jedZfeGe=jg dddZheGe=ji dddZjeGe=jk dddZleGe=jme-ddddZneGe=joe-de[d+edefdZpeGe=jqd+efdZreGe=jse-dU̦d+edefdZseGe=jte-d+edefdZtd Zu d^d!ed"ede jdzde jdzd#edzdBedzde jdzdDe jydzde\fd$ZveGe=jwe- d<d!ed"ed#edzdBedzdDe jydzdef d%ZxeGe=jyg d^d!e jd"e jde jde jd#e jdzdBe jdzde jdzdDe jydzde\fd&ZzeGe=j{e-d'ed,e[d(e\defd)Z|eGe=j}e-dd*Z~eGe=je- ded@ed'ed+e[d,e\d-e\def d.ZeGe=jjP dfd&edaeZed/eZe[d0efd1Zd2ZeGe=jjP dgd3Zd4Zee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jd5ZeGe=je-d6ZeGe=je-dJd7d8ZeGe=je-dJd7d9Zee=jZee=jZee=jZddl.Z ddlZ ddlZ d:ZedS(hN)CallableSequence)Enum)wraps)TypeVar) ParamSpec)SymBoolSymFloatTensor)_add_op_to_registry_convert_out_paramsglobal_decomposition_table meta_table) OpOverload)_prim_elementwise_meta$ELEMENTWISE_PRIM_TYPE_PROMOTION_KIND) BoolLikecorresponding_complex_dtypecorresponding_real_dtypeelementwise_dtypesELEMENTWISE_TYPE_PROMOTION_KIND FloatLikeIntLikemake_contiguous_strides_forNumbersuggest_memory_format TensorLike)_maybe_convert_to_dtype_maybe_resize_out_resize_output_check_safe_copy_out out_wrapper)_broadcast_shapes_maybe_broadcast)_config) ScalingType SwizzleType)_pytree_T_PatenIMPLMetac||zdz |zSNabs L/var/www/icac/venv/lib/python3.11/site-packages/torch/_meta_registrations.pyceil_divr77s EAI! c||zdz |z|zS)z$Rounds up x to nearest multiple of yr1r2xys r6round_upr=;s UQY1  !!r8returncfd}|S)Nc\tfd}tj|S)Nc4tt|dSN)r r)opfns r6registerz0register_meta..wrapper..registerDs  B 3 3 3 3 3r8)r pytree tree_map_)rDrErCs` r6wrapperzregister_meta..wrapperAsD  $ $ 4 4 4 4 4 2&&& r8r2)rCrHs` r6 register_metarI@s# Nr8type_promotionctj|d|i\}fd|D}t|}t|dtjiS)Ntype_promotion_kindc0g|]}t|Sr2)r).0r; result_dtypes r6 z$elementwise_meta..Vs$ C C C #A| 4 4 C C Cr8rJ)utilsrr$rrDEFAULT)rJargs_rOs @r6elementwise_metarUMsp . *OA| D C C Cd C C CD T "D "  BJ  r8ctjtjtjtjtjtji}|||SrB)torch complex32halfcfloatfloatcdoubledoubleget)dtype from_complexs r6toRealValueTyperaas<  ek u|L   E5 ) ))r8cxttg|RtjkfddS)NcddS)Nzoutput with shape z# doesn't match the broadcast shape r2)broadcasted_shape self_shapesr6z)check_inplace_broadcast..nsgZggTeggr8)tupler#rW_check)re args_shaperds` @r6check_inplace_broadcastrjjsX/ HZHHHII LZ'gggggr8Fc \ ttjr,tjdkdttjr,tjdkdt dfDrVt jtj  nAtjt j  fdnptjttj sJtj ttfdttsJtjdkdtj f|d|| S) NrcdSNz:linspace only supports 0-dimensional start and end tensorsr2r2r8r6rfz(meta_linspace_logspace..Pr8cdSrmr2r2r8r6rfz(meta_linspace_logspace..rnr8c3@K|]}t|tVdSrB) isinstancecomplex)rNargs r6 z)meta_linspace_logspace..s, C C:c7 # # C C C C C Cr8cddS)Nzlinspace(): inferred dtype z& can't be safely cast to passed dtype r2)default_complex_dtyper_sr6rfz(meta_linspace_logspace..sz6Kzzsxzzr8cdtjdtjdtjdS)Nz4received an invalid combination of arguments - got (, ))type__name__)endstartstepssr6rfz(meta_linspace_logspace..sYHu++HHs)),HH04U 0DHHHr8cdS)Nz$number of steps must be non-negativer2r2r8r6rfz(meta_linspace_logspace..s%Kr8metar_layoutdevice pin_memory requires_grad)rqrWr rhdimanyrQrget_default_dtypeis_complex_dtyper_ _check_typerempty) r}r|r~baser_rrrrrvs ``` ` @r6meta_linspace_logspacerrs%&&  IIKK1  P P   #u|$$  GGIIN P P    C CsE/B C C CCC 3 % A  # % %! !  =)EE L&u--zzzzz     2022 eU[ ) )))) 5'"" H H H H H H eW % %%%% L!KKLLL ; #    r8ctjjtjkfdtj|dkodk d|jS)NcdjS)Nz2take(): Expected a long tensor for index, but got r_indexsr6rfzmeta_take..sRU[RRr8rcdS)Nz*take(): tried to take from an empty tensorr2r2r8r6rfzmeta_take..s<r8)rWrhr_long _check_indexnumel new_emptyshape)selfrs `r6 meta_takers L uz!RRRR   ZZ\\Q  55;;==A#56<< >>%+ & &&r8rcJj}j}tj||kdtjdkodkfdt jj}|S)NcdS)Nz=linalg.cross: inputs must have the same number of dimensions.r2r2r8r6rfzlinalg_cross..Or8r.cfdddS)Nzlinalg.cross: inputs dimension z must have length 3. Got  and size)rotherrsr6rfzlinalg_cross..sG :c : :99S>> : :(- 3 : :r8)ndimrWrhrr#rr)rrrx_dy_d out_shapes``` r6 linalg_crossrs )C *C L s OO L #!4 31 4      "$*ek::I >>) $ $$r8ct|dt|dtj|tjS)Nzlinalg.matrix_exp memory_format)squareCheckInputscheckFloatingOrComplexrW empty_likecontiguous_formatrs r6linalg_matrix_exprs?d/0004!4555  D0G H H HHr8valuesindicesctj|j|j|j}tj|j|jtj}|dkr |jdkrt||j||fS)Nrr_r) rWrrrr_int64rrmaybe_wrap_dim)rrrrs r6 cummaxminrsu [DKtz J J JFk$*T[ LLLG zz||qTY!^^sDI&&& 7?r8clt||jtj|tjSNr)rrrWrr)rrs r6 logcumsumexprs/3 """  D0G H H HHr8c|j}t|}||z }tt|}dt|D} |D]} d| | <gg} } |D]5} | | s| |  | | 6| | z}t| } ||d| }|fdd||| dz}||}dgt|j|dz}| |}| d}||d<t|}tt|D]}|||||dz<| |tj d t|D}d}|dz }|dkr=||dz|||<||||z}|dz}|dk=t||D]&}|d||z z|||<'|||||S) Ncg|]}dSFr2rNrTs r6rPz_exec_fft..s555A%555r8Tc|SrBr2)r; self_stridess r6rfz_exec_fft.. s <?r8keyreverserrr1rcg|]}dSrr2rs r6rPz_exec_fft..s***1***r8)rlenlistrangeappendstridesortpermuterreshaperresize_rWr as_strided_storage_offset)outr out_sizesrforwardr signal_ndim batch_dims dim_permuteis_transformed_dimdleftright batch_endtmpinput batched_sizes batch_sizebatched_out_sizesi out_strides batch_numelrs @r6 _exec_fftrs 9Dc((K #JuT{{##K55t555 %% $1b%D !!$  KKNNNN LLOOOO,KD I;;==L jyj !CHH****DH999 IJJ//K LL % %ED4 JKK 8999M MM- ( (EAJ!M!]++ 3s88__55#,SV#4!a%  KK!1HKIII+*eDkk***KKQA q&&&1CJJqMM&A KN#yQ00  Q q&&:t $ $GG&)jja*n1E&F&F KN##OOI{C,>,>,@,@AAA Jr8rr exclude_lastct|}||dt|t|z fd|S)Nc|SrBr2)rrs r6rfz_sort_dims..,s l1or8)r)rrrintr)rrr sorted_dimsrs @r6 _sort_dimsr(sls))K;;==L6#k""S%6%6667<< % % % %= r8c$tj|jj|s|St ||}||}t|||||S)Nr) rWrhr_ is_complexclonerrrr)rr normalizationrrrs r6 meta_fft_c2cr3sx L&''' zz||T3''K .. % %C S$ [' J J JJr8ct|tks+t|dkr|ddkr|ddkrdSdS)Nrr1FT)rcufft_max_ndimrs r6use_optimized_cufft_pathrBsD 3xx.  SXX]]s1v{{s1vQR{{utr8c6tj|jjt |}t |}|d}||dzdz}t |}|||<|r|||<t |dkst |dkr||tj |j} |} t |dkr%t|rt| | ||dnt|dkr|n|} t| | | |gdt|dkr.||tj |j} |dd} | r| | } } | | fd d tt t| } | t| | z d}t| | ||d| dt| | z } | |sB| |||kr#| |tj | } | S||tj |jS) Nrrr1cudaxpurTrc|SrBr2)rstridess r6rfzmeta_fft_r2c..ps '!*r8rr)rWrhr_is_floating_pointrr device_hintrrQrrrrrrminrrr)rrronesided input_sizesrlast_dimlast_dim_halfsizeonesided_sizesoutputworking_tensor target_sizesrmax_dims last_dimsrs @r6 meta_fft_r2cr Is L-...tyy{{##K[!!I2wH#H-2Q6+&&N0N80/ (4F""k$&7&75&@&@ U>tzJJ    t   & &+CC+H+H & fnid K K K K K),CA 99>L fnlXJPT U U U U3xx!||!%U%Ftz%R%R"0"" crc(K I)7(//11  ,,,,d!~s;/?/?@@'K(8(88(C(E(EF NNIt**GC ,<,tzJJ   r8) generatorcHt|tj|gSrB)rrWSize)nr rs r6 meta_randpermrs S%*aS// 2 22r8r_rrrc4tj|||||SNrrWr)r r_rrrs r6meta_randperm_defaultrs( ; vf   r8cvdtjkfdtj|||||S)NrcddSNz:random_ expects 'from' to be less than 'to', but got from=z >= to=r2highlowsr6rfzmeta_randint.._S__Y]__r8rrWrhr)rrr_rrrrs` @r6 meta_randintrs[ C L s _____ ; E&J   r8crtjkfdtj|||||S)NcddSrr2rsr6rfz"meta_randint_low..rr8rr)rrrr_rrrs`` r6meta_randint_lowrsV L s _____ ; E&J   r8c4tj|||||Srr)rr_rrrs r6meta_rand_defaultr!s( ; E&J   r8rlastdimctj|jjt |dkrt |}|||d<||t|j}t|r2t|| tj ||dSt|dkrt||ddd|}n | tj }t||||dgdS|}t|dkr'|dd}t|||d}|dd}t |}|||d<||t|j} t| |||dS) NrrrrFrr1r)rWrhr_rrrrrrarrrrrr) rrrr"rrtemprc2c_dimsrs r6 meta_fft_c2rr&s L&'''4F""%% $ #b' 1L1LMM #C ( ( P )@ AA  3xx!||#D#crc(Aw??zz0GzHHVT9s2wiOOO O s88a<<3B3xH xNNNEbcc(C&& $ #b'nnYodj.I.InJJeYUCCCCr8cddlm}||s'tj|dkrt dt |t rr|||}||kr2tj |||S)Nr)free_unbacked_symbolsr1zQmore than one element of the written-to tensor refers to a single memory location) %torch.fx.experimental.symbolic_shapesr(rW_debug_has_internal_overlap RuntimeErrorrqr torr+ expand_copydefault)rsrc non_blockingr( intermediates r6 meta_copy_r2s LKKKKK " !$ ' ' ,1,Md,S,SWX,X,X _   #v@vvdL11 99;;,++-- - -   $ $\499;; ? ? ? Kr8c<t|}t|}||krdn||||z}||d|||||fSr0)rrrrinsert)tensorr result_sizesresult_strides new_strides r6inferUnsqueezeGeometryr9 s &&L&--//**NVZZ\\))|C/@>RUCV/VJQ#z***  ''r8ct||dz}t||\}}||||Sr0)rrr9r)rrg_sizes g_stridess r6meta_unsqueeze_r=sL dhhjj1n - -C/c::GYWi((( Kr8rweight_metabias_activation_opt out_dtypeclt|j}|6|d|dks Jd|d|ddz ksJ|d|d<t|jdks Jdd|df}|/|jt jkr|t jks Jd||||jn| ||}|S) Nrzoutput size mismatchr1rrz*we can only handle the squashed input case9out_dtype is only supported for i8i8->i32 linear operatorr) rrrrr_rWint8int32r as_strided) rr>r?r@rArB output_sizestransposed_stridesrs r6meta_sparse_structured_linearrJ!s: $$L {{1~~1---/E--- ;;q>>UZZ^^a/ / / / /{{1~~L u{  q "N UZZ]]+{ej((Y%+-E-E-E G.F-E-E__&.ekkIj122 Mr8mat1 mat1_metamat2ct|jdksJt|jdksJt|jdksJ|d|ddz ksJ|d|dg}|/|jtjkr|tjks Jd||||jn|}|S)Nrr1rrDrrrrr_rWrErFr)rKrLrMrBrHrs r6meta_sparse_structured_mmrPCs tz??a     y  1 $ $ $ $ tz??a     99Q<<499Q< z matmul!r) r_rWfloat32float16bfloat16rE float8_e4m3fnrr is_contiguousrrFr) rUrVr@rQrBrWrXrYrZis_8bit_input_typer m output_shapes r6meta__cslt_sparse_mmrds =          E       . . .0Q . . . w}   " " "$M " " "%+ E>+ & &&r8)lengthsroffsetsaxisunsafeinitialdatarrrsrtruc|tdfd}|||jS|,|jdd|jddz fz} || Std)Nz?segment_reduce(): indices based reduction is not supported yet.cxtj|jdzdzjdtjS)Nr1rr_rr)rWrrr_r) lengths_shapertrws r6segment_reduce_lengths_tensorz:meta_segment_reduce..segment_reduce_lengths_tensorsA{ DJtaxzz2 2*1     r8rr1z>"  r8ctj|j|f}t|||}||||t jfSNrrQreduction_dimsr_compute_reduction_shaperrWrrrkeepdimrcs r6 meta_max_dimrW  tzC6 2 2C+D#w??L |$$ |5:66 r8c,|dSrrrs r6meta_minrrr8ctj|j|f}t|||}||||t jfSrrrs r6 meta_min_dimr rr8c|rt|j}nt|tj\}}t j||SNrLr)rrr_rr INT_TO_FLOATrWr)rrOrTs r6 meta_anglers^  / ;; ,  ? L   <  D 5 5 55r8ctj|||j|tj|SrB)rW_resize_output_rrcopy_angle)rrs r6meta_angle_outr"s= #tyy{{DK888 99U[&& ' ''r8cdSrBr2)vals r6 assert_asyncr( Fr8cdSrBr2)r assert_msgs r6assert_async_metar-rr8cdSrBr2)ss r6 print_metar2rr8r_rrrrc.tjddS)Nrrrrrs r6make_dep_tokenr7s ;q ( ( ((r8cddlm}t|ttfrt d||||dS)Nr)constrain_range'Constraining SymFloat or Symbool is nyirmax)r)rrqr r ValueError)rrrrs r6sym_constrain_rangerCs[FEEEEE$7+,,DBCCCODcs++++++r8c@t||||SNr)r+rrrr dep_tokens r6functional_sym_constrain_rangerMs#Ts444 r8cXddlm}||tj|dkdSt |t t frtdt|tur6|tj||k|tj||kdS||||dS)Nr)_constrain_range_for_sizerr) r)rrWrhrqr r rrzr)rrrrs r6sym_constrain_range_for_sizerSsPOOOOO {s{ TQY$7+,,DBCCC DzzS ? L % % % ? L % % %d555555r8c@t||||Sr)r+rrs r6'functional_sym_constrain_range_for_sizergs#%%d%=== r8c|SrBr2)rrrs r6functional_assert_async_metarms r8f_namec|dks J|d|d|dks9J|d|dd|dddS)Nrz3: The input tensor must have at least 2 dimensions.rz5: A must be batches of square matrices, but they are  by matrices)rr)rrs r6rrss 88::??? FFF ?? 99R==DIIbMM ) ) ) ss RT ss[_[d[deg[h[hsss * ) ) ) )r8AnamectjjjkfdtjjjkfdtjddkfdtjddkfddS)Nc*djdjdS)Nz:Expected b and A to be on the same device, but found b on z and A on instead.rrrsr6rfz(linearSolveCheckInputs..s0 :{ : :&'h : : :r8c*djdjdS)Nz=Expected b and A to have the same dtype, but found b of type z and A of type rrrsr6rfz(linearSolveCheckInputs..s0 =z = =*+' = = =r8rrcbdddddS)Nz3A must be batches of square matrices, but they are rrrrrrsr6rfz(linearSolveCheckInputs..sF BFF2JJ B B,-FF2JJ B B Br8c ddddddddd S)NzIncompatible matrix sizes for z: each A matrix is rrz but each b matrix is rr)rrrsr6rfz(linearSolveCheckInputs..s HT H H H H)* H H%)YYr]] H H8< "  H Hr8)rWrhrr_r)rrrs```r6linearSolveCheckInputsrs L qx      L ag      L r affRjj      L r diimm#      r8tallow_low_precision_dtypesc,|jtj|p|fd|sItjtjtjtjtjfvfddSdS)NcdS)Nz<: Expected a floating point or complex tensor as input. Got r2r_rsr6rfz(checkFloatingOrComplex..s6^^W\^^r8cdS)Nz*: Low precision dtypes not supported. Got r2rsr6rfz(checkFloatingOrComplex..svPPPPr8) r_rWrhrrr[r]rZr\)rrrr_s ` @r6rrs GE L /^^^^^ &  ek5<u}M M P P P P P       r8arg_namechtj|dkfddS)NrcddS)Nz: The input tensor z! must have at least 2 dimensions.r2)rrsr6rfzcheckIsMatrix..s6YYhYYYr8)rWrhr)rrrs ``r6 checkIsMatrixrs@ L 1 YYYYYr8Brc4tttjr,ddkn+ddkfddS)Nrrc drdndddddddddd S) Nz2: Incompatible shapes of A and B for the equation zAX = BzXA = Bz (rr;rrryr)rrrrsr6rfz#checkInputsSolver..s J J-xxX J J J JffRjj J J/0vvbzz J J<=FF2JJ J J Jr8)rrrWrhr)rrrrs````r6checkInputsSolverrsa   !V L$(Fr affRjj  affRjjAFF2JJ.F       r8resultfn_name result_namec`tjjjkfddS)Nc 8dddjdj S)Nz : Expected z5 and input tensors to be on the same device, but got z on z and input on r)rrrrsr6rfz!checkSameDevice..sV L L; L L L L &  L L=B\ L Lr8)rWrhr)rrrrs````r6checkSameDevicersU  L %       r8UPLOc}tjtdko |dkp|dkfddS)Nr1ULcdS)Nz1Expected UPLO argument to be 'L' or 'U', but got r2)rsr6rfzcheckUplo..sJDJJr8)upperrWrhr)rUPLO_uppercases` r6 checkUplors[ZZ\\N L D QKNc1J^s5JJJJJr8 eigenvalues eigenvectorsr compute_vct|dt|t|j}|r;||}||t |dn|dg}|||t|j }||fS)Nz linalg.eighF row_majorrr) rrrrrrrpoprar_)rrrrvecsvalss r6meta__linalg_eighrsa''' dOOO MME {{5!!  ;EU S S STTTT{{A3 IIKKK ;;uOAG$<$<; = =D :r8ct|dtj|jr|jntj|j}||jdd|S)Nzlinalg.eigvalsrr)rrQrr_rrr)r complex_dtypes r6meta__linalg_eigvalsrsje-...  !%+ . . <  .u{ ; ; ??5;ss+=? A AAr8ct|dtj|jr|jntj|j}||jdd|}||j|}t|dk}||jt|j|||fS)Nz linalg.eigrrrr) rrQrr_rrrrrr)rrrvectorsis_cudas r6meta_linalg_eigrse\***  !%+ . . <  .u{ ; ; __U["-]_ C CFooeko??G%  F*G  0PPP 7?r8r/ct|jtjddS)Nrrr)mTrrWr transpose)r/s r6cloneBatchedColumnMajorrs+ 6< > H HR P PPr8rc t|SrB)r)rrrs r6_cholesky_solve_helperrs #4 ( ((r8ctjjdkfdtjjdkfdtd\}}t |||S)NrcdjdS)Nz-b should have at least 2 dimensions, but has  dimensions insteadrrsr6rfz cholesky_solve.. s^ ^^^r8cdjdS)Nz-u should have at least 2 dimensions, but has rrrsr6rfz cholesky_solve..$s[[[[r8cholesky_solve)rWrhr!_linalg_broadcast_batch_dims_namer)rrrself_broadcasted A_broadcasteds`` r6rrs L Q^^^^ L ! [[[['H a!''#m ""2M5 I IIr8c|dkr tj|tjSt |dt |S)Nrrcholesky)rrWrlegacy_contiguous_formatrrrrs r6r r ,sN zz||qE4RSSSSdJ''' "4 ( ((r8c@t|dt|S)Ncholesky_inverse)rrr s r6r r 5s#d./// "4 ( ((r8 check_errorscDt|dt|d|j}t|}t |d}||}|||||d|dz tj}||fS)Nzlinalg.choleskyFrrr) rrrrrrrrWrF)rrrA_shaper L_stridesrinfoss r6linalg_cholesky_exr=sa*+++1/000gG w<.VZr8rrcdS)Nzbtorch.linalg.householder_product: input.shape[-2] must be greater than or equal to input.shape[-1]r2r2r8r6rfz,linalg_householder_product..Zstr8cdS)Nz`torch.linalg.householder_product: input.shape[-1] must be greater than or equal to tau.shape[-1]r2r2r8r6rfz,linalg_householder_product..^srr8r1c(djdjS)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 rrrsr6rfz,linalg_householder_product..c0 \), \ \OTz \ \r8cdS)Nzltorch.linalg.householder_product: Expected batch dimensions of tau to be equal to input.shape[:-2], but got r2actual_batch_tau_shapesr6rfz,linalg_householder_product..mO6LOOr8c(djdjS)Nz,torch.linalg.householder_product: tau dtype z does not match input dtype rrsr6rfz,linalg_householder_product..us* 939 9 9+0; 9 9r8z torch.linalg.householder_productrFrrrr_r) rWrhrrrr_r empty_stridedrr)rrexpected_batch_tau_shaper s`` @r6linalg_householder_productr&Os  L aZZ L 2%**R..(tt L 2#((2,,&rr  L SX"      zA~~#(;ss#3 !$3B3 "&> >        L U[      6UEJJJ   [*5;%HHHk|    r8c<t|dt|dd||j}||jt |jd||jddt j}||fS)Nz linalg.inv_exF)rrrrrrrrrrrWrF)rrrrs r6linalg_inv_ex_metar)sa)))1o%PPPP AGAMM!'6qw%PPPQQQ KK EKK 8 8E e8Or8LDpivotsinfo) hermitianrr-czt|dt|dtj|jt |jd|j|j}||jddtj }||jddtj }|||fS)Nztorch.linalg.ldl_factor_exFrr#rrr) rrrWr$rrr_rrr)rr-rr*r+r,s r6linalg_ldl_factor_ex_metar/sd89994!=>>>   Z*4:GGGj{    B ^^DJssO59^ = =F >>$*SbS/> ; ;D vt r8)r-c@tdtdtdtjjdkfdjdd}tj|jkfdtjtjj fdtjj j kfdt\}}tj |t|d j j S) Nztorch.linalg.ldl_solvercdjdS)NzMtorch.linalg.ldl_solve: Expected B to have at least 2 dimensions, but it has rr)rsr6rfz'linalg_ldl_solve_meta..! 6& 6 6 6r8rcdjdS)Nzjtorch.linalg.ldl_solve: Expected LD.shape[:-1] and pivots.shape to be the same, but got pivots with shape  insteadrr+sr6rfz'linalg_ldl_solve_meta..% @)/ @ @ @r8cdjS)Nz.s]v|]]r8c(djdjS)Nz!torch.linalg.ldl_solve: LD dtype z does not match b dtype r)rr*sr6rfz'linalg_ldl_solve_meta..s_BH__VWV]__r8Frr#)rrrrWrhrrrQis_integer_dtyper__linalg_broadcast_batch_dimsr$rr)r*r+rr-expected_pivots_shapeB_broadcast_sizerTs``` r6linalg_ldl_solve_metar>scb2333278881b":;;; L !     HSbSM L-     L v|,,]]]] L AG_____7q"==a   *+;uMMMgx    r8Pr)pivotr@ctjjdkfdtj}|d}|d}t ||}||d<|r|}ndg}||d<|}||d<||d<|}|||fS)NrcdjdS)Nz@linalg.lu: Expected tensor with 2 or more dimensions. Got size: r4r5rsr6rfz linalg_lu_meta..sdSTSZdddr8rrr)rWrhrrrrr) rr@sizesrbr kr?rrs ` r6linalg_lu_metarEs L ! dddd MME b A b A Aq AE"I  KK   KK  E"I EAE"IE"I EA a7Nr8LU)r@rctjjdkfdtj}|d}|d}tj|t |djj}| t|||d< |tj }|  |tj }|||fS) NrcdjdS)NzFtorch.lu_factor: Expected tensor with 2 or more dimensions. Got size: r4r5rsr6rfz*linalg_lu_factor_ex_meta..sjYZY`jjjr8rrFrr#r) rWrhrrrr$rr_rrrrr) rr@rrCrbr rFr+r,s ` r6linalg_lu_factor_ex_metarIs L ! jjjj MME b A b A   *5EBBBgx    B IIKKKAq E"I [[ei[ 0 0F IIKKK ;;uEI; . .D vt r8)radjointrJctdtjjjkfdtjjtjkdt dt |dtjddkdtjjddjkfdt\}}tj |t|| jj }| d kr*|s(|r|}|S) Nztorch.linalg.lu_solvec*djdjdS)NzPlinalg.lu_solve: Expected LU and B to have the same dtype, but found LU of type  and B of type r4r)rrFsr6rfz&linalg_lu_solve_meta.. s6 O$&H O O=>W O O Or8cdS)NzElinalg.lu_solve: pivots should be a Tensor of scalar type torch.int32r2r2r8r6rfz&linalg_lu_solve_meta..'sWr8zlinalg.lu_solvercdS)NzYlinalg.lu_solve: Number of pivots per batch should be same as the dimension of the matrixr2r2r8r6rfz&linalg_lu_solve_meta../skr8cdjdS)Nzclinalg.lu_solve: Expected LU.shape[:-1] and pivots.shape to be the same, but got pivots with shape r4r5r6sr6rfz&linalg_lu_solve_meta..5r7r8rr#r)rrWrhr_rrrrrr;r$rrrrconj)rFr+rrrJr=rTrs``` r6linalg_lu_solve_metarRs26777 L AG      L  !WW b1222b!T#4555 L  v{{2&kk  L " %    7q"==a  *+;4xPPPgx F||~~4      #[[]]F Mr8 unpack_data unpack_pivotscDtjjdkfd|r)tj|jtjkdt j}|d}|d}t||}||d<|r|}ndg}|r:||d<|} ||d<||d<|} n,dg} dg} || | fS)NrcdjdS)NzFtorch.lu_unpack: Expected tensor with 2 or more dimensions. Got size: r4r5)rFsr6rfz lu_unpack_meta..UskY[Yakkkr8cdS)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_factorr2r2r8r6rfz lu_unpack_meta..Zs pr8rrr) rWrhrr_rFrrrr) rFr+rSrTrCrbr rDr?rrs ` r6lu_unpack_metarXKs? L 1 kkkk  LEK '      NNE b A b A Aq AE"I LL   LL!  b LL  b b LL   LL!   LL!   a7Nr8modec~dkrd}d}n.dkrd}d}n#dkrd}d}ntjdfd||fS)NreducedTcompleteFrcddS)Nzqr received unrecognized mode 'z=' but expected one of 'reduced' (default), 'r', or 'complete'r2)rYsr6rfz _parse_qr_mode..s"O$OOOr8rWrh)rY compute_qr[s` r6_parse_qr_moderaus~ y                g r8QRr[cTt|dt|dt|\}}|jd}|jd}t ||}|rXt |j}|r|n||d<||}||t|dn|dg}t |j} |s|s|n|| d<|| } | | t| d|| fS)Nz linalg.qrrrFrr) rrrarrrrrr) rrYr` reduced_moderbr rDQ_shaperbR_shapercs r6linalg_qr_metarhs+![!!!1k***,T22I|  A  A Aq Aqw--'.aaQ KK  g:7eTTTUUUU KK  17mmG#;9;!!!GBK GAMM'6w%PPPQQQ a4Kr8sign logabsdetct|dt|dd|j}||dd}||ddt |j}t j|t|d|j|j }||ddt j }||||fS)Nzlinalg.slogdetFrrr#r) rrrrrar_rWr$rrrF)rrrirjrFr+s r6_linalg_slogdetrlsa)***1.666 GE ;;uSbSz " "D E#2#Joag.F.F GGI   *5%88gx    B [[ss5;[ 7 7F B &&r8 full_matrices compute_uvdriverct|dt|dt|jdd}|jd}|jd}t ||}|r|||r|n|gz}||} | |t|d||r|n||gz} || } t|dk} | | t| | n,|dg} |dg} |||gzt|j } | | | fS) Nz linalg.svdrrFrrrr) rrrrrrrrrrar_)rrmrnrorrbr rDU_shaperV_shapeVrSs r6_linalg_svd_metarusl!\"""1l+++agcrcl##J  A  A Aq A #<111== KK  g:7eTTTUUU] 91== KK  a..F* g:7gVVVWWWW KK   KK   J!$OAG,D,D EEA a7Nr8arg1arg2cV|jdd}|jdd}t||}t|}||d|dgz }t|}||d|dgz }||fS)Nrr)rr#rr)rvrwarg1_batch_sizesarg2_batch_sizesexpand_batch_portionarg1_expand_sizearg2_expand_sizes r6r;r;s z#2#z#2#,-=?OPP0112 " 660112 " 66 - --r8c|rt|||t||\}}||jkr|n||}||jkr|n||}||fSrB)rr;rexpand)rvrwrr|r}arg1_broadcastedarg2_broadcasteds r6rrs  1tT4000)EdD)Q)Q&&!DJ..DKK@P4Q4Q!DJ..DKK@P4Q4Q - --r8rcv|jdd}|jdkp|jdz |jko |j|k}|S)Nrr1)rr)rrexpected_batched_rhs_shape vector_cases r6linalg_solve_is_vector_rhsrsK!&SbS!1*/ Q%*$R8R)R r8)rrrrFr+r,ctdtjjjkfdt }|rdn} t | |dt| \} } tj|p| d|r | ddn| } tj| t| | jj } j }tj|t|djj } |ddtj } |ddtj }||||f}| |||f}td |Drjt||D]Y\}}t!||j ||j |t'||d Z|S) Nz linalg.solvec*djdjdS)NzKlinalg.solve: Expected A and B to have the same dtype, but found A of type rMr4r)rrsr6rfz"_linalg_solve_ex..s0 9w 9 9'(w 9 9 9r8rcdS)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)r2r2r8r6rfz"_linalg_solve_ex..'s  Kr8r#Frrc3K|]}|duV dSrBr2rNr;s r6rtz#_linalg_solve_ex..>s& & &Q1D= & & & & & &r8) copy_fromcopy_to exact_dtype)rrWrhr_r unsqueezerr;r$rrrrrFallziprrrr!)rrrrrrFr+r,rB_ B_broad_shaperT result_shaperesult_rLU_pivots_info_rresr]os`` r6_linalg_solve_exr s51n--- L 17     -Q22K' .RQBaT>2223B::M1 L K  *5G="%%-L! *.ms! 9) 9 9 9r8cdjdS)NzMtorch.triangular_solve: Expected A to have at least 2 dimensions, but it has rrrsr6rfz'triangular_solve_meta..tr2r8triangular_solveFrr#rcdS)Nz+triangular_solve: Got an unexpected layout.r2r2r8r6rfz'triangular_solve_meta..s$Qr8)rWrhrrrstridedr;r$rr_r sparse_csr sparse_bsrrr) rrrrrself_broadcast_sizeA_broadcast_sizesolutioncloned_coefficients `` r6triangular_solve_metarbs L Q     L !     4$6777x5=  0LTST0U0U--&$./BeTTT*;     #0!./?5QQQ'8     U% % %U5E)E)E#D))!^^QC00 UQQRRR ' ''r8c~t|dt|d||jdd}||j}||jt |jd||jddt j}|||fS)Nz linalg.detrFrrrr()rdetrFr+s r6_linalg_det_metarsa&&&1l+++ ++agcrcl # #C QW  BNN1775QQQRRR [["U[[ 9 9F F?r8ctjjdkdtjjdkd|rdndtjjjdkfdtjjjdkfdtjjdjdkdtjjjz d kfd tjjjkfd jdkrejdd}jddtj|kfd jddtj|kfd tjjjkfdtjjjkfdt ddt ddtjjtjdjjS)NrcdS)Nz3torch.ormqr: input must have at least 2 dimensions.r2r2r8r6rfzormqr..!Vr8cdS)Nz3torch.ormqr: other must have at least 2 dimensions.r2r2r8r6rfzormqr..rr8rrcddS)Ntorch.ormqr: other.shape[z0] must be greater than or equal to tau.shape[-1]r2left_size_conditionsr6rfzormqr..sq,?qqqr8cddS)Nrz"] must be equal to input.shape[-2]r2rsr6rfzormqr..sc,?cccr8cdS)NzHtorch.ormqr: tau.shape[-1] must be less than or equal to input.shape[-1]r2r2r8r6rfzormqr..rr8r1c(djdjS)Nz[torch.ormqr: Expected tau to have one dimension less than input, but got tau.ndim equal to rrrsr6rfzormqr..rr8c(djdjS)Nzhtorch.ormqr: Expected other to have the same number of dimensions as input, but got other.ndim equal to rrrrsr6rfzormqr..s1 `+0: ` `SXS] ` `r8cdS)NzWtorch.ormqr: Expected batch dimensions of tau to be equal to input.shape[:-2], but got r2rsr6rfzormqr..r!r8cdS)NzYtorch.ormqr: Expected batch dimensions of other to be equal to input.shape[:-2], but got r2)actual_batch_other_shapesr6rfzormqr..sQ6NQQr8c(djdjS)NzPtorch.ormqr: Expected input and tau to have the same dtype, but input has dtype z and tau has dtype rrsr6rfzormqr..s0 O#(; O OCF9 O Or8c(djdjS)NzRtorch.ormqr: Expected input and other to have the same dtype, but input has dtype z and other has dtype rrsr6rfzormqr..s0 S#(; S SEJ[ S Sr8z torch.ormqrrrFrr#) rWrhrrr_rr$rr) rrrrrexpected_batch_shaperr rs ``` @@@r6ormqrrs L aVV L aVV!%,""" L '(CIbM9qqqq L '(EKO;cccc  L " R(ZZ  L SX"      L ej       zA~~${3B3/!$3B3 "&: :       $);ss#3  $(< <        L U[       L u{"     M3u555M5%999   [*5;%HHHk|    r8ctjtdzkfdj}|dzk}|}| }|r/t d|D]}|o|dk}n-t |D]}|o|dk}tj|p|fddS)Nrc4ddzdtS)Nzpadding size is expected to be rz , but got: r)rpaddingsr6rfz,_padding_check_valid_input..s"T!c'TTc'llTTr8r1rc0ddzddzdjS)N Expected r1zD or rzcD (batch mode) tensor with possibly 0 batch size and other non-zero dimensions for input, but got: r5)rrsr6rfz,_padding_check_valid_input..sC Oa O OcAg O OAF O Or8)rWrhrrrr)rrr input_dim is_batch_modevalid_batch_modevalid_non_batch_moders``` r6_padding_check_valid_inputrs! L G CTTTTT  I#'*M$,,Oq)$$ G GA/FEJJqMMQ4F   Gy!! O OA#7#NEJJqMMQ.1Mc%*cc.3ccDIccUZU`ccr8cddS)Nz input (W: z%) is too small. Calculated output W: r2)input_woutput_wsr6rfz_pad1d_common..9sUWUU8UUr8r)rrrrWrhr) rr is_reflection dim_planenbatchnplanerrrrrs ` @@@@@r6 _pad1d_commonrsII E F zQA  Q ug15555LE5 ZZ " "FjjG&H  GO /           LA UUUUU  zQ12229:::r8c&t||dSNTr)rrrs r6meta_reflection_pad1drB t < < <djdS)Nz)"replication_pad1d" not implemented for ''r___str__rsr6rfz(meta_replication_pad1d..M!Xek>Q>Q>S>SXXXr8Fr)rWrhr_boolrrs` r6meta_replication_pad1drHH L uz!XXXX u = = ==r8cd|s'tjt|dkdjdkrdz |\}|zz|r%tj|ko|kfdtjkfdjS)Nr1rcdS)Nz padding size is expected to be 2r2r2r8r6rfz(_pad1d_backward_common..Us0Rr8r.c*ddddjSrr5rsr6rfz(_pad1d_backward_common..brr8c:ddSNz(grad_output width unexpected. Expected: , Got: rr grad_outputrsr6rfz(_pad1d_backward_common..j'e8eeKL\L\]bLcLceer8rWrhrrrrr) rrrrrrrrrs `` @@@@r6_pad1d_backward_commonrRs E T S\\Q&(R(RSSS zQ  LE5jjG&H  GO /           LK$$U+++eeeeee ??5; ' ''r8 grad_inputc(t|||dSrrrrrs r6meta_reflection_pad1d_backwardrps "+ugT R R RRr8c(t|||dS)NFrrrs r6meta_replication_pad1d_backwardrvs "+ugU S S SSr8c ddd}d}t|dj}|dkr$d}dz dz |dz }|\ |}  z z zz |rJtj ko kfdtj ko k fdtj dkp dk fdjd kr| fS|| fS) Nrr1rrc*ddddjSrr5rsr6rfz_pad2d_common..rr8c*ddddjSNzcArgument #6: Padding size should be less than the corresponding input dimension, but got: padding (rxrrr5dim_hrpad_bpad_tsr6rfz_pad2d_common..rr8c ddddS)Nz input (H:  W: z%) is too small. Calculated output H: r2)input_hroutput_hrsr6rfz_pad2d_common..s? = = =g = =$, = =2: = =r8r.rrrrWrhr)rrr dim_slicesrrrr rrrrrrrrrs` @@@@@@@@@@r6 _pad2d_commonr|s E EJ Fug15555 :D qyyA    a !(E5% ZZ # #FjjGjjG&H&H  GO /            GO /           LA &Q        zQ(;<<<(CDDDr8c&t||dSr)rrs r6meta_reflection_pad2drrr8c~tjjtjkfdt |dS)Nc>djdS)Nz)"replication_pad2d" not implemented for 'rrrsr6rfz(meta_replication_pad2d..rr8Fr)rWrhr_rrrs` r6meta_replication_pad2drrr8cZtj|}tj|}||fSrBrWr)grad_wsaved_vsaved_g saved_normsrgrad_vgrad_gs r6meta_weight_norm_backwardr$s. g & &F  g & &F 6>r8c d d d}|j}|dkr dz dz |dz }|\}}}}| } | } | |z|z | |z|ztj k fdtj  k fd||jS)Nrr1rrc:ddSrrrsr6rfz%meta_pad2d_backward..rr8c:ddSNz)grad_output height unexpected. Expected: rrr rrsr6rfz%meta_pad2d_backward..'fHff[M]M]^cMdMdffr8)rrrWrhrr)rrrrrerrrrrrr rrrs` @@@@r6meta_pad2d_backwardr+s E EIJ xxzzQ    Q !(E5%GG&H&H LK$$U+++eeeeee LK$$U+++ffffff >>$* % %%r8c R d ddd}t|djdk}|r)d} dz dz dz |dz }|\|}    zz zz zz|rotj ko k fdtj ko kfdtj ko kfd tjdkp dkp dk fd |r|| fS| fS) Nr.rr1rrc*ddddjSrr5rsr6rfz_pad3d_common.. rr8c*ddddjSr r5r sr6rfz_pad3d_common..rr8c*ddddjS)NzcArgument #8: Padding size should be less than the corresponding input dimension, but got: padding (rxrrr5)dim_drpad_bkpad_fsr6rfz_pad3d_common..sMd%*dd.4ddEJddV[Vaddr8c ,dddddd S)Nz input (D:  H: rz%) is too small. Calculated output D: r2)input_drroutput_drrsr6rfz_pad3d_common..#se K K Kg K K7 K K$, K K2: K K@H K Kr8r)rrrr batch_moderrr1r rr6rrr7rrrr2r3rrrs` @@@@@@@@@@@@@@@r6 _pad3d_commonr9s E E EIug15555qJA      Q 07-E5%v ZZ " "FjjGjjGjjG'H&H&H  GO /            GO /            GO 0 0           LA 7Q7(a-         G(HMNNN(HEFFFr8c&t||dSr)r9rs r6meta_reflection_pad3dr;/rr8c~tjjtjkfdt |dS)Nc>djdS)Nz)"replication_pad3d" not implemented for 'rrrsr6rfz(meta_replication_pad3d..:rr8Fr)rWrhr_rr9rs` r6meta_replication_pad3dr>5rr8c tjt|dkd|jdksJj|jksJdd d |jdkrdz  dz dz |\}}}}}}| } | } |} | |z|z| |z|z| |z|ztjkfdtj k fdtj k fd ||jS) NcdS)Nz padding size is expected to be 6r2r2r8r6rfz%meta_pad3d_backward..Is,Nr8r.rr1r-c:ddSrrrsr6rfz%meta_pad3d_backward..arr8c:ddSr(rr)sr6rfz%meta_pad3d_backward..er*r8c:ddS)Nz(grad_output depth unexpected. Expected: rr)r1rr7sr6rfz%meta_pad3d_backward..irr8r)rrrrrrrr3r2r6rrr1r rr7rrs` @@@@@@r6meta_pad3d_backwardrE?s LW"$N$NOOO :>>>>  uz ) ) ) ) E E E zQ      07-E5%vjjGjjGjjG'H&H&H LK$$U+++eeeeee LK$$U+++ffffff LK$$U+++eeeeee ??5; ' ''r8rpcjtj|d|d}|dkr4|dgtjS|||dz zdzftjS)NcdS)Nz(_pdist_forward requires contiguous inputr2r2r8r6rfz%meta__pdist_forward..ss&Pr8rr1rr)rWrhr`rrr,r )rrFr s r6meta__pdist_forwardrIos L PP ! AAvv~~qc""%%E4R%SSS~~qAE{a/1225586   r8gradpdistctj|dtj|dtj|tjS)NcdS)Nz._pdist_backward requires self to be contiguousr2r2r8r6rfz&meta__pdist_backward..s&Vr8cdS)Nz/_pdist_backward requires pdist to be contiguousr2r2r8r6rfz&meta__pdist_backward..s'Xr8r)rWrhr`rr )rJrrFrKs r6meta__pdist_backwardrO~ss L VV L XX  D0N O O OOr8)rRrQc J ddlm}m}d}d}d} |t j|j||| fr||| ft j dkdt j dkdtj s7t jj j cxko j kncfdj} j | d | d t j d ko d k fd  S) Nr) guard_or_truesym_eqr1rr.cdSNzbatch1 must be a 3D tensorr2r2r8r6rfzmeta_baddbmm..,Hr8cdSNzbatch2 must be a 3D tensorr2r2r8r6rfzmeta_baddbmm..rUr8c8djdjdjS)Nz+Input dtypes must be the same, got: input: z , batch1: z , batch2: r)batch1batch2rsr6rfzmeta_baddbmm..s*~$*~~X^Xd~~pvp|~~r8c :ddddddd SNz@Expected size for first two dimensions of batch2 tensor to be: [rxz ] but got: [rr1].r2 batch2_sizesbscontraction_sizesr6rfzmeta_baddbmm..sY [ [ [& [ [4@O [ [GSTU [ [ [r8)r)rQrRrrWsym_notrrrhr exp_config&skip_dtype_check_in_meta_registrationsr_r)rrYrZrRrQrQrRdim1dim2dim3 batch1_sizesr_r`ras``` @@@r6 meta_baddbmmrisLKKKKKKK ;;q>>D ;;q>>D ;;q>>D}U]66$*tT46H#I#IJJKK/{{D$-.. L"$H$HIII L"$H$HIII  <  J&, 6 6 6 6&, 6 6 6 6 ~ ~ ~ ~ ~ ~   >$))++ & &&r8cBtj|tjSrrirr s r6meta_bernoullirls  D0G H H HHr8?c|SrBr2rrFr s r6meta_bernoulli_rp Kr8cBtj|tjSrriros r6meta_bernoulli_prss  D0G H H HHr8c*tj|SrBrrks r6 meta_poissonru  D ! !!r8c tj| |kdtj|tj} tj|| fS)NcdS)NzJError in fused_moving_avg_obs_fake_quant_cpu: ch_axis must be < self.dim()r2r2r8r6rfz6meta__fused_moving_avg_obs_fq_helper..\r8r)rWrhrrr)r observer_on fake_quant_on running_min running_maxscale zero_pointaveraging_const quant_min quant_maxch_axisper_row_fake_quantsymmetric_quantmasks r6$meta__fused_moving_avg_obs_fq_helperrs]  L$((**\\  D 3 3 3D  T " "D ))r8ctj|dkdtj|dkd|j\|j\tjkfd|Ntj||jkp.|tjko|jtjtjfvd||jn|}|f|S)NrcdS)Nz a must be 2Dr2r2r8r6rfzmeta_mm..~r8cdS)Nz b must be 2Dr2r2r8r6rfzmeta_mm..rr8c "ddddd S)Nz/a and b must have same reduction dim, but got [rxz] X [r]r2)M1M2Nr?sr6rfzmeta_mm..s-[![[r[[PR[[VW[[[r8cdS)NzFout_dtype must be the same as input dtype or fp32 for fp16/bf16 inputsr2r2r8r6rfzmeta_mm..s\r8r) rWrhrrr_r\r]r^r)r4r5rBrOrrrr?s @@@@r6meta_mmrs LA55666 LA55666 GEAr GEB L b[[[[[[[   U]*?G u~>> \ \    (/177YL ;;1v\; 2 22r8c|r.tfdtjDStjjS)Nc3>K|]}|vr j|ndVdS)r1Nr5)rNrdimsrs r6rtz+_compute_reduction_shape..s5UUqatmmTZ]]UUUUUUr8)rgrrrQcompute_reduction_output_shaper)rrrs`` r6rrsRVUUUUUE$)DTDTUUUUUU  / D A AAr8strct|tjjr |jjSt |dr1t |jdr|jjdkr |jjSdS)Nrrzrr)rqrW _subclasses FakeTensor fake_devicerzhasattrr)r5s r6rrsn&%+677 !&&!! FM6 * * M & ( (}!!vr8 input_tensorrrdilation is_transposedgroupsoutput_paddingc6dtdtdtdtdtdtf d}dtdtdtdtdtdtdtfd } |jd d} |jd d|r||jd z} n;|jd } |jd |z|jd krtd |jd | gt|tr|gt z}n,t |d kr|d gt z}t|tr|gt z}n,t |d kr|d gt z}t|tr|gt z}n,t |d kr|d gt z}d} |rXt|tr|gt z} n/t |d kr|d gt z} n|} t t D]} | rH| | || || | | || | | L|| || || | | || d dlm }tj |dd dDfdS)NlnrFrrDrr>c6|d|zz||dz zz dz |zdzS)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 rr1r2)rrFrrDrs r6_formulaz+calc_conv_nd_return_shape.._formula s.QU Q!a%[(1,2Q66r8rCc<|dz |zd|zz ||dz zz|zdzS)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 r1rr2)rrFrrDrrCs r6_formula_transposedz6calc_conv_nd_return_shape.._formula_transposed$ s2"Q!|a!e#a1q5k1B6::r8rr1rzInvalid channel dimensions)sym_orcg|]}|dk Srr2rs r6rPz-calc_conv_nd_return_shape..u s...1Q...r8c@dtddddS)NzGiven input size per channel: z&. Calculated output size per channel: rz. Output size is too small)r)r ret_shapesr6rfz+calc_conv_nd_return_shape..v s6$d$$/8}$$$r8) rrr+rqrrrrr)rrWrh)rr>rrrrrrrr kernel_size out_channelsoutput_padding_listrrrrs @@r6calc_conv_nd_return_shaper s7S7S7S7S7S7S7777";;;;;;;QT;;;;&,qrr"K  abb !D= Q/ |A <?V #|'9!'< < <;<< <#A& 5I&'"")CII% V  )s4yy('7##+)c$ii' W  1:,T*(G$$-:D ) X!  QK=3t99,,01 ng . . 1#1"2SYY">    A % %#1!#4"5D "A  "0  3t99       ##GAJQKN1I'*         a'!*hqk;q>6RS9UU    =<<<<< L.. !"" .../ $ $ $ $ $ r8c\tj|tjkSrBrW _prims_commonr channels_lasttens r6is_channels_lastr~ s!   4 4S 9 9U=P PPr8 running_mean running_vartrainingexponential_average_factorepsiloncvj}||jn|j} ||jn|j} fd} || } |r+| } | }n*d} d}| | |fS)Nctr tjStjr tjStjSr)rrWrr`r)rsr6pick_memory_formatz2meta_miopen_batch_norm..pick_memory_format sJ L ) ) '& &  % %E4K % L L +* *&&r8rr)rrr,)rr>r@rrrrrrsave_mean_shapesave_var_shaperr save_meansave_vars` r6meta_miopen_batch_normr s"I-9,Dl((&,O*5*A[&&v|N'''''   + + . .=O=O=Q=Q . R RC0 **?;; )).99 **400 ))$//  8 ##r8c t||||||||r|nd} d} d} || dkrd| | <|| } | SNr1r)rrr) rr>r@rrrrrr shape_outinput_channels_dimoutput_channels_dimrs r6 meta_convr s*'1T  I+,,11)* %&   + +C Jr8mkldnnc t|||||d|g} || } tj} |dkr tj} | | } | S)NFr-r)rrrWrrchannels_last_3dr,) rr>r@rrrrattrscalars algorithmrrout_memory_formats r6meta_mkldnn_convolution_defaultr s{. &&'8UFB  $$Y//!/      " " % 6 ff#4f55 r8cf|g|jdd|jdRSNrrrr)rr>r@rrrs r6meta_linear_pointwise_defaultr s8%%&Q (:3B3(?&Qa&Q&QRRRr8mklcf|g|jdd|jdRSrr)r packed_weight orig_weightr@rs r6meta_mkl_linearr s>))@,$SbS)@;+. +S+S+S+Sr8c4ddS)Nzw must be a 2D tensor, got r rrsr6rfz/meta_int4mm_packed_weight_cpu.. r r8cdjSNz#expected x to be f32/f16/bf16, got rr sr6rfz/meta_int4mm_packed_weight_cpu.. sC!'CCr8cdjSNzexpected w to be uint8, got rrsr6rfz/meta_int4mm_packed_weight_cpu.. s,T17,T,Tr8cdjS)Nz q_group_size must be int64, got r) q_group_sizesr6rfz/meta_int4mm_packed_weight_cpu.. sK|7IKKr8cdjS)Nz5q_scale_and_zeros must have the same dtype as x, got r)q_scale_and_zerossr6rfz/meta_int4mm_packed_weight_cpu.. seL]Lceer8rr) rWrhrr_r\r]r^rrrrr;rrrs````r6meta_int4mm_packed_weight_cpur sI QUUWW\#S#S#S#STTT QUUWW\#S#S#S#STTT G u}enE E C C C C     Gu{ "$T$T$T$T      %+ - K K K K      #qw . e e e e   {{166!99affQiiqw{???r8ctjkojkfddS)NcndddddddjzS)NzExpected a tensor of dimension z and tensor.size[z] == rxzbut got : dimension z] = rr)rdim_sizerr5sr6rfz check_dim_size.. sU_#____W[___ f f f f ffl[cNd f fgr8)rWrhrr)r5rrrs````r6check_dim_sizer so L > X 6$ > g g g g g g gr8cd}|d|\}} tjt|dvdtjjtjtjtjtjfvfdt|dkr|| } } n3t|dkr|d|d} } n|d|\} } |d |\} } tj|dup|dkd d kr d nd} d } d} d}t||| | d|}t|| | | d|}tj }t|| | | | | dd||||||dkr|||g}n||||g}tj|jj|S)Nctjt|dvfd|d}t|dkr|n|d}||fS)Nr1rcddS)Nz avg_pool2d: 4 must either be a single int, or a tuple of two intsr2rsr6rfz1meta_avg_pool2d..unpack.. ]4]]]r8rr1rWrhrrrHWs` r6unpackzmeta_avg_pool2d..unpack ` HH  ] ] ] ]    FSQAACF!t r8rrr1rcdSNzOavg_pool2d: stride must either be omitted, a single int, or a tuple of two intsr2r2r8r6rfz!meta_avg_pool2d.. ar8c>djdS)Nz""avg_pool2d" not implemented for 'rrrsr6rfz!meta_avg_pool2d.. !Qu{7J7J7L7LQQQr8rr1rrcdSNzdivisor must be not zeror2r2r8r6rfz!meta_avg_pool2d.. *r8rrrrr.rz)rWrhrr_ruint16uint32uint64rrpooling_output_shaperQrpool2d_shape_checkrr)rrrrrcount_include_paddivisor_overrider*kHkWdHdWpadHpadWrr inputHeight inputWidthrrrrs` r6meta_avg_pool2drE sRVM; / /FB L F y aa L EKu|U\RRQQQQ 6{{aRB V  F1IB&))B 7++JD$ LD 9$4$9** %yy{{a//UZZ^^^QF**R..K**R..KBJ' Rr1iPPL&z2tRINNK/66M         $ yy{{a\;7 \;? ; k|#    r8ct|||||||dd| | | | | ||}| }t|||dz |t|||dz | t|||dz | dS)Nr1r.r)r:rr)r gradOutputrr=r>r?r@rArBrrCrDrr mem_formatr nOutputPlanes r6avg_pool2d_backward_shape_checkrJ- s"         $ 99;;DL:tTAX|<<<:tTAX|<<<:tTAX{;;;;;r8cjtjt|dkpt|dkd|d}t|dkr|n|d} tjt|dkp%t|dkpt|dkdt|dkr|n|d} t|dkr| nt|dkr| n|d} tjt|dkpt|dkd|d} t|dkr| n|d} tj|dup|dkd|j}|dkr|d nd}|d }|d }|d }t ||| | d|}t || | | d|}t j|}t||||| | | | | ||||||tj ||j |j | S)Nr1rcdS)NzKavg_pool2d: kernel_size must either be a single int, or a tuple of two intsr2r2r8r6rfz*meta_avg_pool2d_backward..g s]r8rcdSr.r2r2r8r6rfz*meta_avg_pool2d_backward..m r/r8cdS)NzGavg_pool2d: padding must either be a single int, or a tuple of two intsr2r2r8r6rfz*meta_avg_pool2d_backward..s sYr8cdSr3r2r2r8r6rfz*meta_avg_pool2d_backward..z r4r8rrr5rrrz) rWrhrrrr9rQrrJrr_r) gradOutput_rrrrrr;r<r=r>r?r@rArB input_sizerrrCrDrrrHs r6meta_avg_pool2d_backwardrRY sz L KA6[!1!1Q!6]] QB;1$$+a.B L F q@CKK1,@F q0@aa6{{aVAYB6{{a3v;;!+;+;RRB L G .S\\Q.YY 1:Dw<<1$$44'!*D LD 9$4$9** J$yy{{a//Z^^QFR.KR.KBJ' Rr1iPPL&z2tRINNK,U33J#       $ ;k|    r8ctjt|dvd|d}t|dkr|n|d}t|dkr|n|d} tj| pt|dvdtjjtjtjtjtjfvfd|s|n|d} |s|nt|dkr| n|d} |s| nt|dkr| n|d} tjt|dvd|d} t|dkr| n|d}t|dkr| n|d}tjjd vd tj| p|dkd  d} d } d } d} d}t||| | d|}t|||| d|}t|| || d|}t|||| | | | | ||ddd||||||ddjdkr ||||fS |||||fS)Nr1r.cdSNzFavg_pool3d: kernel_size must be a single int, or a tuple of three intsr2r2r8r6rfz!meta_avg_pool3d.. Xr8rr1rcdSNzJavg_pool3d: stride must be omitted, a single int, or a tuple of three intsr2r2r8r6rfz!meta_avg_pool3d.. ryr8c>djdS)Nz""avg_pool3d" not implemented for 'rrrsr6rfz!meta_avg_pool3d.. r1r8cdSNzBavg_pool3d: padding must be a single int, or a tuple of three intsr2r2r8r6rfz!meta_avg_pool3d.. Tr8rr-cdSNz9non-empty 4D or 5D (batch mode) tensor expected for inputr2r2r8r6rfz!meta_avg_pool3d.. Kr8cdSr3r2r2r8r6rfz!meta_avg_pool3d.. r4r8rr5rrz avg_pool3d()T)check_input_sizer) rWrhrr_rr6r7r8rrr9pool3d_shape_checkr)rrrrrr;r<kTr=r>dTr?r@padTrArBrnslicesitimeiheightiwidthotimeoheightowidths` r6meta_avg_pool3dro s2 L KF"XX QB;1$$+a.B;1$$+a.B L  +c&kkV+\\ L EKu|U\RRQQQQ (vayB Fc&kkQ&6&6F1IB Fc&kkQ&6&6F1IB L G TT 1:Dw<<1$$44'!*Dw<<1$$44'!*D L fKK  L 5 0A 5** ZZ]]FjjnnG JJrNNEjjnnG ZZ^^F D"a C CE"7Bb!YGGG !&"dB9 E EF               -2 zQ@AAAHIIIr8ctjt|dvd|d}t|dkr|n|d} t|dkr|n|d} tj| pt|dvd|s|n|d} |s| nt|dkr| n|d} |s| nt|dkr| n|d} tjt|dvd|d}t|dkr|n|d}t|dkr|n|d}tj|jdvd tj| p|dkd |d }|d }|d }|d}t |||| d|}t || || d|}t || || d|}t ||||| | | | | |||||||||d||jS)NrTcdSrVr2r2r8r6rfz*meta_avg_pool3d_backward.. rWr8rr1rcdSrYr2r2r8r6rfz*meta_avg_pool3d_backward.. ryr8cdSr\r2r2r8r6rfz*meta_avg_pool3d_backward.. r]r8r^cdSr`r2r2r8r6rfz*meta_avg_pool3d_backward.." rar8cdSr3r2r2r8r6rfz*meta_avg_pool3d_backward..' r4r8rr5rrzavg_pool3d_backward()) rWrhrrrr9avg_pool3d_backward_shape_checkrr)rrrrrrr;r<rer=r>rfr?r@rgrArBrhrirjrkotime_for_shape_checkoheight_for_shape_checkowidth_for_shape_checks r6meta_avg_pool3d_backwardrz s L KF"XX QB;1$$+a.B;1$$+a.B L  +c&kkV+\\ (vayB Fc&kkQ&6&6F1IB Fc&kkQ&6&6F1IB L G TT 1:Dw<<1$$44'!*Dw<<1$$44'!*D L fKK  L 5 0A 5** jjnnG JJrNNEjjnnG ZZ^^F0D"aSS27Bb!YWW1&"dB9UU#           ', ??5; ' ''r8ctjjdkp jdkfdjddt |z}t j}tj|jj |S)Nr.rcdjS)Nz"Expected 3D or 4D tensor, but got r5rsr6rfz*meta_adaptive_avg_pool2d..P ATZAAr8rrz) rWrhrrrgrQrrr_r)r output_sizercrs` r6meta_adaptive_avg_pool2drL s L Q($)q.AAAA:crc?U;%7%77L/55M ;j{#    r8ctjjdkp jdkfdjddt |zS)Nrr-cdjS)Nz"Expected 4D or 5D tensor, but got r5rsr6rfz*meta_adaptive_avg_pool3d..b r}r8r5)rWrhrrrrg)rr~s` r6meta_adaptive_avg_pool3dr^ sb L Q($)q.AAAA >>$*SbS/E+,>,>> ? ??r8cj}td|D]2tjdkfd3tj|dkp|dkfdtjjjkfdtj}tr tj} j  |S) Nr1rc djddS)Nz{adaptive_avg_pool2d_backward(): Expected grad_output to have non-zero size for non-batch dimensions,  with dimension being emptyr5)grad_outrsr6rfz4meta__adaptive_avg_pool2d_backward..m s1f6>nffVWfffr8r.rcdjS)NzBadaptive_avg_pool2d_backward(): Expected 3D or 4D tensor, but got r5rsr6rfz4meta__adaptive_avg_pool2d_backward..r saUYU_aar8c(djdjSNexpected dtype z! for `grad_output` but got dtype r)rrsr6rfz4meta__adaptive_avg_pool2d_backward..v s_$*__x~__r8r) rrrWrhrr_rrrrrr,)rrrrrs`` @r6"meta__adaptive_avg_pool2d_backwardrg s =D 1d^^   MM!  q  f f f f f     L  TQYaaaa L hn$_____+M,+ >>$* % % ( (} ( E EEr8cbt|dtj|tjS)Nadaptive_avg_pool3d_backwardr)!_adaptive_pool_empty_output_checkrWrr rrs r6"meta__adaptive_avg_pool3d_backwardr~ s.&k3QRRR  D0N O O OOr8rcj}td|D]3tjdkfd4dS)Nr1rc$djddS)Nzc(): Expected grad_output to have non-zero size for non-batch dimensions, but grad_output has sizes rrr5)rrrsr6rfz3_adaptive_pool_empty_output_check.. s>``-8->``PQ```r8)rrrWrhr)rrrrs`` @r6rr s{  D 1d^^     Q  ! #            r8c8 j}tj|dvfdtd|D]2 tj dk fd3tjt |dkdd}d}d}jdkrd}|dz }|dz }|\}}jd kr?|||f}|} |tj } | | fS||||f}tj } | | } |tj  | } | | fS) Nr.rcdjS)Nz:adaptive_max_pool2d(): Expected 3D or 4D tensor, but got: r5rsr6rfz*meta_adaptive_max_pool2d.. ZU[ZZr8r1rc djddS)Nzjadaptive_max_pool2d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes rrr5rrsr6rfz*meta_adaptive_max_pool2d.. 4T',{TTDETTTr8rcdS)NzCadaptive_max_pool2d(): internal error: output_size.size() must be 2r2r2r8r6rfz*meta_adaptive_max_pool2d.. Ur8rr.rr) rrWrhrrrrrrQrr,) rr~rdimHsizeBsizeDosizeHosizeWrrrrrs ` @r6meta_adaptive_max_pool2dr s :D L ZZZZ1d^^   JJqMMA           L KAUU D E E zQ 1    JJtax E NFF zQFF+ ooi((//)5;/??G|E662 3E:: ooi((++-+HH//)5;/??BB'C  G|r8c:j}tj|dvfdtdtjjjkfdt j}j |S)NrcdjS)NzKadaptive_max_pooling2d_backward(): Expected 3D or 4D grad_output, but got: r5rsr6rfz3meta_adaptive_max_pool2d_backward.. sq^i^oqqr8adaptive_max_pool2d_backwardc(djdjSrr)rrsr6rfz3meta_adaptive_max_pool2d_backward.. sc%+ccP[Paccr8r) rrWrhrr_rQrrrr,)rrrrrs`` r6!meta_adaptive_max_pool2d_backwardr s  D L qqqq &k3QRRR L {((ccccc /66M ??5; ' ' * * * G GGr8c: j}tj|dvfdtd|D]2 tj dk fd3tjt |dkdd}d}d}|dkrd}|dz }|}|\}}}|d kr||||f} n|||||f} | } | tj } | | fS) Nr^cdjS)Nz:adaptive_max_pool3d(): Expected 4D or 5D tensor, but got: r5rsr6rfz*meta_adaptive_max_pool3d.. rr8r1rc djddS)Nzjadaptive_max_pool3d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes rrr5rsr6rfz*meta_adaptive_max_pool3d.. rr8r.cdS)NzCadaptive_max_pool3d(): internal error: output_size.size() must be 3r2r2r8r6rfz*meta_adaptive_max_pool3d.. rr8r-rr)rrWrhrrrrr) rr~rdimDrrosizeTrrrrrrs ` @r6meta_adaptive_max_pool3dr sg :D L ZZZZ1d^^   JJqMMA           L KAUU D E E qyy 1    JJt  E(FFF qyyFFF3 E666: //) $ $Cooiu{o;;G <r8cVt|d||jS)Nadaptive_max_pool3d_backward)rrr)rrrs r6!meta_adaptive_max_pool3d_backwardr s)&k3QRRR ??5; ' ''r8cN|td||S)Nz:cannot repeat_interleave a meta tensor without output_size)r+r)repeatsr~s r6meta_repeat_interleave_Tensorr s+WXXX   [ ) ))r8c |jjsJ|jjsJt|t |j|t |jt j}|SNrJ)r_rrUr,rrrR)realimagrs r6 meta_complexr s{ : '''' : ''''  +DJ7788 +DJ77886>F Mr8) fill_valuercl|||ftjSr)rrrWr)rrrs r6nonzero_staticr s) >>4,EJ> ? ??r8ctjtjdtj||fd|ftj|jS)NcdS)NaYThe 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.r2r2r8r6rfznonzero..' s Sr8r1r_r) rW_check_not_implementedrcmeta_nonzero_assume_all_nonzeror$rrrrrs r6nonzeror" sy  2 S S   txxzz" DJJLLj{    r8c , tjtdg}tD]d\Ftjjtjtjtjtjfvdjtjtjfvrlj}t|tj j zj kfdtj D]dtj j j zkfd||de8|O|f|tjtj kfdddlm}t%|jtj kr-dtj k-d}d}D]|dkrd} |dkrd }nd }|sg}g}tD]1\*||2tD]1\*||2||gggtD]_\} Dr!j | *j | Kt%j `fd } zz} dd lm} | dkr| S| } t3j| \}}t%|t%tt|kr}t3j| j |}t3j|}t3j|t3j|}| | |} | S) NcdS)Nz#at least one index must be providedr2r2r8r6rfz#meta_index_Tensor..7 s(Mr8cdS)Nz?tensors used as indices must be long, int, byte or bool tensorsr2r2r8r6rfz#meta_index_Tensor..? sYr8cdjS)N)too many indices for tensor of dimension rrsr6rfz#meta_index_Tensor..F sS SSr8c :djddjdzS)NzThe shape of the mask at index z0 does not match the shape of the indexed tensor r5)rrjrDrsr6rfz#meta_index_Tensor..K sP!h!h!hPQ!h!hJN*!h!h`ade`e!h!hr8r1c:djdtdS)Nrz (got ry)rr)rrsr6rfz#meta_index_Tensor..V s$\DI\\SQX\\\\\r8rFrTczz}t|}dgtz|tt|jtz <|||S)zI This follows restride_src in TensorAdvancedIndexing.cpp r)rrrrrG)rrr after_shape before_shapereplacement_shapes r6 _restride_srcz(meta_index_Tensor.._restride_src s00;>t{{}}%%KL#PS Q Q K L!!C OOc+6F6F$FFGug...r8guard_or_false) rWrhr enumerater_rrrErrrrrrrselect torch._refs_refsrr$rrr)rrrQ3compute_elementwise_output_logical_to_physical_perm apply_permr invert_permrGr)rrrrrefsstatehas_contiguous_subspacertransposed_indicesrrrrrestrided_selfpermrT perm_shaper8rrrrrrDrs`` @@@@@@@r6meta_index_Tensorr5 s Lg M MNNN#%Fg&&!!5   L  EIuz5:NNYY   {uz5:666--//KK" Ndi/SSSSuz**88A& A$*QU*;;hhhhhhhh MM'..A"6"67777 8 e$$$$ MM% G L G  !\\\\\ (4('233G g,, " "t g,, " " E# ' ' A::  aZZ} !#' # %!'** 1 1HAu  A"))%000!'** 1 1HAu} A"))%000||D!!$!LK#%((22 U =  5""4:c?3333##DJsO4444 $U[ 1 1   / / / / / / / ..(99KG H HCDDDDDD~djjlla'(( #]4((NGWWGD! DzzT%D **++++%ci66 6zBB %j%2CD2I2IJJ nnSXXZZ44 Jr8c d} d} d} | dr'||} | dr'||} | dr||} | | | fSNrr1rrr) grad_output_input_weight_bias_sizes_optrrr transposedrr output_maskbackend_grad_inputbackend_grad_weightbackend_grad_biass r6meta_convolution_backwardr s 1~C)33FKKMMBB1~E*44W\\^^DD1~C(22>BB  35F GGr8cd}d}|||f}tjdkdtjdkdtjddkfdtjddkfdtj|d|ko|d|kd ||S) Nr1rr.cdSrTr2r2r8r6rfzmeta_addbmm.. rUr8cdSrWr2r2r8r6rfzmeta_addbmm.. rUr8rc`ddddS)Nz8batch1 and batch2 must have same number of batches, got rrrrYrZsr6rfzmeta_addbmm.. s3p6;;WX>>pp`f`k`klm`n`nppr8c ddddddddd S)Nz#Incompatible matrix sizes for bmm (r1r;rrryrrsr6rfzmeta_addbmm.. sl 6&++a.. 6 66;;q>> 6 6;;q>> 6 6$*KKNN 6 6 6r8cdS)Nz.self tensor does not match matmul output shaper2r2r8r6rfzmeta_addbmm.. s@r8)rrrWrhrr)rrYrZrRrQrerfs `` r6 meta_addbmmr sg ;;q>>D ;;q>>D ;;d| $ $D L"$H$HIII L"$H$HIII L A&++a..(ppppp L A&++a..(      L ! 51!5@@ >>$))++ & &&r8c P||SrBr)rrkwargss r6meta_randint_liker s >>$))++ & &&r8) grad_scale found_infc t||||||fD]-tjttfd.dS)Nc(dtSNz'exponent must be a tensor list but got rzlsr6rfz#meta__fused_adam_.. Gd1ggGGr8rWrhrqr)rgradsexp_avgs exp_avg_sqsmax_exp_avg_sqs state_stepslrbeta1beta2 weight_decayepsamsgradmaximizerrrs @r6meta__fused_adam_r s]&E8[/; O   q$   G G G G      r8c ||||||fD]-tjttfd.d}||||||||||fS)Nc(dtSrrrsr6rfz"meta__fused_adam..&rr8cd|DS)Nc6g|]}tj|Sr2r)rNrs r6rPz=meta__fused_adam..empty_like_list..*s#999 ##999r8r2) tensor_lists r6empty_like_listz)meta__fused_adam..empty_like_list)s99[9999r8r)rrrrr r r r r rrrrrrrrs @r6meta__fused_adamrs&E8[/; O   q$   G G G G    ::: !! $$((  r8crtjdkdtjdkdtjjtjufdtjjtjufdtjddkfdddftj S) NrcdS)Nza must be a 2D tensorr2r2r8r6rfzmeta__int_mm..8'>r8cdS)Nzb must be a 2D tensorr2r2r8r6rfzmeta__int_mm..9rr8cdjS)Nzexpected self to be int8, got r)r4sr6rfzmeta__int_mm..<:::r8cdjS)Nzexpected mat2 to be int8, got r)r5sr6rfzmeta__int_mm..@rr8r1rc ddddddddd S)Nz'Incompatible matrix sizes for _int_mm (rr;r1rryrr3sr6rfzmeta__int_mm..Dsl ,affQii , ,!&&)) , ,66!99 , , vvayy , , ,r8r)rWrhrr_rErrrFr3s``r6 meta__int_mmr"5s" LA>>??? LA>>??? L 5::::: L 5::::: L q QVVAYY      ;;q 166!99-U[; A AAr8cjtjdkdtjjtjufdd}ddz}|dz||dzzd|dzftj S) NrcdSNzw must be a 2D tensorr2r2r8r6rfz2meta__convert_weight_to_int4pack..Nrr8cdjSrrrsr6rfz2meta__convert_weight_to_int4pack..Q8qw88r8rr1 r)rWrhrr_rrrrFr inner_k_tilesr rDs` r6 meta__convert_weight_to_int4packr-Ls LA>>??? L 5;8888 q A q A A ;; F -"$ %  Q   k   r8cNtjdkdtjjtjufdd}d}||dzftjS)NrcdSr%r2r2r8r6rfz:meta__convert_weight_to_int4pack_for_cpu..brr8cdjSNzexpected w to be int32, got rrsr6rfz:meta__convert_weight_to_int4pack_for_cpu..er'r8rr1r)rWrhrr_rFrrrr+s` r6(meta__convert_weight_to_int4pack_for_cpur2`s LA>>??? L 5;8888 q A q A ;; AF k   r8ctjdkdtjdkdtjjtjtjtjfvfdtjjtjufd d ddzj S) NrcdSNzx must be a 2D tensorr2r2r8r6rfz*meta__weight_int4pack_mm..qrr8rcdS)Nzw must be a 4D tensorr2r2r8r6rfz*meta__weight_int4pack_mm..rrr8cdjSrrr sr6rfz*meta__weight_int4pack_mm..u?ag??r8cdjSr1rrsr6rfz*meta__weight_int4pack_mm..yr'r8rr(r rWrhrr_r\r]r^rFrrrs`` r6meta__weight_int4pack_mmr;os LA>>??? LA>>??? L EM5=%.AA???? L 5;8888 ;;qvvayy!&&))a-qw; ? ??r8ctjdkdtjdkdtjjtjtjtjfvfdtjjtjufd d djS)NrcdSr5r2r2r8r6rfz2meta__weight_int4pack_mm_for_cpu..rr8cdSr%r2r2r8r6rfz2meta__weight_int4pack_mm_for_cpu..rr8cdjSrrr sr6rfz2meta__weight_int4pack_mm_for_cpu..r8r8cdjSrrrsr6rfz2meta__weight_int4pack_mm_for_cpu..r'r8rr) rWrhrr_r\r]r^rrrrs`` r6 meta__weight_int4pack_mm_for_cpurA~ LA>>??? LA>>??? L EM5=%.AA???? L 5;8888 ;;qvvayy!&&))17; ; ;;r8ctjdkdtjdkdtjjtjtjtjfvfdtjjtjufd d djS)NrcdSr5r2r2r8r6rfz;_weight_int4pack_mm_with_scales_and_zeros..rr8cdSr%r2r2r8r6rfz;_weight_int4pack_mm_with_scales_and_zeros..rr8cdjSrrr sr6rfz;_weight_int4pack_mm_with_scales_and_zeros..r8r8cdjSr1rrsr6rfz;_weight_int4pack_mm_with_scales_and_zeros..r'r8rrr:)r;rrqScaleqZeross`` r6)_weight_int4pack_mm_with_scales_and_zerosrJrBr8r4r5c||zdz |z|zSr0r2r3s r6 kai_rounduprLs UQY1  !!r8c0 |dkr||kr+d}d}d}dddd fd fd}||||||S|dzd krH||zd krAd}d}d}dddd fd } fd  d fd  fd|||||||SdSdSdS)Nrr(r)rcHt||zd}t||S)NrrL)rDkrsrkr_sr_roundedup4s r6kai_k_roundedupz3get_kai_packed_weight_size..kai_k_roundedups)$/rBw#:#: "1&6777r8cf|||}|dzdks Jd||dzzzzzS)Nrrzk_internal must be evenr2) rDnrrPrQ k_internalrSkai_num_bytes_biaskai_num_bytes_multiplier_rhskai_num_bytes_sum_rhss r69kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0z]get_kai_packed_weight_size..kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0sd-_QB77 "Q1,,,.G,,,1_23+,))r8cLt|||z}|||||zSrBrO)r rDrUrPrQnum_rowsrZs r67kai_get_rhs_packed_size_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0z[get_kai_packed_weight_size..kai_get_rhs_packed_size_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0sA'q"--3OO2r2r8r*rc||zdksJ| zdksJ|zdksJt|||z}||||||zSrorO) r rDrUrPrQblr\kai_bl_multiple_of;kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0kai_nr_multiple_ofs r69kai_get_rhs_packed_size_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0z]get_kai_packed_weight_size..kai_get_rhs_packed_size_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0sRA~~~~//A5555//A5555&q"--3QQ2r2rr8c||zdksJ| zdksJ|zdksJ } ||} ||}|||zz zzSror2)rDrUrPrQr_num_bytes_multiplier_rhsnum_blocks_per_rownum_bytes_per_blockr`#kai_get_bf16_datatype_size_in_bytesrbkai_num_blocks_per_rowrWkai_num_bytes_per_blockrYs r6raz_get_kai_packed_weight_size..kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0sRA~~~~//A5555//A5555,O+N+P+P(%;%;Ar%B%B"&=&=0''#(+==+,()r8cdS)Nrr2r2r8r6rhzGget_kai_packed_weight_size..kai_get_bf16_datatype_size_in_bytessqr8c@|zdksJt|||zSrorO)rDr_r`s r6riz:get_kai_packed_weight_size..kai_num_blocks_per_rows///A5555"1b))R//r8c*|zdksJ|dz|zS)Nrrr2)r_rer`s r6rjz;get_kai_packed_weight_size..kai_num_bytes_per_blocks*//A5555a#;;;r8r2)n_bitsrK groupsizekai_nrkai_krkai_srr]rcr`rhrarZrSrbrirWrXrjrYs @@@@@@@@@@@r6get_kai_packed_weight_sizerts {{ >>FFF$% !+, (!"  8 8 8             KJ1fff ^q Q]a%7%7FFF$% !!" !" !#                    ,    0 0 0 0 0 < < < < <ML1fffi Q{\! %7%7r8crtjjtjufdtjjr||kr|jtjks-||krg|dzdkr^||zdkrU|jtjkr@td|||} t|tjS | z}||| z } |tjS)NcdjSrr)weightssr6rfz2meta__dyn_quant_pack_4bit_weight..s>w}>>r8r*rrr) rWrhr_rbackendskleidiai is_availabler[r^rtrrr)rw scales_zerosr@ block_size in_features out_featurespacked_weight_sizes` r6 meta__dyn_quant_pack_4bit_weightrs6 L $>>>> ~++-- M { " "|'9U['H'H  $ $R1$$j(A--"en448 |[*    %7!8!8  LLL <+=+=+?+?? djjll*   /u{  C CCr8cTtjdkdtjjtjkpjtjkokfdd}||jS)NrcdS)Nzinput must be a 2D tensorr2r2r8r6rfz-meta__dyn_quant_matmul_4bit..1s)Dr8c$djddS)NzPexpected input to be f32 or bf16 (bf16 requires block_size == in_features), got z with block_size=z and in_features=r)r|r}inpsr6rfz-meta__dyn_quant_matmul_4bit..5s= Z9 Z Z/9 Z ZLW Z Zr8rr)rWrhrr_r\r^rr)rpacked_weightsr|r}r~rs` `` r6meta__dyn_quant_matmul_4bitr)s La!D!DEEE L em # G I ' EJ+,E        A ==L = : ::r8ctjdkdtjjtjtjtjfvfdtjdkdtjjtjufd d djS)NrcdSr5r2r2r8r6rfz*meta__weight_int8pack_mm..@rr8cdjSrrr sr6rfz*meta__weight_int8pack_mm..Cr8r8cdSr%r2r2r8r6rfz*meta__weight_int8pack_mm..Err8cdjS)Nzexpected w to be int8, got rrsr6rfz*meta__weight_int8pack_mm..Hs7ag77r8rr) rWrhrr_r\r]r^rErr)r;rq_scaless`` r6meta__weight_int8pack_mmr>s LA>>??? L EM5=%.AA???? LA>>??? L 5:7777 ;;qvvayy!&&))17; ; ;;r8ctjdkfdtjdkfdtjddkfdtjt jjfdtjt jjfdtj|dkd tjd vfd d }d }jdd }jdd }ttj ||}| ||g |S) Nrc4ddS)Nz1cdist only supports at least 2D tensors, X1 got: r rx1sr6rfz$meta_cdist_forward..QOBFFHHOOOr8c4ddS)Nz1cdist only supports at least 2D tensors, X2 got: r rx2sr6rfz$meta_cdist_forward..Urr8rc`ddddS)Nz4X1 and X2 must have the same number of columns. X1: rz X2: r)rrsr6rfz$meta_cdist_forward..Ys2frwwr{{ffY[Y`Y`acYdYdffr8cdjS)Nz3cdist only supports floating-point dtypes, X1 got: rrsr6rfz$meta_cdist_forward..]PbhPPr8cdjS)Nz3cdist only supports floating-point dtypes, X2 got: rrsr6rfz$meta_cdist_forward..arr8rcdS)Nz)cdist only supports non-negative p valuesr2r2r8r6rfz$meta_cdist_forward..cs!Lr8rcdS)Nz(possible modes: None, 0, 1, 2, but was: r2) compute_modesr6rfz$meta_cdist_forward..fsI<IIr8r) rWrhrrrQis_float_dtyper_rrbroadcast_shapesextendr) rrrFrr1r2 batch_tensor1 batch_tensor2rcs `` ` r6meta_cdist_forwardrMs L A OOOO L A OOOO L  rwwr{{"fffff L RX&&PPPP L RX&&PPPP LaLLMMM L'IIII B BHSbSMMHSbSMM.}mLLMMLR!!! << % %%r8c@|jd}|jd}|jd}|jdd}|jdd} ttj|| } | } | ||gt j| } |dks|dks |dks| dkrtj|S| t|jkr| | }tj |tj S)Nrrrr) rrrWrcopyrmathprod zeros_likerrr) rJrrrFcdistc1rrrrr{tensor1_expand_size batch_products r6meta_cdist_backwardrqs  "B "B "BHSbSMMHSbSMM 6}m T TUU.3355Bx(((I233M Qww"''R1WW (:(:###d28nn,, YY* + +  Be.E F F FFr8c rtjjtjtjfvfdtjjtjtjfvfdtjt jjfdd} |rtj| dkd| dz} | d} tj|tkdtjj dkfdtj  kfd fd d fd } td krd} } |tkr+| d}n d}n| | |}|ttfvs|s)d} nd} | } jd}|tkrC|rtj|dkd|dz}|jd}n'| }| | | |fS)NcdjS)Nz(expected indices to be long or int, got r)rsr6rfz$meta_embedding_bag..J7=JJr8cdjS)Nz(expected offsets to be long or int, got r)rssr6rfz$meta_embedding_bag..rr8cdjS)Nz/expected weight to be floating point type, got r)r>sr6rfz$meta_embedding_bag..sP&,PPr8rr1cdSNz1include_last_offset: numBags should be at least 1r2r2r8r6rfz$meta_embedding_bag..sGr8cdS)Nz@embedding_bag: per_sample_weights only supported with mode='sum'r2r2r8r6rfz$meta_embedding_bag..sVr8cdjdS)Nz1expected per_sample_weights to be 1D tensor, got r r)per_sample_weightssr6rfz$meta_embedding_bag..sbHZH_bbbr8c^dddS)Nz%expected per_sample_weights.numel() (z$ to be the same as indices.numel() (ryr)rrsr6rfz$meta_embedding_bag..sFI8J8P8P8R8RII6=mmooIIIr8cP|||o|ddkSNrr1r)r/r~r padding_idxis_fast_path_index_selects r6is_fast_path_index_select_scalez;meta_embedding_bag..is_fast_path_index_select_scales0 % %c6; ? ? XELLQROOWXDX r8c|jtjks|jtjko7|ddko|ddko|dkSr)r_rWr[rYr)r/rrs r6rz5meta_embedding_bag..is_fast_path_index_selects_ Y%+ % @ej)@  1 "  a  A% a  r8c>|||||S|||SrBr2)r/r~rrrrs r6 is_fast_pathz(meta_embedding_bag..is_fast_paths5  223v{SS S,,S&+FF Fr8cpucdSrr2r2r8r6rfz$meta_embedding_bag..sOr8)rWrhr_rrrQrrrMODE_SUMrrrMODE_MAX MODE_MEANr)r>rrsscale_grad_by_freqrYsparserinclude_last_offsetrnum_bagsrr offset2bagbag_size max_indices fast_path_sumnumBagsrrs``` ` @@r6meta_embedding_bagrs L %*ei00JJJJ L %*ei00JJJJ L V\**PPPP ||AH M G G    A    h A 7 7F% H  V V      #q ( b b b b      $ $ & &'--// 9                GGGGGG 7u$$&&w||A77 $$W\\^^44 8  !++Hfkk!nnEEKK!++A..KK$ V-?UU Ix( ( ( ( **7<<??;;JJ **1--J$$X..-" 8  "  qLOO1 !++GV\!_EEKK!++HMMOO<rrsrSrrrrs r6meta_embedding_bag_forward_onlyrsm0B1#'111-FJ+7u$$$$W\\^^44 :x 44r8cr|r|S|jjs |jjr|jS|r tjS|jSrB)r_rrrWr)rr_promote_int_to_longs r6_get_reduction_dtypersH   {$ (>{ z ;r8rct||d}tj|j|}t |||}|||S)NT)rr)rrQrrrr)rrrr_rrcs r6 meta_nansumrsS(u$OOOL   T 2 2D+E4AAL ??<|? < <.3slr8rcddS)Nz"Repeats cannot be negative, found rr2)rrepsr6rfzmeta_repeat..8sKKKKKr8rc2g|]}||zSr2r2)rNr padded_sizers r6rPzmeta_repeat..?s&LLL1;q>GAJ.LLLr8) rWrhrrrrgrrr)rrnum_new_dimensions target_sizerrrs ` @@@r6 meta_repeatr/s L G  "llG$$  3 1H K K K K K    W 2++eDJ.?.??KLLLLLc'll8K8KLLLK >>+ & &&r8c|SrBr2rs r6 meta_zero_rCrqr8cnt|tjrt|j|j|SrB)rqrWr rjrrrs r6meta_binop_inplacerHs1%&&9 EK888 Kr8cd}d}d}||r||rtd||r||stdt|tjrt |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. ct|trtj|jSt|t SrB)rqrrQr:r_rrss r6 is_integericz.meta_binop_inplace_alpha..is_integericks7 c: & & ,)#)44 4c7++ +r8ct|trtj|jSt|t SrB)rqrrQrr_rrs r6 is_floaticz,meta_binop_inplace_alpha..is_floaticqs7 c: & & .' 22 2c9-- -r8ct|trtj|jSt|t SrB)rqrrQis_boolean_dtyper_rrs r6 is_booleanicz.meta_binop_inplace_alpha..is_booleanicws7 c: & & -)#)44 4c8,, ,r8z]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+rqrWr rjr)rrrQrrrs r6meta_binop_inplace_alpharYs$,,, ... ---|D jj//  k   |D ,,u"5"5  m   %&&9 EK888 Kr8c:t||tjSrrUrrRrrrQs r6meta_binop_alphars&  e$C$K   r8c 8t|tjSrr)rrs r6 meta_roundrs"  <D   r8c\tjtjjfdt tjr2tjtjjfddStjt tfddS)NcdjS)Nz7: Expected input tensor to have an integral dtype. Got r)rrsr6rfz#shift_dtype_check..s7__SWS]__r8cdjS)Nz6: Expected shift value to have an integral dtype. Got rrrsr6rfz#shift_dtype_check..swaaVYV_aar8cdS)Nz): Expected shift value to be an int. Got r2rsr6rfz#shift_dtype_check..swNNNNr8)rWrhrQr:r_rqr r)rrrs```r6shift_dtype_checkrs L tz**_____#u|$$    "39 - - a a a a a      sG $ $ N N N N N     r8c\td||t||tjS)NrshiftrrrUrrRrs r6 meta_rshiftsr6he,,,  e$C$K   r8c\td||t||tjS)Nlshiftrrrs r6 meta_lshiftsr r r8c6||jSrBrrs r6 meta_zerors >>$* % %%r8c|SrBr2rrs r6 meta_fill_rrqr8c*tj|SrBrrs r6 meta_fillr  D ! !!r8c|SrBr2rs r6 meta_relu_rrqr8c:t||tjSrrrs r6meta__add_relurs&  e$C$K   r8?UUUUUU?c*tj|SrBrrnoiselowerrrr s r6meta_rrelu_with_noisers  D ! !!r8cRtj|tj|fSrBrrs r6 meta_rrelu_with_noise_functionalr!s%  D ! !5#3E#:#: ::r8c|SrBr2)rrrrr s r6meta_rrelu_with_noise_r#s  Kr8c*tj|SrBrrrr accumulates r6meta_index_putr'rr8c:t|j|j|SrBrjr)rrvalues r6meta_masked_fill_r+sDJ 333 Kr8c||tj|}|Sr)rrr,rQr)rrr~ masked_scales r6meta__masked_scaler.sF>>$))++..111$772L r8ctj|jtjtjfvdtjjjkfdS)NcdS)NzMask must be bool or uint8r2r2r8r6rfz&meta_masked_scatter_.. s9Ur8c(djdjS)NzEmasked_scatter: expected self and source to have same dtypes but got rr)rrfsr6rfz&meta_masked_scatter_..s(:*::+1<::r8)rWrhr_rr)rrrfs` `r6meta_masked_scatter_r2ss L uz5;//1U1U L fl" : : : : : Kr8ct||\}}tj|tj}t |||Sr)r$rWrrr2)rrrfrs r6meta_masked_scatterr4sA"$--JD$  d%2I J J JF f 5 55r8c,||SrBr)rrrCs r6meta_masked_scatter_backwardr6s >>%  r8c|SrBr2r%s r6meta_index_put_r8!rqr8c  ddlm}m}tj|dkdtj|dkd|}| |d |d|d} d} || ftj|| d | d fd|r|jtjkp|jtj ko|tj k} tj||jkp| d |  |} n| } |scatjdkd tj|fd | S) Nr)sym_andrRr.cdSrTr2r2r8r6rfz)common_meta_baddbmm_bmm..)rUr8cdSrWr2r2r8r6rfz)common_meta_baddbmm_bmm..*rUr8rr1c :ddddddd Sr\r2r^sr6rfz)common_meta_baddbmm_bmm..7sURSURR RR+7?RR>J1oRRRr8cdS)Nzfout_dtype only supported for torch.float32 output with float16/bfloat16 inputs or same as input dtypesr2r2r8r6rfz)common_meta_baddbmm_bmm..@s|r8cdS)Nzself must be a 3D tensorr2r2r8r6rfz)common_meta_baddbmm_bmm..Hs6Pr8c8ddS)Nz*Expected an input tensor shape with shape z but got shape: r)r~ self_baddbmmsr6rfz)common_meta_baddbmm_bmm..Ks%sss^j^o^o^q^qssr8) r)r:rRrWrhrrr_r]r^r\rr,)rYrZis_bmmrArBr:rRrhres_rowsres_colssupported_out_dtyperr_r`rar~s ` @@@@r6common_meta_baddbmm_bmmrF&s/EEEEEEEE L"$H$HIII L"$H$HIII;;==L;;==L aB#AAHAHx*K L|A++VVLOEU-V-VWW R R R R R R  / LEM ) KV\U^-K)5=(    % <)< | |   !!+..11)<<!!+..  l. \%%''1,.P.PQQQ F<$$&& 4 4 s s s s s   Mr8c$t||dS)NTrF)rrMs r6meta_bmmrIQs "4t 4 44r8c(t||d|S)NT)rBrH)rrMrBs r6meta_bmm_dtyperKVs "4ty I I IIr8c||z}||z}|dkr-t|dkt|dkkr|dz}|Sr)r)r;r<qr]s r6div_rtnrN[sK QA AA Avv4A;;$q1u++-- Q Hr8ct||z|z||dz zz dz |r|dz ndz|dz}|r|dz |z||zkr|dz}|Sr)rN) inputSize kernelSizerrrrr outputSizes r6pooling_output_shape_pad_lrrSes    *q.) *   '-vzzA  /      Nf $ E(9 9 9 !OJ r8c tj|dkdtjdkfdtjdz zdzdzkfdt|||S)NrcdS)Nzstride should not be zeror2r2r8r6rfz&pooling_output_shape..s&Ar8cdS)Nz'pad must be non-negative, but got pad: r2padsr6rfz&pooling_output_shape..s#RS#R#Rr8r1rcdddS)NzApad should be at most half of effective kernel size, but got pad=z, kernel_size=z and dilation=r2)rrQrXsr6rfz&pooling_output_shape..s; @PS @ @% @ @5= @ @r8)rWrhrS)rPrQrXrrrs `` ` r6r9r9s L1AABBB LRRRRSSS L a8+a/A55       ':sC9  r8c j } tjdkodkfdtjdkodkfdtjdkodkfdddkoddk}|tjkr8tj|dko|oddkfd nXtj|dkrddkr|p |dko|oddkfd tjdzkodzkfd tj dko dk  fd dS) NrcddS)Nz5kernel size should be greater than zero, but got kH: , kW: r2)r=r>sr6rfz$pool2d_shape_check..sVVVRTVVr8cddS)Nz0stride should be greater than zero, but got dH: , dW: r2)r?r@sr6rfz$pool2d_shape_check..sQ2QQRQQr8cddS)Nz9dilation should be greater than zero, but got dilationH: , dilationW: r2) dilationH dilationWsr6rfz$pool2d_shape_check..soIoodmoor8r1rrr.c2dS)NzExpected 4D (batch mode) tensor expected for input with channels_last layout with optional 0 dim batch size for input, but got: rrsr6rfz$pool2d_shape_check..s"RCH::<<RRr8c2dS)NzYExpected 3D or 4D (batch mode) tensor with optional 0 dim batch size for input, but got: rrsr6rfz$pool2d_shape_check..s~pupzpzp|p|~~r8c ddddS)NzKpad should be smaller than or equal to half of kernel size, but got padW = z , padH = z, kW = z, kH = r2)r=r>rArBsr6rfz$pool2d_shape_check..s>>>>!%>>.0>>9;>>r8c .ddddddd SNzGiven input size: (r;z). Calculated output size: (z). Output size is too smallr2)rCrDrrIrrsr6rfz$pool2d_shape_check..sZ#k##K##*##$0##3?##BM###r8)rrWrhrr)rr=r>r?r@rArBrarbrrCrDrrrr valid_dimsrIs`````````````` @r6r:r:sc" 99;;DL L Q26VVVVV L Q26QQQQQ LA ')a-ooooo A!#: 1 (:J+++ AI ;* ;A!); R R R R     QY <5::a==A- <* A ?j?UZZ]]a-? ~ ~ ~ ~    L a4+B!GtO > > > > > > >  Lq.\Q. # # # # # # # # #r8rhrer=r>rfr?r@pTpHpW dilationTrarbrirjrkrlrmrnrcc 8 j}tjdko dkodkfdtjdko dkodkfdtj dko dko dk fdtj|dvfdt|D]@|dkrdkrtjdkfdA|r-tjko kokfd tjd z kod z kod z k  fd tjd ko d kod kfd dS)NrcdddS)Nz5kernel size should be greater than zero, but got kT: z, kH: r\r2)r=rer>sr6rfz$pool3d_shape_check..s4 , , , , ,') , ,r8cdddS)Nz0stride should be greater than zero, but got dT: z, dH: r^r2)r?rfr@sr6rfz$pool3d_shape_check..s$ Wr W W W WSU W Wr8cdddS)Nz9dilation should be greater than zero, but got dilationT: z , dilationH: r`r2)rarlrbsr6rfz$pool3d_shape_check..s; V# V V2; V VJS V Vr8r^cdjS)Nz/: Expected 4D or 5D tensor for input, but got: r5)rrsr6rfz$pool3d_shape_check..s7XX5;XXr8r-cJdjddS)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!)rr)rrrsr6rfz$pool3d_shape_check..sGMM-2[MM,1JJqMMMMMr8c .ddddddd S)Nzinput image (T: r5rz ) smaller than kernel size (kT:  kH:  kW: ryr2)rjrirkr=rer>sr6rfz$pool3d_shape_check..s]<5<rjrirksr6rfz$pool3d_shape_check..sh K K K K K%' K K02 K K;= K KFH K Kr8r1c:dddddddddSrgr2)rjrirkrhrmrlrnsr6rfz$pool3d_shape_check..!sv (' ( (E ( (G ( (f ( ((/ ( (27 ( (:A ( (DJ ( ( (r8)rrWrhrr)rrhrer=r>rfr?r@rirjrkrlrarbrirjrkrlrmrnrrcrrs````````````````````` @r6rdrds0 :D L Q$26$b1f       L Q$26$b1f        LA 9)a-9IM       L XXXXX 4[[    199a  JJqMMA             RK :GrM :fl             L Q" 6a26"q&B,          L  3v{3w!|          r8c|j}t|||||||| | | | | |||||||||t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |dS)Nrr.rr1rrdr)rrrrhrer=r>rfr?r@rirjrkrlrarbrirjrkrlrmrnrrs r6max_pool3d_backward_shape_checkrz)s$2 :D            +0;dQh888;dQh666;dQh888;dQh7777D$(G4447D$(E2227D$(G4447D$(F33333r8c|j}t||||||||| | | ddd| | |||||dt|||dz |t|||dz |t|||dz |t|||dz |dS)Nr1Trr.rry)rrrhrer=r>rfr?r@rirjrkrirjrkrlrmrnrrs r6rvrvgs* :D                -2;dQh888;dQh666;dQh888;dQh77777r8cPd}|d|\}}tjt|dvdt|dkr||} } n|d|\} } |d|\} } |d|\} }|d }|d }|d }t j|}|tjkr-tj|d kd nQ|tjkr+tj|dvdntjddt||| | | |}t||| | ||}t|||| | | | | ||||||||||fS)Nctjt|dvfd|d}t|dkr|n|d}||fS)Nr!cddS)Nz max_pool2d: r#r2r$sr6rfzEmax_pool2d_checks_and_compute_shape..unpack..r%r8rr1r&r's` r6r*z3max_pool2d_checks_and_compute_shape..unpackr+r8rr,cdS)NzOmax_pool2d: stride must either be omitted, a single int, or a tuple of two intsr2r2r8r6rfz5max_pool2d_checks_and_compute_shape..r/r8rrrrr5rrrcdS)NzMnon-empty 4D (batch mode) tensor expected for input with channels_last layoutr2r2r8r6rfz5max_pool2d_checks_and_compute_shape..scr8rcdS)Nz9non-empty 3D or 4D (batch mode) tensor expected for inputr2r2r8r6rfz5max_pool2d_checks_and_compute_shape..Or8FcdS)NzAUnsupported memory format. Supports only ChannelsLast, Contiguousr2r2r8r6rfz5max_pool2d_checks_and_compute_shape..sWr8) rWrhrrrQrrrrr9r:)rrrrrrr*r=r>r?r@rArBrarbrrCrDrrrs r6rrsVM; / /FB L F y aa 6{{aRB&))B 7++JD$!6*h77Iy**R..K**R..KBJ/66M+++ IIKK1  c c     %1 1 1 IIKK6 ! O O      W W   ( Rr9iXXL&z2tRIVVK       $  k 11r8cP t|||||\} tjjjkfd| j fd} | | |t j} tjjjj | S)Nc(djdjS)NzExpected dtype z for `gradOutput` but got dtype rrsr6rfz7meta_max_pool2d_with_indices_backward..sa$*aakN_aar8ct|dz t|dz t|dz dS)Nr.rr1)r)rrIrrrs r6_check_dim_sizez>meta_max_pool2d_with_indices_backward.._check_dim_sizesQq$q,777q$q,777q$q+66666r8rz) rrWrhr_rrQrrrr)rrrrrrrrrrrrIrrrs`` @@@@r6%meta_max_pool2d_with_indices_backwardrs , k67Hi     L k''aaaaa L 9D77777777 OK   OG/55M ; j{#    r8ct||||||\}}}|dkr|dnd} tj|} |dkr|||g} n| |||g} t j| |j|j| t j| t j |j| fSr) rrrrQrrWrr_rrrs r6meta_max_pool2d_with_indicesrs , {FGXy    %yy{{a//UZZ^^^QF/66M yy{{a\;7 \;? +<'     +<'      r8c tjjdvfdj}t|dz |D]2 tj dk fd3tjt dkdtjt |dkdd }d  d  |d krd}nd }tjjjkdtjjdkfdd}d }d tj||kdtj||kdtj dk fdtj|ddzd z k fdtj|d d zd z k fdd kr|||d|d g} n||d|d g} tj| jj tj| tj j fS)NrcdjS)Nz:fractional_max_pool2d: Expected 3D or 4D tensor, but got: rrsr6rfz,meta_fractional_max_pool2d..@sXTYXXr8r.rc:dddS)Nz_fractional_max_pool2d: Expected input to have non-zero size for non-batch dimensions, but got rz emptyr)rrsr6rfz,meta_fractional_max_pool2d..Gs7^7;yy{{^^TU^^^r8rcdS)NzNfractional_max_pool2d: kernel_size musteither be a single int or tuple of Intsr2r2r8r6rfz,meta_fractional_max_pool2d..N 2r8cdS)NzOfractional_max_pool2d: output_size must either be a single int or tuple of Intsr2r2r8r6rfz,meta_fractional_max_pool2d..Srr8r5rrrr1cdS)Nz6Expect _random_samples to have the same dtype as inputr2r2r8r6rfz,meta_fractional_max_pool2d..asHr8cdjS)Nz1Expect _random samples to have 3 dimensions got, r)random_samplessr6rfz,meta_fractional_max_pool2d..esYNDWYYr8cdS)Nz=Expect _random_samples.size(0) no less then input batch size.r2r2r8r6rfz,meta_fractional_max_pool2d..mrr8cdS)Nz.qNr8cddS)Nz/Expect _random_samples.size(2) equals to 2 got .r2)rsr6rfz,meta_fractional_max_pool2d..ss!WST!W!W!Wr8c dddS)Nz%fractional_max_pool2d: kernel height rz' is too large relative to input height r2) input_heightrsr6rfz,meta_fractional_max_pool2d..ws} A}}o{}}r8c dddS)Nz$fractional_max_pool2d: kernel width r1z& is too large relative to input width r2) input_widthrsr6rfz,meta_fractional_max_pool2d..{sz{1~zzmxzzr8r) rWrhrrrrr_rrrr) rrr~rrinput_channels input_batchr crrrrs `` ` @@@r6meta_fractional_max_pool2dr<sQ L VXXXX 9D 4!8T " "   IIaLL1  ^ ^ ^ ^ ^     L KA 2 2  L KA 2 2 YYr]]N99R==L))B--K qyyiill   L n**HH Lq YYYY AAAAAA L [OO L ^NN LaWWWWXXX LAQ'!+|;}}}}} LAQ'!+{:zzzzz  xxzzQ^[^[^L A A?  *;    +;    r8cTtjt|dvd|d}t|dkr|n|d}t|dkr|n|d}tj| pt|dvd|s|n|d} |s|nt|dkr| n|d} |s|nt|dkr| n|d} tjt|dvd|d} t|dkr| n|d} t|dkr| n|d}tjt|dvd|d}t|dkr|n|d}t|dkr|n|d}tj|jd vd |jd kr|d nd}|d }|d}|d}|d}t ||| | ||}t ||| | ||}t |||| ||}t |||||| | | | | ||||||||||d|jd ko!tj|tj k}|jdkrQ| d}| o| tj }||||f}n|||||f}| |}| |tj }|r@|tj }|tj }||fS)NrTcdSNzMmax_pool3d: kernel_size must either be a single int, or a tuple of three intsr2r2r8r6rfz.meta_max_pool3d_with_indices.._r8rr1rcdSNzQmax_pool3d: stride must either be omitted, a single int, or a tuple of three intsr2r2r8r6rfz.meta_max_pool3d_with_indices..cr8cdSNzImax_pool3d: padding must either be a single int, or a tuple of three intsr2r2r8r6rfz.meta_max_pool3d_with_indices..[r8cdSNzJmax_pool3d: dilation must be either a single int, or a tuple of three intsr2r2r8r6rfz.meta_max_pool3d_with_indices..ryr8r^cdSr`r2r2r8r6rfz.meta_max_pool3d_with_indices..rar8r-rr5rrzmax_pool3d_with_indices()rrr)rWrhrrrr9rdrQrrrr`rrr,)rrrrrrrer=r>rfr?r@rirjrkrlrarbrrhrirjrkrlrmrnrinput_channels_last_checkrrrs r6meta_max_pool3d_with_indicesrs L KF"__ QB;1$$+a.B;1$$+a.B L  +c&kkV+cc (vayB Fc&kkQ&6&6F1IB Fc&kkQ&6&6F1IB L G [[ B7||q  gajB7||q  gajB L H \\ I ]]a// Xa[I ]]a// Xa[I L fKK %zQUZZ^^^AFjjnnG JJrNNEjjnnG ZZ^^F BIy I IE"7BB 9MMG !&"b"i K KF            #+2  aXE7>>%BXX zQ$)OOA$6$6!)7799 9 '5506    eWf5 WeWf= //) $ $Cooiu{o;;GCff5#9f::**5+A*BB <r8ctjt|dvd|d}t|dkr|n|d} t|dkr|n|d} tj| pt|dvd|s|n|d} |s| nt|dkr| n|d} |s| nt|dkr| n|d} tjt|dvd|d}t|dkr|n|d}t|dkr|n|d}tjt|dvd|d}t|dkr|n|d}t|dkr|n|d}tj|jd vd |d }|d }|d }|d}|d }|d }|d}t |||||| | | | | ||||||||||||d|jdko!t j|tjk}|jdkrJ| d}| o| tj}| |j }|r | tj}|S)NrTcdSrr2r2r8r6rfz7meta_max_pool3d_with_indices_backward..rr8rr1rcdSrr2r2r8r6rfz7meta_max_pool3d_with_indices_backward..rr8cdSrr2r2r8r6rfz7meta_max_pool3d_with_indices_backward..rr8cdSrr2r2r8r6rfz7meta_max_pool3d_with_indices_backward.. ryr8r^cdSr`r2r2r8r6rfz7meta_max_pool3d_with_indices_backward..(rar8rr5rrz"max_pool3d_with_indices_backward()r-rr)rWrhrrrrzrQrrrr`rrr,)rrrrrrrrrer=r>rfr?r@rirjrkrlrarbrhrirjrkrlrmrnrrrs r6%meta_max_pool3d_with_indices_backwardrs L KF"__ QB;1$$+a.B;1$$+a.B L  +c&kkV+cc (vayB Fc&kkQ&6&6F1IB Fc&kkQ&6&6F1IB L G [[ B7||q  gajB7||q  gajB L H \\ I ]]a// Xa[I ]]a// Xa[I L fKK jjnnG JJrNNEjjnnG ZZ^^F   R Er""G   b ! !F#            ,/6  aXE7>>%BXX zQ$)OOA$6$6!)7799 9 '5506    --JI]]1G]HH r8gridc"tjjjkfdtjjtjkojtjkfdtjjdjdkfdtjjdjdz kfdtdjD]*tjjdkfd+dS) Nc(djdjS)NzNgrid_sampler(): expected input and grid to be on same device, but input is on z and grid is on rrrsr6rfz+check_grid_sampler_common..ds0 A\ A A37; A Ar8c(djdjS)NzTgrid_sampler(): expected input and grid to have torch.strided layout, but input has z and grid has )rrsr6rfz+check_grid_sampler_common..ks0 C C C59[ C Cr8rc(djdjS)NzZgrid_sampler(): expected grid and input to have same batch size, but got input with sizes  and grid with sizes r5rsr6rfz+check_grid_sampler_common..rs0 O %  O OBF* O Or8rrc.djdz djS)Nz+grid_sampler(): expected grid to have size rz, in last dimension, but got grid with sizes )rrrsr6rfz+check_grid_sampler_common..ys. ?%*q. ? ?26* ? ?r8c djddS)NzYgrid_sampler(): expected input to have non-empty spatial dimensions, but input has sizes rrr5rsr6rfz+check_grid_sampler_common..rr8)rWrhrrrrrr)rrrs``@r6check_grid_sampler_commonrast L  #      L  %F$+*F      L A$*Q-'      L 2%*q.(     1ej ! !   KNQ            r8ceZdZdZdZdZdS)GridSamplerInterpolationrr1rN)r{ __module__ __qualname__BILINEARNEARESTBICUBICr2r8r6rrsHGGGGr8rinterpolation_modectjjdkojjkfdtjjdko|tjjk ddS)Nr-c(djdjS)Nzdgrid_sampler(): expected 5D input and grid with same number of dimensions, but got input with sizes rr5rsr6rfz'check_grid_sampler_3d..s+ 149K 1 1$(J 1 1r8cdS)Nz.rr8)rWrhrrrr*)rrrs`` r6check_grid_sampler_3drs L a3EJ$)3      L J!O M"&>&F&LL  ON r8c|d}|r!tj|tj}nd}tj|tj} || fSNrr)rWrrr rrrr padding_mode align_cornersrinput_requires_gradr grad_grids r6grid_sampler_2d_backward_metarsX&a.%e5;RSSS   U5LMMMI  ""r8ct||t||||jd}|jd}|jd}|jd}|jd} |||||| fS)Nrr1rr.)rrrr) rrrrrrCout_Dout_Hout_Ws r6grid_sampler_3drs{eT***%'9::: AA AA JqME JqME JqME ??Aq%6 7 77r8rct||t||||d}|r!tj|tj}nd}tj|tj} || fSr)rrrWrr rrs r6grid_sampler_3d_backwardrseT***%'9:::%a.% !?     U5STTTI y  r8c|d}|stj|}||d<tj|g|Ri|S)Nr_)r^rQ get_dtyperWr)rrrSrr_s r6fullrsS JJw  E , ++F7O ;t -d - - -f - --r8cn|tjkrtj|dudtjd||jn||||jn||}|jrN||| | n;||| d| d|Stj||||||}|d|S)NcdS)Nz9memory format option is only supported by strided tensorsr2r2r8r6rfzzeros_like..rr8rrTr)rW sparse_coorhrr_r is_sparsesparse_resize_and_clear_r sparse_dim dense_dimr _coalesced_r+rr.fill_)rr_rrrrrs r6rrsA!!! T ! O O   k % $**5"(.4;;f!     > E  ( ( T__..0@0@      ( (dhhjj! D D D  / ! ! # "  CIIaLLL Jr8rc|tj}|tj}| tj}tj|||||SrrWrget_default_devicerrrr_rrrrs r6 meta_onesrZ }')) ~)++ ~ ; E&J   r8c|tj}|tj}| tj}tj|||||Srrrs r6 meta_zerosr-rr8c*tj|SrBrQclone_preserve_strides)rr/rrs r6meta_select_scatterrC  ' - --r8c*tj|SrBr)rr/rr}r|steps r6meta_slice_scatterrHrr8 dim_post_expr wrap_scalarc|dkr|sJd}| }|dz }||ks||krJd|d|d|d|dkr||z }|S)Nrr1zdim z out of bounds (rxryr2)rrrrrs r6rrNs~ .C ! Cc S3YYY)R)R)RS)R)RC)R)R)RYYY Qww } Jr8cP|dkrdn |j|Srr)rrs r6ensure_nonempty_sizerZs"1 11!'#,.r8cptd}td}tj||kdt |D]Ckr;tjt t kfdDdS)Nr1cdS)NzDIndex tensor must have the same number of dimensions as input tensorr2r2r8r6rfz$gather_shape_check..dsVr8c:ddjdjdzS)Nz!Size does not match at dimension z expected index  to be no larger than self  apart from dimension r5)rrrrsr6rfz$gather_shape_check..js6\A\\u{\\W WWRUWWXr8)rrrWrhrr)rrr self_dims index_dimsrs``` @r6gather_shape_checkr_sDHHJJ""IUYY[[!$$J LZVV9   88 L$UA..2FtQ2O2OOXXXXXXX   r8cpddlm}t||}|dk}|sQt jjt jkpjt j kfdt||| j S)NrrcdjS)Nz8gather(): Expected dtype int32/int64 for index, but got rrsr6rfzmeta_gather..xs\u{\\r8) r)rrrrrWrhr_rrrrr)rrr sparse_gradr wrapped_dimis_index_emptys ` r6 meta_gatherr osDDDDDD dhhjj11K#^EKKMMQ$677N 5 K5: % A )A \ \ \ \    4e444 >>%+ & &&r8c|r@|dkrdS|dkrdS|dkrdS|dkrdS|d krd Stjd d dS|d krdS|dkrdStjd ddS)Nr REDUCE_ADDrREDUCE_MULTIPLYmean REDUCE_MEANamaxREDUCE_MAXIMUMaminREDUCE_MINIMUMFcdS)Nz=reduce argument must be either sum, prod, mean, amax or amin.r2r2r8r6rfz#get_operator_enum..sSr8addmultiplycdS)Nz/reduce argument must be either add or multiply.r2r2r8r6rfz#get_operator_enum..s$Ur8r_)reduce_use_new_optionss r6get_operator_enumrs e  <   $$    =   ##   ##  S S     e  <  " "$$ UUUVVVr8c*ddlm}||dkr@tj|jtjkp|jtjkfd|(tj|j|jkfddSdS)Nr)rQcdS)Nz((): Expected dtype int32/int64 for indexr2 method_namesr6rfz,scatter_gather_dtype_check..s{LLLr8cdS)Nz0(): Expected self.dtype to be equal to src.dtyper2rsr6rfz,scatter_gather_dtype_check..s{TTTr8)r)rQrrWrhr_rr)r rrsrc_optrQs` r6scatter_gather_dtype_checkr#sCCCCCC}U[[]]a'((  K5: % A )A L L L L    J'- ' T T T T     r8c"t|dSr0)rrs r6ensure_nonempty_dimr%s sA;;r8czddlm}|dkrdStjt t kdd}t }t|D]1}t|}|kr|t|krd}n2|s<:t|D]*}t|}|t|krd}n+vtjt t kdtj| fddStj| fddS) NrrcdSNzCIndex tensor must have the same number of dimensions as self tensorr2r2r8r6rfz%scatter_shape_check..rr8FTcdSr(r2r2r8r6rfz%scatter_shape_check..sYr8cDdjdjddjzS)NExpected index rrz and to be no larger than src r5)rrrr"sr6rfz%scatter_shape_check..s7ZekZZdjZZYsYY'-YYZr8c4djdjdzS)Nr+rrr5)rrrsr6rfz%scatter_shape_check..s,ZekZZdjZZ,s,,-r8) r)rrrWrhr%rrr) rrrr"ris_wrong_shaperr index_d_sizes ```` r6scatter_shape_checkr/s-DDDDDD~ekkmmq()) LDHHJJ''+>uyy{{+K+KKUU N#DHHJJ//I9  +E155 88  .tQ77 7 7!N E 8 g1y!!  A/q99L27A>>>>!%?  + +/B599;;/O/O O Y Y       Z Z Z Z Z Z Z        - - - - - -     r8ct||}td|||t|||||t ||dSdS)Nscatter)rrr#r/r)rrrr/rrr s r6scatter_meta_implr2se dhhjj11Ky$s;;;k5#666'?33333r8c\t||||d||jSNrr2rrrrrr/s r6meta_scatter_addr7s,dCU333 >>$* % %%r8c,t||||d|Sr4r2r6s r6meta_scatter_add_r:sdCU333 Kr8ct|tjr|nd}t|||||||jSrB)rqrWr r2rrrrr src_or_valuergr/s r6 meta_scatterr>sH%\5<@@ J,,dCdCV444 >>$* % %%r8cht|tjr|nd}t||||||SrB)rqrWr r2r<s r6 meta_scatter_r@s9%\5<@@ J,,dCdCV444 Kr8queryrr* dropout_p is_causalreturn_debug_maskr~c |d}|d}|d} |d} |d} tj|} tj||| ftj|j} |rX| dkrdnd}t j| |z }| dkrd}n| dkrd}tj||| |f|j|j}n!tjd|j|j}tj j rtj st|d krCtjd tjd }tjd tjd }nBtjdtjd }tjd tjd }| | dd| | |||f S) Nrr1rr.r@rr2r)rrWrrr[rrceilr_versionhiprrzrrr8)rBrr*rCrDrEr~r num_headsmax_seqlen_batch_qhead_dimmax_seqlen_batch_k attention logsumexp blocksize_c max_seqlen_k debug_maskseedoffsets r6(meta__scaled_dot_product_flash_attentionrXsAJ 1 IAzz!}}H! ''I Y 23k|I  L%]]cc y!3k!ABB  $ $LL 3 & &L[ $6 E+<   [%+elKKK  }DUZ4466D+e:L:LPU:U:U{2UZ???Ruz&AAA{Ael6BBBRu|FCCC      r8 res_shape.chtjkrtj}nt gdfddfdD}fdt t D}tj|jj  |}|S)N)rr1rr.c8|SrBr)idxrBs r6rfz,alloc_with_matching_layout..Zs%,,..*=r8Trc g|] }| Sr2r2)rNr\rYs r6rPz.alloc_with_matching_layout..\s>>>S)C.>>>r8c:g|]}|Sr2r)rNr dim_orders r6rPz.alloc_with_matching_layout..]s%KKK++KKKr8r) rgrrWrsortedrrrr_rr)rBrYrpermuted_shape final_permuter_s`` @r6alloc_with_matching_layoutrcRs U[Y&&u%% LL====t   ?>>>I>>>KKKKU3y>>5J5JKKK k %+el   '-   Jr8 attn_biascompute_log_sumexpc |d} |d} |d} |d} |d} | | | | f}t||}tj| | | dftj|j}tjdtjd}tjdtjd}||dd| | ||df SNrr1rrrr2rrrcrWrr[rr)rBrr*rdrerCrDrEr~rr(S_QS_KVD_VrYr logsum_exprVrWs r6(meta__scaled_dot_product_cudnn_attentionrmes  1 A 1 A **Q--C 88A;;D **R..CAsC I $UI 6 6C AsAk|J ;rF ; ; ;D [5:f = = =F         r8c |d}|d} |d} |d} |d} || | | f} t|| }tj|| | ftj|j}tjdtjd}tjdtjd}||dd| | ||df Srgrh)rBrr*rdrCrDrEr~rH_QrirjrkrYrrlrVrWs r65meta__scaled_dot_product_fused_attention_overrideablerps  1 A **Q--C **Q--C 88A;;D **R..CCc"I $UI 6 6C C k|J ;rF ; ; ;D [5:f = = =F         r8rrR cum_seq_q cum_seq_kmax_qmax_k philox_seed philox_offsetctj|}tj|}tj|}|||fSrBr)rrBrr*rrRrqrrrsrtrCrDrurvr~grad_qgrad_kr"s r6'meta__scaled_dot_product_flash_backwardrzsA, e $ $F  c " "F  e $ $F 66 !!r8 attn_maskc*|d}|d}|d} tj|} tj|| |ftj|jdd} | | fS)Nrr1rr)rrWrrr[rr) rBrr*rCrDr{r~rrMrNrQrRs r60meta__scaled_dot_product_flash_attention_for_cpur}sAJ 1 IA ''I    k|i1oo  r8c Dtj|d|j|j} tj|d|j|j} tj|d|j|j} | | | fS)Nrrr1r.r)rWempty_permutedrr_r) rrBrr*rrRrCrDr{r~rxryr"s r69meta__scaled_dot_product_flash_attention_for_cpu_backwardrs& ! k| F  !  iz F  ! k| F 66 !!r8 dropout_maskcd}|\||\} } ||\} } j\| j\} } } fd} fd}dks | krdkr |S| S)Nc|dkr|ddfS|dkrd}t|dz D]}||j|z}|||d|d|ddfS|d fS) Nr.rTrr1r5rrF)rrrrviewr)r;rrs r6 ensure_4dzBmeta__scaled_dot_product_attention_math_for_mps..ensure_4d.s 5577a<<;;q>>4' ' UUWWq[[J15577Q;'' ) )agaj( 66*affRjj!&&**affRjjII4O Oe8Or8cj} r|}f} rodkr|d}nAt jdd|jddz}||}||fS)Nr.rr5r1r)rrview_asrsqueezerr) rattnrrmax_seq_lengthnum_headq_q_sizerB unsqueezeds r6sdpa_vector_fast_mpszMmeta__scaled_dot_product_attention_math_for_mps..sdpa_vector_fast_mps@sll28$$  %++e$$C||Z6>JKK  (yy{{a||AU["-..AaC@yy''Dyr8cxd}j}|f}||fS)Nr*r)blocksrr1r head_sizerrrs r6sdpa_vector_2pass_mpszNmeta__scaled_dot_product_attention_math_for_mps..sdpa_vector_2pass_mpsNsBll28$$||Z669$UVV L  r8iir5)rBrr*r{rCrDrr~rk_rTv_k_sizerrrrrrrrrs` @@@@@@@r6/meta__scaled_dot_product_attention_math_for_mpsr#s   Yu%%NB IcNNEB Ie  EB.0h+J&)#%8 Av~q           !!!!!!!!! $FVOO$8N8N$$&&&##%%%r8c.|dd}|dd}|dd}|d}|d} |d} |d} tj|| | | |j|j} tjjr&tj r |r| nd} n|rtj | dz dznd} tj|| | ftj |j}| dd} tjdtj d }tjdtj d }| |||fS) Nr1rrrrrr*r2r)rrrWrr_rrKrLrrzrrJr[r)rBrr*rdrerCrDr~rrrMKvr logsumexp_dimrlrVrWs r6,meta__scaled_dot_product_efficient_attentionrZsv OOAq ! !E --1  C OOAq ! !E 1 A 1 A 2I BB +aIrU\ R R RC }LUZ4466L 06Q 2DK !b&))B..!  I}%k|J --1  C ;rF ; ; ;D [5:f = = =F  D& ((r8grad_input_maskc |d} |d}|d}|d}|d}|d}tj| |||fd|j|j}tj| |||fd|j|j}tj| |||fd|j|j}d}|| dr~|d}|dzdkr|n |dz|dzz }t |}||d<tj||j|j}|d d|f}||||fS) Nrr1rr.rrrr).)rrWrr_rrr)rrBrr*rdrrRrurvrCrrDr~rrMrsrO head_dim_vrtrxryr" grad_biaslastDimlastDimAligned new_sizess r6+meta__scaled_dot_product_efficient_backwardrs(AJ 1 I JJqMMEzz!}}HAJ HHQKKE  ! Yx0k| F  ! Yx0iz F  ! Yz2k| F I!3..$$$+bLA$5$57R<'TV,;V))** & " K Y_Y5E   c8G8m, 669 ,,r8ctj|}tj|}tj|}|||fSrBr)rrBrr*rrRrurvrdrqrrrsrtrCrDr~rxryr"s r6'meta__scaled_dot_product_cudnn_backwardrsA. e $ $F  c " "F  e $ $F 66 !!r8window_size_leftwindow_size_right seqused_k alibi_slopesc||dn|dz }||dn|}||dn|}|d}|d}tj|}|*tj|||ftj|j}n=|d}tj||ftj|j}| rX|dkrdnd}tj||z }|dkrd}n|dkrd}tj||||f|j |j}n!tjd|j |j}d \}}tj j ratj rCtjd tjd }tjd tjd }nBtjd tjd }tjd tjd }|||||fS) Nrr1rrrrGrHrINNr2rr)rrrWrrr[rrrJr_rKrLrrzrr8)rBrr*rqrrrsrtrCrDrEr~rrrrrrNrPrMrOrQrRtotal_qrSrTrUrVrWs r6meta__flash_attention_forwardrs?4#,"3A9J9JQ9NJ*3*;A(1(9!u 2Izz"~~H ''IK $6 7+<   **Q--K   EL    L%]]cc y!3k!ABB  $ $LL 3 & &L[ $6 E+<   [%+elKKK LD& }DUZ4466D{2UZ???Ruz&AAA{Ael6BBBRu|FCCC   r8ctj|}tj|}tj|}|||fSrBr)rrBrr*rrRrqrrrsrtrCrDrurvr~rr grad_querygrad_key grad_values r6meta__flash_attention_backwardr0sA0!%((J$$H!%((J x ++r8 cu_seqlens_q cu_seqlens_k max_seqlen_qrTcustom_mask_typecausal_diagonalseqlen_k window_sizec|d}|d}|d}|d}|d}tj|||||j|j}||ddz n|}|}||J|}||n|}| rt j|dz dznd}tj|||ftj|j}tjdtjd}tjdtjd}||||||fS) Nrr1rrrr*r2r) rrWrr_rrrJr[r)rBrr*r@rrrrTrCrrer~rrrrrrrMrrlogsumexp_batch_dimactual_max_seqlen_qactual_max_seqlen_krrlrVrWs r6!meta__efficient_attention_forwardrOsZ,  1 A 1 A  A 2I BB +aIrU\ R R RC7C7O,++A..22VW'''**6*B,,4FM %*++b00A i7k|J ;rF ; ; ;D [5:f = = =F  D&*=?R RRr8bias_requires_gradnum_splits_keyshared_storage_dqdkdvc|rtj|jd|jdkdtj|jd|jdkdtjg|jddd|jd|jdR|j|j}|d d}|d d}|d d }n.sBr8r.cdS)Nz3embedding dim must match for `shared_storage_dqdkdvr2r2r8r6rfz4meta__efficient_attention_backward..Ir8rrrrr5rr).r2r) rWrhrrr_rrrrr)rrBrr*r@rrrrTrRrCrurvrrr~rrchunkrrrrrrrs r6"meta__efficient_attention_backwardrs2- KNcil * B B     KNcil * I I    Eek!B$ E EEKO EU[_ E E+<    \\"a(( <<A&&\\"a(( %e,, #C((%e,,  ))B--$+bLA$5$57R<'TV,;V%% & " K DKPPP c8G8m, K5<888 xY 66r8scale_ascale_b scale_resultuse_fast_accumcr d}tjdkodkfdtj|jo|jfdt dkst dkr9d} d} d } tj| p | fd tj| p | fd tjd d zdkfdtjdd zdkod d zdkfdj\d jtjkojtjkp)jtj kojtj k}  d krX d kr@tjjtj kojtj kdnQ| rjtj krd } dznd} d}t| }t|ddz}|t|z|z|t|z|z krj krRtj dtj dn\tjdfdn?tjjtj kojtj kdtjdkodkfddkrd d krpdd krWd kr>tj o dndkrld dcxkrtdkr+nn(d tdkrndkr^d dcxkrtdkrnnd krntjdfd||nj}tjdd |jS)Ncx|tjtjtjtjtjfvSrBrWr_ float8_e5m2float8_e4m3fnuzfloat8_e5m2fnuzfloat4_e2m1fn_x2rs r6is_fp8_or_fp4_typez2_check_scaled_mm_sizes..is_fp8_or_fp4_type2      !  !  "    r8rc\ddSNz%Inputs must be 2D but got self.dim()=z and mat2.dim()=rrMrsr6rfz(_check_scaled_mm_sizes..-` ``TXT\T\T^T^``r8c(djdjSNz?Expected both inputs to be fp8 or fp4 types but got self.dtype=z and mat2.dtype=rrsr6rfz(_check_scaled_mm_sizes.. zRVR\zznrnxzzr8rrc>|d|dko |ddkSrr2rs r6 is_row_majorz,_check_scaled_mm_sizes..is_row_major"!9vay(;VAY!^ ;r8c2|ddko |ddkSrr2rs r6 is_col_majorz,_check_scaled_mm_sizes..is_col_major!9>3fQi!m 3r8cf|ddkp|ddkSrr tensor_2ds r6 has_zero_dimz,_check_scaled_mm_sizes..has_zero_dim/>>!$$)CY^^A->->!-C Cr8c2dSNz#self must be row_major, got stride rrsr6rfz(_check_scaled_mm_sizes..I$++--IIr8c2dSNz#mat2 must be col_major, got stride rrMsr6rfz(_check_scaled_mm_sizes..rr8r1r)rc4ddSNzBExpected self.size(1) to be divisible by 16, but got self.size(1)=r1rrsr6rfz(_check_scaled_mm_sizes.. gY]YbYbcdYeYeggr8cdjSNz?Expected both dimensions of mat2 to be divisible by 16 but got r5rsr6rfz(_check_scaled_mm_sizes..bVZV`bbr8cdS)NzNFor tensorwise scaling, both scale_a and scale_b must be float (fp32) tensors.r2r2r8r6rfz(_check_scaled_mm_sizes.. shr8r*rHrcdS)Nzscale_a must be contiguousr2r2r8r6rfz(_check_scaled_mm_sizes..+8r8cdS)Nzscale_b must be contiguousr2r2r8r6rfz(_check_scaled_mm_sizes../rr8Fc jddddd S)NzTInvalid blockwise scaling configuration. For blockwise scaling, scale_a should have  elements, got z, scale_b should have rr)expected_a_sizeexpected_b_sizerrsr6rfz(_check_scaled_mm_sizes..4sdbFUbbfmfsfsfufubb/>bbOV}}bbbr8cdS)NzKFor rowwise scaling, both scale_a and scale_b must be float (fp32) tensors.r2r2r8r6rfz(_check_scaled_mm_sizes..=ser8c\ddS)NzLFor non-tensorwise scaling, scale tensors must be 2D, but got scale_a.dim()=z and scale_b.dim()=rrrsr6rfz(_check_scaled_mm_sizes..Bs2~Y`YdYdYfYf~~nunynyn{n{~~r8cdS)Nz@Both scale_a and scale_b must be contiguous for rowwise scaling.r2r2r8r6rfz(_check_scaled_mm_sizes..Ns^r8cddddtdd dtddtdddtdd zdtddd d dd d d dd d zSNz}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>). For (BlockWise1x128, BlockWise128x128), scale_a should be (rxrHz), zscale_b should be (z<). For (BlockWise1x128, BlockWise1x128), scale_a should be (z). Got scale_a.size()=(rr1z) and scale_b.size()=(ryr7r)_krbr rrsr6rfz(_check_scaled_mm_sizes..bsmrCDrrbcrrWXrr\ddfhk[l[lrrrpS0A0AppXaQTEUEUppTUppYabdfiYjYjppp p US0A0AUUQUU/6||AUUBI,,q//UU/6||AUUBI,,q//UUU Ur8r)rWrhrr_rrrrfloat8_e8m0fnur_rr\r7r`rr)rrMrrr@rrBrrrrris_blockwise_scaling block_size_k block_size_mn num_k_blockspadded_num_k_blocks _out_dtyperrrrbr s```` @@@@@r6_check_scaled_mm_sizesr s    L  a+DHHJJ!O````` L4:&&I+=+=dj+I+Izzzzz 4F""k$&7&75&@&@ < < < 4 4 4 D D D  L ' ' =<<+=+= I I I I     L ' ' =<<+=+= I I I I     IIaLL2  " g g g g     IIaLL2  " =tyy||b'8A'= b b b b    2 IIaLL !55:MU%99   !449MU%88  ==??a  GMMOOq$8$8 L .Q7=EM3Qhh    "` } 333 " !V! M#B 55L"*<";";a"? M : ::=PP M : ::=PP   ?22MMOO66 ))++88 ))++88   L .Q7=EM3Qee    L "9w{{}}'9~~~~~     Q1$$LLOOq((LLOOq((LLOOq(( ))++G0E0E0G0G^^  Q1$$LLOOw||AKKKK(2s:K:KKKKKKLLOOx3'7'777 Q1$$LLOOw||AKKKK(2s:K:KKKKKKLLOOq((            (3J ;tyy||TYYq\\DK X X XXr8c .t||||||||SrB)r )rrMrrr@rrBrs r6meta_scaled_mmr ss' " dGWdL)^  r8scale_recipe_ascale_recipe_b swizzle_a swizzle_bc   !"#$d} d} tjdkodkfdtj| jo| jfdjdjdjd| jr| jrd} | zd|D}d |D}r d Dn t jg r d D n t jg td kstd k rUd}d}d}tj|p |fdtj|p |fdtj ddzdkfdtj ddzdko ddzdkfddttdttfd}dttdttfd}dttdttfd}dttdttfd}dttdttfd}dttdttfd}dttdttfd}|||rtjd dkoSd dko5djtj kodjtj kdn|||rtjdjdkoqd koSdjtj ko8d kodjtj kfd n|||rjdjtj kodjtj k}d##jdkoo#jdd!zko[#ddkoB#dkp)#jddko#ddk}d$$jdkoo$jdd!zko[$ddkoB$dkp)$jddko$ddk}tj|o|o|#$fd"n}|||r~djtj kodjtj k}td!z d#d##jdkoo#jdd!zko[#ddkoB#dkp)#jddko#ddk}d$$jdkoo$jdd!zko[$ddkoB$dkp)$jddko$ddk}tj|o|o|#$fd$n|||r~djtj kodjtj k}td!z d#d##jdkoo#jdd!zko[#ddkoB#dkp)#jddko#ddk}d$$jdkoo$jdd!zko[$ddkoB$dkp)$jddko$ddk}tj|o|o|#$fd%ng|||rtjjr_t#jdd&jdz t#jdd&jdz!t j"ntjdd!tt#jdd&d#z tjdd!tt#jdd&d#z!t j"tjd kokdjtjkoPd !ko2djtjkod"ko d"k !" fd'n|||rtd!tt#dd#z td!tt#dd#z!t j"tjd ko݉djtjko‰d dkodjtj kod !kokdjtjkoPd dko2djtj kod"ko d"k !" fd(ntjd)fd*||nj}tj|j+S),Ncx|tjtjtjtjtjfvSrBrrs r6rz5_check_scaled_mm_sizes_v2..is_fp8_or_fp4_typerr8c |tjfvSrB)rWrrs r6 is_fp4_typez._check_scaled_mm_sizes_v2..is_fp4_types/111r8rc\ddSrrrsr6rfz+_check_scaled_mm_sizes_v2..rr8c(djdjSrrrsr6rfz+_check_scaled_mm_sizes_v2..rr8rr1c,g|]}t|Sr2r&rNsis r6rPz-_check_scaled_mm_sizes_v2..???"k"oo???r8c,g|]}t|Sr2rrs r6rPz-_check_scaled_mm_sizes_v2..rr8c,g|]}t|Sr2r'rs r6rPz-_check_scaled_mm_sizes_v2..999[__999r8c,g|]}t|Sr2rrs r6rPz-_check_scaled_mm_sizes_v2..r r8rrc>|d|dko |ddkSrr2rs r6rz/_check_scaled_mm_sizes_v2..is_row_majorrr8c2|ddko |ddkSrr2rs r6rz/_check_scaled_mm_sizes_v2..is_col_majorrr8cf|ddkp|ddkSrrrs r6rz/_check_scaled_mm_sizes_v2..has_zero_dimrr8c2dSrrrsr6rfz+_check_scaled_mm_sizes_v2..rr8c2dSrrrsr6rfz+_check_scaled_mm_sizes_v2..rr8r)c4ddSrrrsr6rfz+_check_scaled_mm_sizes_v2..rr8cdjSrr5rsr6rfz+_check_scaled_mm_sizes_v2..rr8recipe_arecipe_bct|dko>t|dko+|dtjko|dtjkSr)rr& TensorWiser)r*s r6 is_tensorwisez0_check_scaled_mm_sizes_v2..is_tensorwisesRH ":MMQ&:QK;#99:QK;#99  r8ct|dko>t|dko+|dtjko|dtjkSr)rr&RowWiser-s r6 is_rowwisez-_check_scaled_mm_sizes_v2..is_rowwisesRH "7MMQ&7QK;#667QK;#66  r8ct|dko>t|dko+|dtjko|dtjkSr)rr& BlockWise1x32r-s r6is_mxz(_check_scaled_mm_sizes_v2..is_mxsRH "=MMQ&=QK;#<<=QK;#<<  r8ct|dkojt|dkoW|dtjkoA|dtjko+|dtjko|dtjkS)Nrrr1)rr& BlockWise1x16r,r-s r6is_nvz(_check_scaled_mm_sizes_v2..is_nvs~H ":MMQ&:QK;#<<:QK;#99:QK;#<< : QK;#99  r8ct|dko>t|dko+|dtjko|dtjkSr)rr&BlockWise1x128r-s r6is_1x128_1x128z1_check_scaled_mm_sizes_v2..is_1x128_1x128sRH ">MMQ&>QK;#==>QK;#==  r8ct|dko>t|dko+|dtjko|dtjkSr)rr&r9BlockWise128x128r-s r6is_1x128_128x128z3_check_scaled_mm_sizes_v2..is_1x128_128x128sUH "@MMQ&@QK;#==@QK;#??  r8ct|dko>t|dko+|dtjko|dtjkSr)rr&r<r9r-s r6is_128x128_1x128z3_check_scaled_mm_sizes_v2..is_128x128_1x128sRH ">MMQ&>QK;#??>QK;#==  r8cdS)Nz\For Tensorwise scaling, both scale_a and scale_b must be single element float (fp32) tensorsr2r2r8r6rfz+_check_scaled_mm_sizes_v2.. svr8c djddddjdddd S)Nz'For Rowwise scaling, scale_a must have rz elements (got: z), and scale_b must have r1ry)rr)rMrrrsr6rfz+_check_scaled_mm_sizes_v2..)s~ddjmdd]def]g]m]m]o]odd/3z!}ddNUVWjN^N^N`N`dddr8rHczdddzdjddjdddzdjddjdS)Nz>For 1x128 x 1x128 blockwise scaling, scale a must have shape [rxrH] (got: ) and stride [1, )scale b must have shape [ryrr)rorrsasbsr6rfz+_check_scaled_mm_sizes_v2..Fsz01zz56#XzzGIxzzbczzmomvzz01zz56#XzzGIxzzbczzmomvzzzr8rczdddzdjddjdddzdjddjdS)Nz]For 128x128 x 1x128 blockwise scaling, L4 = {round_up(K / 128, 4)}, scale a must have shape [rxrHrCrDrEryrFroL4rrrGrHsr6rfz+_check_scaled_mm_sizes_v2..esz02zz673hzzHJzzcezzoqoxzz01zz56#XzzGIxzzbczzmomvzzzr8czdddzdjddjdddzdjddjdS)Nz]For 1x128 x 128x128 blockwise scaling, L4 = {round_up(K / 128, 4)}, scale a must have shape [rxrHrCrDrEryrFrJsr6rfz+_check_scaled_mm_sizes_v2..s|01||56#X||GIx||bc||momv||02||673h||HJ||ce||oqox|||r8r*cdddddddtjddjddjdd dddd S) Nz!for MX scaling scale_a must have  (got: r) and scale_b must have z). Scales must have types z (for self: , mat_b: z) Must have swizzle type (got self: ry)rrWrr_expected_scale_a_elemsexpected_scale_b_elemsexpected_swizzlerrrrsr6rfz+_check_scaled_mm_sizes_v2..ss8NssW^_`WaWgWgWiWiss-CssLSTUJL\L\L^L^ss"'"6ssDKAJDTss_fgh_i_oss/?ssMVVWLssdmmncosssr8cddddddddddddS) Nz!for NV scaling scale_a must have rNrrOz). Must have swizzle type rQrPryrrRsr6rfz+_check_scaled_mm_sizes_v2..si8NiiW^_`WaWgWgWiWiii-CiiLSTUJL\L\L^L^ii$4iiBKA,iiYbcdYeiiir8Fcddddtdd dtddtdddtdd zdtddd d d dd d d d d dd d d zSrr)rorrrrsr6rfz+_check_scaled_mm_sizes_v2..sm?@mm^_mmSTmmX``acfWgWgmmmkHQ,<,<kkC@P@PkkPQkkU]^_adUeUekkk k WHQ,<,<WWWW+21:??1+=+=WWAHQRASASWW+21:??1+=+=WWAHQRASASWWW Wr8r)rWrhrr_rr' NO_SWIZZLErrrrr&rr\r=rKrLr7SWIZZLE_32_4_4rr_rr)%rrMrrrrr@rBrrrrrK_packed_multiplierrrrr.r1r4r7r:r=r?types_ok scale_a_ok scale_b_okr rorKrrrSrTrUrGrHs%``` ` `` @@@@@@@@@r6_check_scaled_mm_sizes_v2r^s   222 L  a+DHHJJ!O````` L4:&&I+=+=dj+I+Izzzzz 1 A 1 A 1 A{4:!;;tz#:#:!  ?????N?????N 99y999   "  99y999   " 4F""k$&7&75&@&@ < < < 4 4 4 D D D  L ' ' =<<+=+= I I I I     L ' ' =<<+=+= I I I I     IIaLL2  " g g g g     IIaLL2  " =tyy||b'8A'= b b b b    D$5 kAR     k!2 d;>O     D- k9J     D- k9J     T+%6 ${BS     tK'8 DDU     tK'8 DDU     = 8 8p  L   ""a'6AJ$$&&!+6AJ$ 56AJ$ 5vv     Z 7 7g  L  #q(6AJ$$&&!+6AJ$ 56AJ$$&&!+6AJ$ 5    ^NN ; ;[    EM1Wgaj6F%-6W B q THQK18+TIIaLLA%TYYq\\Q&R28A;!+;+Q ! PQ@Q  B q THQK18+TIIaLLA%TYYq\\Q&R28A;!+;+Q ! PQ@Q  L6Z6J     nn = =}    EM1Wgaj6F%-6W !c'1%%BB r!UHQK18+UIIaLLA%UYYq\\R'SBHQK1,<,R1QRAR  B q THQK18+TIIaLLA%TYYq\\Q&R28A;!+;+Q ! PQ@Q  L6Z6J     nn = =^    EM1Wgaj6F%-6W !c'1%%BB q THQK18+TIIaLLA%TYYq\\Q&R28A;!+;+Q ! PQ@Q  B r!UHQK18+UIIaLLA%UYYq\\R'SBHQK1,<,R1QRAR  L6Z6J    U>> 2 2? } >)1$*Q-)D)DtzRS})T&)1$*Q-)D)DtzRS})T&#.#9  )1$*Q-)E)ETZ]B//II*&*2$*Q-)E)ETZ]B//II*&$/#= L   ""&<<5AJ$(<<5AJ$$&&*@@5AJ$(<<5aL$44 5 aL$44    U>> 2 2# %-a%5%5!RRS8T8T%T "%-a%5%5!RRS8T8T%T "*9  L   ""&<< 5AJ$(;; 5AJ$$&&!+ 5AJ$ 5 5AJ$$&&*@@ 5 AJ$(;; 5 AJ$$&&!+ 5AJ$ 5 5aL$44 5aL$44    $ L           (3J ;q!:dk B B BBr8rcontraction_dimsc 6t|||||||| |||  S)N)r@rBrrr)r^) rrMrrrrrrr@rr_rs r6meta_scaled_mm_v2ras< %   %    r8c`t|||||d||jSNT)rr5rrrr/rgres r6meta_scatter_reduce_twores3dCVTJJJJ >>$* % %%r8c0t|||||d|Srcr9rds r6meta_scatter_reduce__tworgs"dCVTJJJJ Kr8cdtjdcxkodkncfddkr&tj|tjjStjd|tjjS)Nrrc2dS)NzAThe probability distributions dimensions must be 1 or 2, but got rrsr6rfz"meta_multinomial..saTYT]T]T_T_aar8r1r)rWrhrrrrr)r num_samples replacementr s` r6meta_multinomialrls L EIIKK1aaaa yy{{a{;ejNNNN ; 1 {%*U\   r8cd}|D]}||z}|Sr0r2)vsr]vs r6multiply_integersrps& A  Q Hr8ctjtkfddztjtkfdtjtdddDotdDfddd\}}||gRS)Nc.ddtS)Nz%It is expected output_size equals to , but got size r)num_spatial_dimsr~sr6rfz'upsample_common_check..s#k8HkkY\]hYiYikkr8rc.ddtS)Nz$It is expected input_size equals to rsr)expected_input_dimsrQsr6rfz'upsample_common_check..s#l7Jll[^_i[j[jllr8c3"K|] }|dkV dSrNr2rNrs r6rtz(upsample_common_check..s& * *aAE * * * * * *r8c3"K|] }|dkV dSrxr2rys r6rtz(upsample_common_check..s&2N2NQ1q52N2N2N2N2N2Nr8cddS)NzDInput and output sizes should be greater than 0, but got input size z and output size r2)rQr~sr6rfz'upsample_common_check..s)A AA3>AAr8)rWrhrr)rQr~rtrchannelsrvs``` @r6upsample_common_checkr}s L K,,kkkkk+Q. L J..lllll  L * *:abb> * * ***Ns2N2N+2N2N2N/N/N A A A A A ""1"~FH H +{ + ++r8cptjdkp(tddfdt |d}|tj S)Nrr1c2dS)Nz>Non-empty 3D data tensor expected but got a tensor with sizes rrsr6rfz$upsample_nearest1d..*_QVQ[Q[Q]Q]__r8rtr rWrhrrprr}rr,rQr)rr~scalesfull_output_sizes` r6upsample_nearest1dr$ L A/ QRR0@AA____- kA ??+ , , / /1%88 0  r8ctjdkp(tddfdt |d}|}tj}j \}}}}j j dkr|dkr tj }| |}|S) Nrr1c2dSNz>Non-empty 4D data tensor expected but got a tensor with sizes rrsr6rfz$upsample_nearest2d..:rr8rrrrr)rWrhrrprr}rrQrrrrzr contiguous) rr~scales_hscales_wrrrrT n_channelss ` r6upsample_nearest2dr4s L A/ QRR0@AA____- kA__- . .F/66M +Az1a |F""zA~~/   ]  ; ;F Mr8r~rQrrcvt||dtjjdkfdt dD]9tjkfd:|tj S)NrrrcdjSNzFExpected grad_output to be a tensor of dimension 4 but got: dimension rrsr6rfz-upsample_nearest2d_backward..`kYdYikkr8c RddddS)NzCExpected grad_output to have the same shape as output; output.size() = z but got grad_output.size(rrrrsr6rfz-upsample_nearest2d_backward..es]J !JJ'7':JJ-.JJ4?4D4DQ4G4GJJr8r) r}rWrhrrrrr,rQr)rr~rQrrrrs` @@r6upsample_nearest2d_backwardrNs-K! LAkkkk1XX     Q  #3A#6 6             , , / /1+>> 0  r8cptjdkp(tddfdt |d}|tj S)Nrr1c2dS)Nz>Non-empty 5D data tensor expected but got a tensor with sizes rrsr6rfz$upsample_nearest3d..wrr8r.rrr)rr~scales_drrrs` r6upsample_nearest3drqrr8ctj|tj|tj}}||t|tsJt|tsJ|j}|} t||}t||}||| ||| t||t||||fS||fS)Nr)rr) rWrrrqrrrrrr!) rstabler descendingrrrorr out_strides r6 meta_sortrs  D ! !5#3D #L#L#LqA g1&*-----':.....G XXZZ "6955#GY779j111Iz222F3333G4444w a4Kr8c tjjdkfdtjjjkfddwtjjdkfdtjkfdtjjjkfdtjjdkfdd zztjkfd tjt fd fDd dS) NrcjdSNz != 2r) input_gatessr6rfz%rnn_cell_checkSizes..;3C0J0J0Jr8c&jdjSN != r5) hidden_gatesrsr6rfz%rnn_cell_checkSizes..s;$>>,*<>>r8r1cjdS)Nz != 1r) input_biassr6rfz%rnn_cell_checkSizes..sjo3L3L3Lr8c6dSrr) gates_sizersr6rfz%rnn_cell_checkSizes..s!z''));;z;;r8c&jdjSrr5) hidden_biasrsr6rfz%rnn_cell_checkSizes..sz'@@[->@@r8cjdSrr) prev_hiddensr6rfz%rnn_cell_checkSizes..rr8rc pdddddd S)Nrrz * z // z (aka ry)rr)expected_prev_hidden_numelfactorrrrsr6rfz%rnn_cell_checkSizes..sL;$$&&K,<,.sC   H * *      r8cdS)Nz%expected all inputs to be same devicer2r2r8r6rfz%rnn_cell_checkSizes..s7r8)rWrhrrrrr)rrrrrrrrs``````@@r6rnn_cell_checkSizesrs L!Q&(J(J(J(JKKK L\//>>>>>!!!$$J Z_)+L+L+L+LMMM     * , ; ; ; ; ;       1 1 @ @ @ @ @    L!Q&(J(J(J(JKKK!,!1!1!!4!4z!AV!K L99 L     #J [I     87 r8ct||||d|tj|tj}tj|tj}tj|tj}|||fS)Nrr)rrWrr)rrcxrr workspacehycys r6_thnn_fused_lstm_cell_metarsq \:{ArRRR Eweight_stride0 weight_bufhxrrY hidden_size proj_size num_layers batch_firstdropouttrain bidirectional batch_sizes dropout_stateis_input_packed seq_length mini_batchbatch_sizes_sumnum_directionsout_sizerr cell_shaperr reserve_shapereserves r6 _cudnn_rnnrs&+&&!+O%%  ^ +a.'2FU[^^ A '2FU[^^ A '.QQQN%NNyy H $h&?@  EZX%> ? ?j(^*CD  __Y ' 'F~-z;GJ z [5< 0 0 0 \\* % % zN2JI J JB%AAAMoom5;o??G 2r7J ..r8c|r |jdn |jd}|r |jdn |jd}| }|r|||gn|||g}||}|tjd|j}n||j}|tjd|j}n||j}tjd|jtj}||||fS)Nr1rrr)rrrWrrr)rw0w1w2w3hx_cx_rrrYrr has_biasesrrrrroutput_chanelsrrrrrs r6mkldnn_rnn_layerr s&$/BQEKNJ#.BQEKNJ N  6Z00*n 5 __Y ' 'F { [5< 0 0 0 ]]39 % % { [5< 0 0 0 ]]39 % % Ael%+FFFI 2r9 $$r8c|jdkr%tjdkpdkfddStj|dkfddS)NrrcdS)Nz4: Expected reduction dim -1 or 0 for scalar but got r2rrsr6rfz'zero_numel_check_dims..7swYYTWYYr8cddS)Nz: Expected reduction dim z to have non-zero size.r2rsr6rfz'zero_numel_check_dims..<swUUUUUr8)rrWrr)rrrs ``r6zero_numel_check_dimsr3s yA~~  1H !r Y Y Y Y Y      IIcNNa  U U U U U     r8c|5t||}t||dStj|dkfddS)NrcdS)Nz@: Expected reduction dim to be specified for input.numel() == 0.r2r$sr6rfz%check_argmax_argmin..Hst]]]r8)rrrrWrhr)rrrs` r6check_argmax_argminrAsq S$((**--dC..... JJLLA  ] ] ] ]     r8ctd||tj|j||fnd}t |||}||t jS)Nargmaxr)rrQrrrrrWr)rrrrrs r6argmax_argmin_metarLs[$,,,   coSFF4 P PD $T4 9 9E >>%u{> 3 33r8cl|tjkr tj}tjd||||S)Nr2r)rWjaggedrr)rr_rrrs r6 scalar_tensorrTs<  ; %v*   r8ct||d}|dkrdn||}tj|dktj||kdt |j}t|dkr|||<||||tj fS)NTrrr1cdS)Nzk not in range for dimensionr2r2r8r6rfztopk_meta..fs)Gr8r) rrrrWrhrrrrr)rrDrlargestr` sliceSizetopKSizes r6 topk_metar`s dhhjjd ; ; ;CXXZZ1__$))C..I La Li!G!GHHHDJH 8}}q >>( # #T^^HEK^%P%P PPr8c| | Jd|}|} tj|| j| j| jS)Nz;segment_reduce(): Either lengths or offsets must be defined)r_rr)rrWrr_rr) rJrrwrgrrrsrtrv data_contig grad_contigs r6meta__segment_reduce_backwardrnso  '"5"5E#6"5"5//##K//##K  !!    r8c4ddlm}t|d|dkr|nd}t j||dk||kfdt|jd|jdzdz}|r.|dkr| d| || |t j fS)Nr)r:Trr1cdS)Nz9kthvalue(): selected number k out of range for dimension r2rsr6rfzkthvalue_meta..sQCQQr8r) r)r:rrrrWrhrrr4rr)rrDrrr:dimSizers ` r6 kthvalue_metars>===== dhhjjd ; ; ;C $ QdiinnnAG LQW %%QQQQ DSD!DJsQwyy$99 : :E488::>> S! >>% $..ek."J"J JJr8c||n|}tj|dkd|}|,tj||kd|,tj||kdtj||kdtj||kdtj|dkdtj||d|d zd zkd dS) NrcdSNr2r2r8r6rfz(checkLSTMBackwardSizes..s"r8cdSrr2r2r8r6rfz(checkLSTMBackwardSizes..r8cdSrr2r2r8r6rfz(checkLSTMBackwardSizes..r r8cdSrr2r2r8r6rfz(checkLSTMBackwardSizes..r8cdSrr2r2r8r6rfz(checkLSTMBackwardSizes..r r8cdSrr2r2r8r6rfz(checkLSTMBackwardSizes..srr8rr1rcdSrr2r2r8r6rfz(checkLSTMBackwardSizes..sRr8)rWrhrrr)grad_hygrad_cyrrr defined_gradexp_sizes r6checkLSTMBackwardSizesr s2%177wL L!!##q(**555  ""H W\\^^x/<<< W\\^^x/<<< Lh& 333 Lh& 333 LA%zz222 L""hqkHQK&?!&CCZZPPPPPr8c||dSt|||||tj|t}tj|t}|r|ddnd}|||fS)NNNNrrF)r)r rWrlegacy_contiguous_memory_formatr) r r rrrhas_bias grad_gatesgrad_cxrs r6#_thnn_fused_lstm_cell_backward_implr s7?7GRY???!!@Jr1PQQQG4<F q%000$I w ))r8c`d}d}d}|dr'||}|ds|dre||d|df}||d}|||fS)Nrr1rrr)rrrrr grad_weightrs r6linear_backwardr sJKI1~;!++FKKMM:: 1~BQB",,l.?.?.C.CV[[QS__-UVV  **<+<+.is_channels_lasts!"88==ATTTr8c$r+tdkr tjStjStjr tjStjr tjSdS)Nrr)rrWrrr`preserve_format)rrsr6rz.meta_pixel_shuffle..pick_memory_formats  D ! ! )4  F**..**   e.E  F F )* *   e.C  D D )( ( ) )r8rrr)rrrr,) rupscale_factorrrHrWrrrrs ` @r6meta_pixel_shuffler s DJ! 2.>2Q RVW W W We$*eeUcee!X W W UUU ) ) ) ) ) ) 2>N:;A B. (B B. (B-$*SbS/-1-b-"--I .. # #C &&1133& 4 4C Jr8cL||j}||j}||j}||j}||j}||j}|||||||fSrBr)rweight0weight1weight2weight3rcx_tmprhy_cy_grad_output_r_opt grad_hy_r_opt grad_cy_r_optrrYrrrrrrrrdiff_xdiff_hxdiff_cxdiff_w1diff_w2diff_bs r6mkldnn_rnn_layer_backwardr/ s4__U[ ) )FmmCI&&Gv|,,G ..G ..G   w} - -F 7GVVWg EEr8) out_int32rcttj||r tjn tjtjS)Nr_r)rWrrFrr)r boundariesr0 rs r6meta_bucketizer4 s7   &7ekkEK-   r8dcdtdkr*tjfdtdkr(rt jdtjt tfdtjdkfdtjt tfd tjt tfd tjkfd tj j j S) Nzhistc()rcdjdS)Nz%"histogram_cpu" not implemented for 'rrrsr6rfzmeta_histc..sLekLLLr8rz%_histc_cuda with floating point inputc,dtS)Nz#: argument 'bins' must be int, not rbinsrsr6rfzmeta_histc..s7KKtDzzKKr8rcdS)Nz: bins must be > 0, but got r2r9 sr6rfzmeta_histc..sg#Q#Q4#Q#Qr8c,dtS)Nz%: argument 'min' must be Number, not r)rrsr6rfzmeta_histc..7LLcLLr8c,dtS)Nz%: argument 'max' must be Number, not r)rrsr6rfzmeta_histc..#r= r8cdS)Nz: max must be larger than minr2)rsr6rfzmeta_histc..%s%N%N%Nr8r) rrWrhrrQrrqrrrrr_)rr: rrrs````@r6 meta_histcr@ s|G5U""  # # % % L L L L   5V##(?(?(A(A# %&MNNN L4!!KKKKK LQQQQQRRR L3LLLLL L3LLLLL LNNNNOOO ;tEL D D DDr8ct|d}tjdkp2t dddDfd|tj S)Nrrrc3"K|] }|dkV dSrxr2)rNrs r6rtz,meta_upsample_bimode2d_aa..7s&!H!Ht$(!H!H!H!H!H!Hr8r1c2dSrrrsr6rfz+meta_upsample_bimode2d_aa..8rr8r) r}rrWrhrrrr,rQr)rr~rrrrs` r6meta_upsample_bimode2d_aarD )s- kA L Hc!H!Huzz||ABB7G!H!H!HHH____ ??+ , , / /1%88 0  r8cft||dtjjdkfdt dD]1tjjkfd2|tj S)NrrrcdjSrrrsr6rfz4meta_upsample_bimode2d_aa_backward..Mrr8c RddddS)NzD Expected grad_output to have the same shape as output; output.size(rz but got grad_output_size(rrsr6rfz4meta_upsample_bimode2d_aa_backward..RsR9DE99K[\]K^9999!,!1!1!!4!499r8r) r}rWrhrrrrr,rQr)rr~rQrrrrrs` @@r6"meta_upsample_bimode2d_aa_backwardrH ?s-K! LAkkkk1XX    a $4Q$7 7 9 9 9 9 9 9      , , / /1+>> 0  r8c6tj|dkdtj|dkdtj|jjdtj|jjddS)Nr1cdS)Nz%found_inf must be a 1-element tensor.r2r2r8r6rfz<_amp_foreach_non_finite_check_and_unscale_.._(Or8cdS)Nz%inv_scale must be a 1-element tensor.r2r2r8r6rfz<_amp_foreach_non_finite_check_and_unscale_..brK r8cdS)Nz!found_inf must be a float tensor.r2r2r8r6rfz<_amp_foreach_non_finite_check_and_unscale_..f3r8cdS)Nz!inv_scale must be a float tensor.r2r2r8r6rfz<_amp_foreach_non_finite_check_and_unscale_..jrN r8)rWrhrr_r)rr inv_scales r6*_amp_foreach_non_finite_check_and_unscale_rQ \s LQ O O LQ O O L)33 L)33r8c*tj|SrBr)rnanposinfneginfs r6 nan_to_numrV orvr8c|jtjtjtjtjhvsJd|jd|j}t||}t||}||kr|St| }t| }||||c||<||<||||c||<||<| |||S)Nz>torch.transpose_: in-place transposition is not supported for z layout) rrWr sparse_cscr sparse_bscrrrrrr)rdim0rendimsrrs r6rrus ;         ^]]]     IE $ & &D $ & &D t||   D $++-- F!'vd|F4L&,!$ZdDJT T6""" Kr8c(|j}|jrF|}|}|dkr|dksJd|d|dn&|dksJd|dt |d|dkrdndS) NrrzEt_ 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 r1)rrrrrr)rr[ rrs r6t_r] s IE ~  __&& NN$$ Q9>>> L! L L/8 L L L$2>>> xxzzQ MU M M M dAEAIIqq1 5 55r8)r0 rsidesorterctjtjdkpjddjddkfdtjdupjjkfdtj|dkp| d|r tjn tj}t tjr!tj|tj Stj d|j S) Nr1rc\dtjdtjS)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 rr)rsorted_sequencesr6rfz#meta_searchsorted..s> 5378M3N3N 5 5#'tz"2"2 5 5r8cddtjdtjngS)Nz[torch.searchsorted(): boundary and sorter must have the same size, but got boundary tensor z and got sorter tensor rb )rc r_ sr6rfz#meta_searchsorted..sK B#'(=#>#> B B%+%7tFL!!!R B Br8rcdS)Nzetorch.searchsorted(): side and right can't be set to opposites, got side of left while right was Truer2r2r8r6rfz#meta_searchsorted..s $r8r2 r2r) rWrhrrrFrrqr rrrr)rc rr0 rr^ r_ r_s`` ` r6meta_searchsortedrf sF L O !""a' 9  " %CRC 8         L$?//6<?      L #e) $ $ % 5EKK%+E$ %%K U-D    {2U?3IJJJJr8c~tjtjtjtjfvfddS)NcdS)Nz/Unsupported input type encountered for isin(): r2rsr6rfz3_check_for_unsupported_isin_dtype..sI%IIr8)rWrhr complex128 complex64rs`r6!_check_for_unsupported_isin_dtyperk sC L ej%"2EODDIIIIr8cZ|||df}|S)Nrr)rr num_weightsrrr s r6meta_embedding_dense_backwardrn s0''k6F6Fr6J6J(KLLK r8c || r#t||||||||| | St||||||||| | SrB)r+_embedding_bag_sparse_backward!meta_embedding_bag_dense_backward) rJrrsrrmaximum_indicesrm rrYrrrs r6meta_embedding_bag_backwardrs st 22             1             r8c 4tjjtjtjtjtjfvfd|tkrtj|du| df} | S)NcdjS)Nz$Unsupported input type encountered: r)rJsr6rfz3meta_embedding_bag_dense_backward..$sCtzCCr8r1) rWrhr_r]r^r\float64rrr) rJrrrrr rm rrYrrindex_grad_weights ` r6rq rq s L u}enemU]SSCCCC x _D0111 TYYq\\'BCC r8c|d}tj|tkdtj|dktj|dk|d}tj|dktj|d|k||f} | S)Nr1cdS)NzHembedding_bag_backward: per_sample_weights only supported for mode='sum'r2r2r8r6rfz@meta_embedding_bag_per_sample_weights_backward..9rr8rr)rrWrhrrr) rJr>rrsrrYrembedding_featuresrjrs r6.meta_embedding_bag_per_sample_weights_backwardr{ ,s1 L ZZ Lq!!! L!#$$$,,q//K L"### LQ#55666 ^^[N + +F Mr8) assume_uniqueinvertctjt|tpt|tdt|tstj||j}t|tstj||j}t |jt |jtj|tj S)NcdS)Nz.Irr8rr) rWrhrqr r5rrk r_rr)elements test_elementsr| r} s r6 meta_isinr Ds L8V$$I =&(I(INN h ' 'G<1EFFF mV , ,L ]8?KKK %hn555%m&9:::  HEJ 7 7 77r8r ctj|dkdt|tj\}}tj||S)NrcdS)Nz,polygamma(n, x) does not support negative n.r2r2r8r6rfz meta_polygamma..Ys!Or8rr)rWrhrrrr)r rrTrOs r6meta_polygammar VsW LaOOPPP( ;HOA|  D 5 5 55r8c td)Nz.Tensor.item() cannot be called on meta tensors)r+rs r6meta_local_scalar_denser as G H HHr8c*tj|SrBrrs r6silur frvr8cjt|tj\}}tj||Sr)rrrrWr)rrTrOs r6sigmoidr ls=) ;HOA|  D 5 5 55r8c|dk}|dk}|r|r@|d|d|dg}nVtj|d|dkd|d|dg}n|rktj|d|dkd|d|dg}n~tj|d|dkd|d|d|dg}|p|j}tjjrUd|jz}|d|zdz |z|z}||kr|d|z|dg} n|dg} tj|| ||j } ntj |||j } | S) Nrrr1cdSNz matrix batch sizes have to matchr2r2r8r6rfz2_create_grouped_mm_output_tensor..6Xr8rcdSr r2r2r8r6rfz2_create_grouped_mm_output_tensor..r r8cdS)Nzbatched dimension has to matchr2r2r8r6rfz2_create_grouped_mm_output_tensor..s6Vr8r)r) rrrWrhr_rKritemsizer$rr) rKrMoffsrB mat1_is_2d mat2_is_2dr alignment size_paddedrrs r6 _create_grouped_mm_output_tensorr vs1qJqJC  5 ! diillDIIaLLAHH L !  ! ,.X.X    ! diimm4HH  C L !  ! ,.X.X    ! diill3HH L !  ! ,.V.V    ! diillDIIbMMBH'TZI } I),, |i/!3 AIM  # #"1+ 3[!DJJ%q)J! j $+   k()DKHHH Jr8mat_amat_br c dtjdudukdduodu} | rWtjjr tjn tj} tjj| ko j| kfdnAtjjtjkojtjkfdtjdvodvfddk} dk} | r| s@tj d dkd | rHd } d }tj| fd tj|fd d}|d|dtjjtj krjtj kp)jtj kojtj kfdjtj kojtj kdfd }| r| r j dnd}|dd||dd|tj|dud| s| rwtjdufdYtjdkfdtjjtj kfdntjdudtj|dudtj|dup|tjkdt|S)NcdS)Nz,Either both scale factors are given, or noner2r2r8r6rfz)_meta_grouped_mm_common..s>r8c*djdjdS)Nz5Expected inputs of E4M3 FP8 type but got mat_a.dtype= and mat_b.dtype=rrr r sr6rfz)_meta_grouped_mm_common..s!xEKxxjojuxxxr8c*djdjdS)Nz1Expected inputs of BF16 type but got mat_a.dtype=r rrr sr6rfz)_meta_grouped_mm_common..s!t ttfkfqtttr8)rr.c\ddS)Nz3Multiplicands must be 2D or 3D but got mat_a.dim()=z and mat_b.dim()=rr sr6rfz)_meta_grouped_mm_common..s-qeiikkqqdidmdmdodoqqr8rrrcdS)Nz3contraction dimension of mat_a and mat_b must matchr2r2r8r6rfz)_meta_grouped_mm_common..rr8cZ|}|ddko |ddkSNrr1rrmat mat_strides r6rz-_meta_grouped_mm_common..is_row_majors,Jb>A%=*R.A*= =r8cZ|}|ddko |ddkSr rr s r6rz-_meta_grouped_mm_common..is_col_majors,Jb>Q&=:b>A+= =r8cBdddS)NzNExpected mat_a tensor to be row major in the last two dimensions, got strides rr)r sr6rfz)_meta_grouped_mm_common..s)zejeqeqesestvtwtwexzzr8cBdddS)NzQExpected mat_b tensor to be column major in the last two dimensions, got strides rr)r sr6rfz)_meta_grouped_mm_common..s)}hmhththvhvwywzwzh{}}r8cDdz dz}dz dkrQtdjdz kr)t j|zdkfddSdkrTdz tdjkr,t jdz |zdkfddSt jdfddS)Nr1r)rc(ddddS)Nr stride along % dim to be multiple of 16 bytes, got rr2end_dimmat_namer sr6rfzF_meta_grouped_mm_common..check_valid_strides..sBAHAAGAAjtu|j}AAAr8c4dddz ddz dS)Nrr r1r rr2r sr6rfzF_meta_grouped_mm_common..check_valid_strides..s\IHIIGaKIInxzADEzEoFIIIr8Fc ddjdS)NzInvalid strides/sizes, got z for strides and z for sizes.r5r sr6rfzF_meta_grouped_mm_common..check_valid_strides..sijiiSVS\iiir8)r element_sizerrrrWrh)r r r r r s`` @@r6check_valid_stridesz4_meta_grouped_mm_common..check_valid_stridess''))a-#**,,, ZZ\\ gk "a ' 'Jw,?3 sy1%D D - -  L7#i/14AAAAAA      A % %*Wq[*AS sy!F F + +  L7Q;')3q8IIIIII      Liiiii     r8r r c*djdjdS)NzhFor FP8 scales must both be float32, or for MXFP8 both scales must be float8_e8m0fnu. Got scale_a.dtype=z and scale_b.dtype=rrrsr6rfz)_meta_grouped_mm_common..sEr@G@Mrrbiborrrr8r1c dkrtjfd rDtjkfddStjdkfdtjjdjzkfddStjddkfd tjjdjdkfd  rtjjjdz kfd j\}}d }t||z d  t|d tjjdkojd zk fddStjdkfdtjjdjdzkfddS)NrcddS)Nrz to be contiguous.r2 scale_namesr6rfz>_meta_grouped_mm_common..check_scale..sF FFFr8c.ddjdjS)NzKFor MXFP8, scale must have same number of dimensions as target tensor, but  has mat.ndim= and scale.ndim=rr r~r sr6rfz>_meta_grouped_mm_common..check_scale..sN!onx!o!oILIQ!o!ochcm!o!or8r1c:dddS)Nrz to be 1D tensor, but got D tensor.rr~r sr6rfz>_meta_grouped_mm_common..check_scale..( hJ h hRWR[R[R]R] h h hr8rcNddjzdjddS)Nrz to have rrz elements.r5)r r~scale_multiplierr scaled_dimsr6rfz>_meta_grouped_mm_common..check_scale..!s^!UJ!U!U:AVYiAi!U!Uz{FGH{I!U!U!Ur8rcddS)Nrz( to be contiguous in the last dimension.r2r sr6rfz>_meta_grouped_mm_common..check_scale..&s\ \\\r8cHddjddjddS)Nrz batch dimension to be r, got rr5r sr6rfz>_meta_grouped_mm_common..check_scale..*s4p pp39UV<pp_d_jkl_mpppr8c.ddjdjS)Nz0For MXFP8, 3d tensor should have 2d scales, but r r rr sr6rfz>_meta_grouped_mm_common..check_scale..1sF!TS]!T!Tmpmu!T!THMHR!T!Tr8r*rrHc:djddzdjS)NzFor MXFP8, expected mat.shape=z to have scale shape of (,z ), but got r5)G blocked_K blocked_Nr r~sr6rfz>_meta_grouped_mm_common..check_scale..;sY!Y!Y!Yef!Y!Yiru~i~!Y!YLQLW!Y!Yr8c:dddS)Nrz to be 2D tensor, but got r rr sr6rfz>_meta_grouped_mm_common..check_scale..@r r8cNddjdzdjddS)Nrz non-batch dimension to be r1r rr5)r r~r r sr6rfz>_meta_grouped_mm_common..check_scale..Ds]!FJ!F!FSVS\]^ak]kSl!F!FtytABuC!F!F!Fr8)rrWrhr`rrrr=) r r~r r r rorr|r r r is_mxfp8s ````` @@@r6 check_scalez,_meta_grouped_mm_common..check_scale sRwwyyA~~ ''))FFFF L swwyy0oooooo L q(hhhhhL A#)J*?BR*RRUUUUUUUU  LL$$)\\\\ KNcil2pppppp L chl2TTTTTT "iGAq!!#J (Z ; ;I (C 0 0IL A!+W A)iBW0WYYYYYYYY L q(hhhhhL A#)A N*CCFFFFFFFr8rrrcdS)Nz:Scale result tensor provided, but it is not supported yet.r2r2r8r6rfz)_meta_grouped_mm_common..Ornr8c^dddS)Nz/Offsets tensor not provided, but is needed for zD/zD multiplicand layouts.rr sr6rfz)_meta_grouped_mm_common..Us0yeiikkyyUZU^U^U`U`yyyr8c4ddS)Nz.Offsets tensor must be 1D, but got offs.dim()=rrr sr6rfz)_meta_grouped_mm_common..ZsVVVVr8cdjdS)Nz7Offsets tensor must be integer (int32) tensor, but got rrr sr6rfz)_meta_grouped_mm_common..^s_RVR\___r8cdS)NzJOffsets tensor provided, but is not needed for 3D/3D multiplicand layouts.r2r2r8r6rfz)_meta_grouped_mm_common..cs`r8cdS)Nz2Bias tensor provided, but it is not supported yet.r2r2r8r6rfz)_meta_grouped_mm_common..hsDr8cdS)Nz4If output dtype provided, it must be torch.bfloat16.r2r2r8r6rfz)_meta_grouped_mm_common..msFr8r)rWrhrKrLrr_r_r^rrr\rrrFr )r r rrr r@rrBrscaled fp8_dtype mat_a_is_2d mat_b_is_2drrr r r r s````` @r6_meta_grouped_mm_commonr s L Dgo.>>D 8WD%8F   -2]->WE))EDW  K9 $ A )A x x x x x     K5> ) KekU^.K t t t t t    L v7%))++"7qqqqq ))++"K))++"K  k  JJrNNejjnn , I I     > > > > > >  L   z z z z     L   } } } }   0''''''w2 ]em + N 0N  !55:MU%99 r r r r r     MU1 1 6 !55  : : : : : : z"-+-+-DJqMMST   Iwq2BCCC Iwq2BCCC D P P    k    y y y y y      L aVVVV    L ek)____     DL ` `    L  DD  LT8Y%.8FF ,E5$ J JJr8c 0t||dd||d|S)N)rrr r@rrB)r )r r r r@rBs r6meta_grouped_mmr ss3 #        r8c N|p tj}t||||||||| S)N)rrr r@rrBr)rWr^r ) r r rrr r@rrBrs r6meta_scaled_grouped_mmr sB+U^I "    !%    r8r; half_to_floatc|r!|jtjtjfvsJt j|tjj\}}|s|n|}tj||tj }|S)Nrr2 ) r_rWrYr^rQrrrRrr)r;rr computation_dtyperOrs r6softmaxr s7w5:u~66666&+&> uDL'''#|(5K<<:KL  1L@W X X XC Jr8c   tjtdzdkfd|jt tdz} |z tj |k fdt dDr|}t D]d z dz z  dkr2|  |j z} dzdkr.|d|j dzz}|Std }t |D]otdzdzz  z z dzz}tj|dk  fd| |ptj ||j |j |j t|S) Nrrc(dtS)Nz1Length of pad must be even but instead it equals rrWsr6rfz'_constant_pad_nd_meta..sNCHHNNr8c0dtddS)Nz`Length of pad should be no more than twice the number of dimensions of the input. Pad length is z while the input has z dimensions.r)l_inprXsr6rfz'_constant_pad_nd_meta..s/25c(( r8c3VK|]$}t|tjo|dkV%dSrx)rqrQIntWithoutSymInt)rNrFs r6rtz(_constant_pad_nd_meta..s8 I I:a/ 0 0 ;Q!V I I I I I Ir8r1c Xdzdddzdzd S)NzThe input size z, plus negative padding rr1zG resulted in a negative output size, which is invalid. Check dimension z of your input.r2)rrl_diffrXpad_idxsr6rfz'_constant_pad_nd_meta..sfMk&1*&=MM7|MM"%gk"2MM17!MMMr8)r_rrr)rWrhrrrrnarrowrrrrr_rrr) rrXr*l_padc_input new_shapenew_dimrrr r r s ` @@@@@r6_constant_pad_nd_metar s L C1 NNNN +K   E HHME U]F L       I IS I I III vu%% T TA519q=)G7|a!..G }gmA&6W&E7Q;!##!..Aw}Q/?#gPQkBR/RSS}}[&)**I 5\\ " "c((q1uk*fqj)CL83w{;KK qL M M M M M M M M   !!!! ;k|)+E22    r8rrrc|dks Jd|j}|j}|jdkr |df}n)|jdkr|d|df}n g||dR}|j}|||S)Nrz'weight' must be 2-Drr1r)rrrr_r) r>rrrr weight_shape indices_shaperrBs r6 embeddingr s ::<<1   4   ._f s$ =J    r8rIr"funcr s r6_create_unary_float_meta_funcr  94]]  ] Ir8c  |js|js|jrtd|dkr/||j|dfS|d} |d}|d}|| ||}| r3| r|| ||}n.|| |||}n|d}||fS)NzP_native_multi_head_attention fake implementation does not support nested tensorsrr1) is_nestedr}rrrr)rBrr* embed_dimr qkv_weightqkv_bias proj_weight proj_biasr need_weightsaverage_attn_weights mask_typerT output_dimr attn_weightss r6 native_multi_head_attention_faker s   #- 5? ! ^    {{}} ,,eooa.@.@AA 1 A 1 A !!!$$J __Q: . .F *  >!??1a33LL!??1h1==LLq)) L !!r8c^t|td}|S)Nc:t||tjSrr r:s r6r z*_create_binary_float_meta_func.._fR s& q!@!M    r8r r s r6_create_binary_float_meta_funcr Q r r8ctfd}jd}||_ttt||}|S)NcT|g|Ri|}t|j|j|SrBr))rrSrrrDs r6_fnz#_register_inplace_meta.._fnv s;b'''''' CI666 r8rT)rr{rIgetattrr+)rDr inplace_names` r6_register_inplace_metar u sk 2YYY k$$$LCL 4-l33 4 4S 9 9C Jr8cBtjjjkfdg}ttrGjdkr'tjjjkfd|t|dtj iS)Nc(djdjS)Nrz for `end`, but got dtype r)r|r}sr6rfzlerp.. sT%+TTTTr8rc(djdjS)Nrz for `weight`, but got dtype r)r}r>sr6rfzlerp.. sb%+bbTZT`bbr8rJ) rWrhr_rqrrrrUrrR)r}r|r>rSs``` r6lerpr" s L sy TTTTT 3. s  r8r)rWrhrQr:r_rUrrRr$ s r6addcdivr* sr L  "7= 1 1 6&w}55    w0O0W   r8ci}dD]#}t|}|D]}||vr ||||<$|D]\}}t|tjjr&t|t sJ|tjj j |tj | dr"|tdvrt|d|jr| dvrd| vrt||d| vrt"||Hd| vrt$||{d | vrt&||t(||dS) N)r post_autograd pre_autogradCompositeImplicitAutogradrz 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::rot90aten::_to_copyaten::empty_stridedaten::constant_pad_ndaten::as_strided_scatterzmkldnn::zmkl::zonednn::z quantized::)ritemsrqrW_opsHigherOrderOperatorrpy_impl_C DispatchKeyr-%_dispatch_has_kernel_for_dispatch_keyrr+is_view2_meta_lib_dont_use_me_use_register_meta_for_mkldnnimpl/_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_tablerzregistryopo op_overloadrDs r6 activate_metarH sb:99-d3 9 9C---+3C=#C( 9/44666N6N R k5:#A B B  +z222226 EH0566r::: 8 9 9      ;  + N8@@@"";;;A   N         [--////BGG UWXXXXK,,....?DD[RTUUUU{//1111BGG UWXXXX+"2"2"4"444EJJ8<<["MMMMm6N6Nr8rr rB)NNNFrr1rr)Tr)r)rT)FF)TT)r[)FTN)TFF)TF)r)rmN)r>r)r2rrF)r2rFTN)FrFNFr)NF)rF)rrFN)NNNNN)rNNr1)NNF)rAFFN)NrAFFN)rAFNN)NrAFNN)rAFN)FN)FNNNN)NNNF)NNNNF)NrFNN)NNNN)rTT)NNrN)r5 rr)r)rFF)rA)NTTN(rcollections.abcrrenumr functoolsrtypingrtyping_extensionsrrWtorch._prims_commonrrQr r r torch._decompr r rr torch._opsr torch._primsrrrrrrrrrrrrrtorch._prims_common.wrappersrrr r!r"rr#r$torch.fx.experimentalr%rctorch.nn.functionalr&r' torch.utilsr(rFr)r*opsr+libraryLibraryrC rrrrr7r=rIrUrarjlinspacelogspacerrtaker.rrrrcummaxcumminrrrrrrr_fft_c2crrr_fft_r2cr randperm generator_outrrrrandintrrlow_outrrandr!_fft_c2rr&rr2r9 unsqueeze_r=_sparse_semi_structured_linearrr_rJ_sparse_semi_structured_mmrP_sparse_semi_structured_addmmrT_cslt_sparse_mmrd index_reducerk index_reduce_rm index_selectrqsegment_reducer~r unary_outrrrrrrrrr _assert_asyncrmsgr_printr_make_dep_tokenrr_functional_sym_constrain_rangerr(_functional_sym_constrain_range_for_sizer_functional_assert_asyncrrrrrrrr _linalg_eighrr_linalg_eigvalslinalg_eigvalsr linalg_eigrrrrr r rr& linalg_inv_exr)linalg_ldl_factor_exrgr/linalg_ldl_solver> linalg_lurElinalg_lu_factor_exrIlinalg_lu_solverR lu_unpackrXra linalg_qrrhrlri _linalg_svdrur;rrrlinalg_solve_triangularrrr _linalg_detrrrrreflection_pad1drreplication_pad1drrreflection_pad1d_backwardrreplication_pad1d_backwardrrreflection_pad2drreplication_pad2dr_weight_norm_interface_backwardr$reflection_pad2d_backwardrreplication_pad2d_backwardr+r9reflection_pad3dr;replication_pad3dr>reflection_pad3d_backwardreplication_pad3d_backwardrE_pdist_forwardr[rI_pdist_backwardrObaddbmmri bernoullirl bernoulli_rprFrspoissonru_fused_moving_avg_obs_fq_helperrmmrrrrrmiopen_batch_normr convolutionrr: _has_mkldnnr> r_convolution_pointwiser_linear_pointwiserhas_mklr@ r _mkl_linearrrA rqconv2d_pointwiseqconv_pointwiser5rbinary binary_tensorrqlinear_pointwiserrlinear_dynamic_fp16linear_relu_dynamic_fp16rrB r max_pool2drint4mm_packed_weight_cpurr avg_pool2drErJavg_pool2d_backwardrR avg_pool3droavg_pool3d_backwardrz_adaptive_avg_pool2dr_adaptive_avg_pool3dr_adaptive_avg_pool2d_backwardr_adaptive_avg_pool3d_backwardrradaptive_max_pool2drrradaptive_max_pool3drrrrepeat_interleaverrrrrrr _unsafe_indexrconvolution_backwardraddbmmr randint_liker _fused_adam_ _fused_adamw_r _fused_adamr_int_mmr"_convert_weight_to_int4packr-#_convert_weight_to_int4pack_for_cpur2_weight_int4pack_mmr;_weight_int4pack_mm_for_cpurArJrLrt_dyn_quant_pack_4bit_weightr_dyn_quant_matmul_4bitr_weight_int8pack_mmr_cdist_forwardr_cdist_backwardr_embedding_bagr_embedding_bag_forward_onlyrrnansumrmedian nanmedianr dim_valuesrYrr logical_not_rrepeatrzero_rmul_Scalardiv_ logical_and_ logical_or_ logical_xor_radd_sub_rrsubrrounddecimalsrr __rshift__r __lshift__r zerorrrfillrrelu_r _add_relurrrelu_with_noiserrrelu_with_noise_functionalr!rrelu_with_noise_r# index_put_unsafe_index_putr' masked_fill_r+ _masked_scaler.masked_scatter_r2masked_scatterr4masked_scatter_backwardr6 index_put_r8rFbmmrIrKrNrSr9r:rdrzrvr max_pool2d_with_indices_backwardrmax_pool2d_with_indicesrfractional_max_pool2drmax_pool3d_with_indicesr max_pool3d_with_indices_backwardrrrrgrid_sampler_2d_backwardrrrrronesrzerosrselect_scatterr slice_scatterrrrrgatherr rr#r%r/r2 scatter_addr7 scatter_add_r:r1r/r*rg value_reducer>scatter_r@#_scaled_dot_product_flash_attentionrXrc#_scaled_dot_product_cudnn_attentionrm0_scaled_dot_product_fused_attention_overrideablerp,_scaled_dot_product_flash_attention_backwardrz+_scaled_dot_product_flash_attention_for_cpur}4_scaled_dot_product_flash_attention_for_cpu_backwardr*_scaled_dot_product_attention_math_for_mpsr'_scaled_dot_product_efficient_attentionr0_scaled_dot_product_efficient_attention_backwardr,_scaled_dot_product_cudnn_attention_backwardr_flash_attention_forwardr_flash_attention_backwardr_efficient_attention_forwardr_efficient_attention_backwardSymIntrr  _scaled_mmr r^ _scaled_mm_v2rascatter_reducetwotwo_outrescatter_reduce_rg multinomialrlrpr}r_upsample_nearest_exact1dr_upsample_nearest_exact2dr"_upsample_nearest_exact2d_backwardr_upsample_nearest_exact3drr values_stablerr_thnn_fused_lstm_cellrrrrrrargminrrtopkr_segment_reduce_backwardrkthvaluerrr r r r pixel_shuffler r/ bucketize Tensor_outr4 histcr@ _upsample_bilinear2d_aa_upsample_bicubic2d_aarD _upsample_bilinear2d_aa_backwardrH rQ rV rr] searchsortedrf rk embedding_dense_backwardrn _embedding_bag_backwardrs _embedding_bag_dense_backwardrq *_embedding_bag_per_sample_weights_backwardr{ isinr polygammar _local_scalar_denser r r r r _grouped_mmr _scaled_grouped_mmr _softmaxr constant_pad_ndr r _jagged_to_padded_dense_forwardr r _native_multi_head_attentionr r special_airy_aispecial_bessel_y0special_bessel_y1special_modified_bessel_i0special_modified_bessel_i1special_modified_bessel_k0special_modified_bessel_k1!special_scaled_modified_bessel_k0!special_scaled_modified_bessel_k1special_chebyshev_polynomial_tspecial_chebyshev_polynomial_uspecial_chebyshev_polynomial_vspecial_chebyshev_polynomial_w&special_shifted_chebyshev_polynomial_t&special_shifted_chebyshev_polynomial_u&special_shifted_chebyshev_polynomial_v&special_shifted_chebyshev_polynomial_wspecial_hermite_polynomial_hspecial_hermite_polynomial_hespecial_laguerre_polynomial_lspecial_legendre_polynomial_pr r" r' r* lerp_addcmul_addcdiv_torch._refs.nn.functionaltorch._refs.specialrH r2r8r6rK s^ ........'''''' ######++++++++++ "!!!!!UUUUUUUU                          <;;;;;;;77777788888888)))))) WT]]Yt__ y~*/-*?*?PV*W*W' %a)X""" 8BF#3"4hr2v6F"FG    3(*** t}-..     =555/.5p !49=122  ' '32 '!)4+<+@ABB %'%%%%CB%$t%&& II'&I  [$+/4;+> P Xy!!"!!)4+<+@ABB IICBI333lV$s)4 %t}'89:: KK;:K$s) %t}'89:: 8 8 ;:8 vt}*++"&3333,+3t}$%% *      &%  $dl&6788  *  98&  $,"6788  *  98& !49=122 %)$tPT32  %t}'89:: $Dv$DDI$Dc$DC$D$D$D;:$DNtz!""#"0 ( ( (t&''('t233 "&$(     4-  4Z  {T! 43Bt.// %)   {T! 0/0t122 $(    {T!32Dt#$$$("+<+<,+< \+< 4-+< D= +< {T! +<  +< +<+<+<+<+<%$+<\t ()) I I I  I  I  I L I  I I  I I I*) It!)**          L        +*  t ())''*)'t"*++ "!!  W W W  W  Wd] W d] W d] W  W  W  W W W,+ WF $("4566 76tx| $("4566 76tx|tz!""66#"6tz~((( t!)**  +* t!%&&  '& t{"##  $# t#+,,    ))))-,)t'/00,,,10,t3;<<=< t0899666:96&t<DEEFE t,01121 FC  F #    N(,     !%    $VSC  F $     "            C!)4+<+HIJJ ]N++sT,+KJ"$,d.A.EFGG BB6BBBHGB !! ]N++ 6   ,+"!  QQFQQQQt*++ ))F)4)F))),+)t"##  J J JF J4 JF J J J$# Jt} ))6)$)6))))t$%% ))6)$)6)))&%) t&.//&T0/" $,d.M.QR .f.6.f....dt!)**&+*)143L3PQRR T8V$$     666 !" %$SR&%-t/D/HIJJ   '''' ''  '  '''KJ'T&(:;<< S#s/3ffff>T8U=<4(0$2J2NOPP T8V$$      666 !" %$QPD$,d.B.FGHH   4444 44  4  4 444IH4nt~ S#s $$$ $$ $  666 !" $$$$PtTz!2*&(:;<< S#fCffn8M=<4$,d.B.GHII V[$11 'v '%(F"G ' ' '21JI '$t'((  "" """ $J """)("J. . . 49d3i  ....". . . *. 66> ....(fVt$%%   777 7 7  7  7 TM 7  7 TM7 4-7 6666 )*777&%7t,4d6R6VWXX          $ YX2t$%% S#4((( +(+( +( +( +( +(  +(  66> +(+(+()(&%+(^t'((  )( tz   WW W W W  W  W  WWWWt>#;#;#;Lt$%% ==&%=t%&& >>'&>(((<t-.. \SS/.St.// \TT0/T2E2E2Ejt$%% ==&%=t%&& >>'&>(0  &. &1 '/ '2  \&&&:<G<G<G~t$%% ==&%=t%&& >>'&> &. &1 '/ '2  \$($($(Nt"##    f    v    $#  t#$$ PvPVPPfPQWPPP%$P $dl&6788 /0''''98':&(:;<< &*IIII=<I t$%%&%t~  III! I  $dl&6788 """98"t3;<<***=<*.tw 33U[4/3333,BBB    *.2rr,r Lr IOr#Y_ r 3i#o r  r rIOd*rrrrjQQQt%-.."$,"$ L"$ , "$,% "$ $ "$  "$!&"$"$"$"$/."$Jt'((, L , I  #Y  3i I )(> 8G@9>9N9N&&::6]59#:BCCDC,]59#5=>>SS?>S  x :?-:O:O 66; ; 7 uy}0 1 1   2 1  :?9N9N&&::6]59#5=>>]59#3;<<]59#3:;;00<;=<?>0d]59#5<==]59#5CDDED>=6]59#5=>>]59#5<==>=?>8]59#5<==]59#5CDD!!ED>=!F]59#7?@@]59#<DEE  FEA@ =BM>@@?>@F)1233HH43H: #T[_566 ./q''''76'0!()**''+*'!)4+=+EFGG !    HG 2 ()** !!!!!+*!H ~ BBB*012232&89::  ;: ()** @ @+* @0122 < <32 <>?@@ < <A@ <"3"3"3""""kkk\0122D!'$DDD32D4+,--;;.-;(()** < <+* <t"*++ & &,+ &Ft#$$ GG%$G*t"*++  `5`5`5,+`5Ft/78855985     #T[_566 =$====76= #T^%;<==))>=)    !      Xy!!"!  t ())*)t{"##''$#'&tz!""#"     !   !          ****Z       "DJ$7899:9    "&(>?@@A@&(>?@@A@ty !!&&"!& !4:#456676  $)"2344""54"tz!""#"t~$%% F&% %&'' RV"""('" 0122RV;;;32; &'((KO)( &(>(FGHH"""IH"t '(()( t!)**+*t#$$  %$ t"## 66$#6 t+,,!!-,!t&''('((((Vtx  55! 5tx~JJJ   6   ;;;h#-YY Y Y Y Y Y Y Y Y Y Y YYYY Y !Y" #Y$ %Y&'Y( )Y*+Y,-YYYYx;4;4;4|38 383838 38 38 38 38 38 38 38 38 38 3838 38 !38"#38$ %38&'38383838lI2I2I2Xt4<==((>=(Vt+344    ###54#Lt)122QQ32Qht+,, UI    ddd-,dNt455 \bb65bJ% V% 6% % % % Pt v3$t,455##65#$t#$$ 88%$8"t,-- \;''!!('.-!, !"##..$#.t&''    )))(')X !49=122    32( "DJN344    54(t"*++..,+.t!)**...+*.   C d    ///    t{"## ' ' '$# '6      - - - - b4444t'((&&)(& t !!"!     !  &&&&     "  89:: #<< < < < <  <  < 4<<<<;:<~ S#X&89::#(( ( ( (} (  (  ((( 4<(((;:(VEFGG $#'' ' ' '} '  '  '' 4<'''HG'T 9(""" " "  "  "  """ " """"" 4<"""  ". 8#       }  4<  : A $!"!"!" !" !"  !"  !"  !"!"!"}!" 4<!"!"!"  !"H?@AA $"&3&3& 3& 3& 3&} 3&  3&  3&4-3& 4<3& 66>3&3&3&BA3&l<=>>)))) )) )) ))} ))  )))) 4<))))))?>))X ="4-4-4- 4- 4-  4- } 4-  4-4-4-4-4-$Z4-4- 4<4-4-4-  4-n 9*!""" " "  "  "  """""" " """ 4?@@      A@    ,,,* $d&D&LM     $d&D&LM. (0 /7"! # ,-u|+,dl  dl   : $d&D&LM         &%%%Pt)122    32 t&''4/4/('4/nt$,--$%$%.-$%N        #T[%89::444;:4t!)**+*ty !! Q Q Q"! Qt,-- LP   .-   %t}';<== Xy!! K K K"!>= K #("9 Q Q Qt7?@@ * *A@ *t#+,, 0 0-, 0t!)**+*>t-566FF76FD%t~'@ABB 27uCB | EEEE4 !)4+F+NO  &5=>??   @?8t>FGGHG$')<=>> """?>"uy~())*)8uy~ !!66"!6$t !!    -K-K-K-K"!-K`t,--.-t+,,' ' ' -,' Tt12232,t>??@?.ty 8=e 8 8 8 8 8 t~ 6c66F6666t'((I&III)(Ity "v"&""""t| 6&6V6666'''^(,$( PKPK PK PK\D PK\D PK 4- PK 4- PK,%PK{T!PKPKPKPKPKft   $(    4- 4-  {T!   ! &'()) !% $(,$(  < <\\  ,   ,  ,%{T!*)6t}  v C      t#$$ 333%$3lt~ $ 88 8 88 8  8  8888,t3;<<  * *  * &\ *c * * * *=< * t0899 -"-"-":9-"`d2333d4555d4555d=>>>d=>>>d=>>>d=>>>dDEEEdDEEEtBCCCtBCCCtBCCCtBCCCtJKKKtJKKKtJKKKtJKKKt@AAAtABBBtABBBtABBB   ty $t| ./ t| ./, ty)) ! !$, / / ! !$, / /  BNBNBNJ r8