i -{ddlZddlmZmZddlmZddlmZddlm 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 jxjzZ=e j|jdddZ@eAd\ZBZCZDdZEdZFdeee;e:fgee;e:fffdZGde!fdZHdZIdZJeGe=je=jge-ddde jddfd ZNeGe=jje=jjge-d!ZReGe=jje=jjge-d"d#d$ZSeGe=je-d%ZTeGe=jje=jje=jje=jjge-d&d'd(ZWeGe=jje=jjge-d)ZXd*ZYd2d+ed,eZe[d-e\fd.Z]eGe=jje=jjge-d/Z_dZ`d,eZe[fd0ZaeGe=jje=jjge-d1ZceGe=jjdd2d3ZfeGe=jje jdddd4d5ZheGe=jje=jjge-e jdddd4d6ZjeGe=jje=jjge-e jdddd4d7ZmeGe=jje=jjge-ddddd4d8ZoeGe=jje=jjge-d+ed,eZe[d9e[d:e[fd;ZqeGe=jjd2d<Zsd=ZteGe=jjd>ZveGe=j d3d?ed@edAedBedzdCexdzdDe jdzf dEZzeGe=j d4dFedGedHedDe jdzfdIZ|eGe=jdJdJddKd?edFedGedHedDe jdzf dLZ~eGe=j d5dMe j$dNe j$dBedzdOedzdDe jdzdPe\dQe[dRe[dSe[fdTZeGe=jjdUdVd+ed,e[dWedXe j$dYexdZe\defd[ZeGe=jjdUdVd+ed,e[dWedXe j$dYexdZe\defd\Ze-eGe=j jd]ZeGe=jjddddddd^d_edYexd`edzd'edzdaedzdbe[dce\defddZeGe=jje=jjge-deZeGe=jjd2dfZeGe=jje=jjge-dgZeGe=jjd2dhZeGe=j"jdiZeGe=j"jdjZeGe=j(jdkZeGe=j(j,dlZeGe=j0jdmZeGe=j4jddddddndoZeGe=j8jd6dpZeGe=j:jd3dqZeGe=j>jd6drZeGe=j@jdsZeGe=jDj,dtZd+eduexfdvZd+edwedxexfdyZ d7dzeduexd{e\fd|Zd8dweduexd}exfd~Zdwedede\duexfdZ d9dexded?edexfdZdexfdZeGe=jVje=jVjXge-ddd:dwedexde\fdZeGe=j\je=j^jge-d?edefdZeGe=jbge-ddd?efdZdedefdZeGe=jhe-d+edwede\defdZeGe=jje-d2d+edwede\defdZeGe=jle-d2d+ede\defdZeGe=jne-d2d+ede\defdZeGe=jpjd;dwede\de\fdZeGe=jrje=jrjge-d?ededefdZeGe=jtjd2dwede\fdZeGe=jxje=jxjge-ddddddd+ede\de\deeeeffdZeGe=j~je=j~jge-dddededede\def dZeGe=jje=jjge-ddddUddwede\deeeeffdZeGe=jje=jjge-ddddUdddwede\de\deeeeffdZeGe=jje=jjge-dUdddededede\de\def dZeGe=je-ddd ddwede\de\dexdzfdZdededeeZe[eZe[ffdZdededxexdzdeeeffdZd?edede\fd„ZeGe=jdUddddddÜdwedede\de\dedzdedzdedzdedzdeeeeeffdĄZeGe=jje=jjgdUdddŜdwedede\de\de\dedzdefdȄZeGe=je-dddU˫ d?d+edwede\de\de\deeeff d̈́ZeGe=jjd΄ZeGe=je- d@d?ededede\de\def dτZdЄZdфZeGe=je-d҄ZeGe=je-dӄZdԄZeGe=je-dիdքZeGe=je-dիdׄZd؄ZeGe=je-dلZeGe=je-dڄZeGe=jjdۄZeGe=jje=jje=jje=jjge-dիd܄Zd݄ZeGe=je-dބZeGe=je-d߄ZeGe=jje=jje=jje=jjge-dիdZeGe=je-dAd+ededefdZeGe=je-ded+edededef dZeGe=jje=jjge-dU˫dJdJddZeGe=jje=jjge-dd2dZeGe=jjdBdZeGe=jjdBdZeGe=j je=j jge-d4dZeGe=jj d;dZ eGe=je-dU˫d4dDe jdzfdZ dZ dCdZ d4de j$d@e j$deZe[e[zdeZe[e[zdeZe[e[zde\de[deZe[e[zdzfdZdZeGe=j jde j$d@e j$dBe j$dzde j$dzde j$dzde\dedefdZeGe=j$jde j$d@e j$dBe j$deZe[deZe[deZe[de\deZe[de[fdZe j(j*re j|jdddZeGe jxj.j0jdZeGe jxj.j4jdZe j(j8rOe j|jdddZeGe jxj<j>dZ e j|jdddZ!eGe jxjDjFjeGe jxjDjHjeGe jxjDjHjJdZ&eGe jxjDjFjNeGe jxjDjFjPdZ)eGe jxjDjTjeGe jxjDjTjJdZ+eGe jxjDjTjNeGe jxjDjTjPd Z,eGe jxjDjZjeGe jxjDj\jd Z/e j|jd ddZ0eGe jxjbjd dDdZ3eGe jxjbjhdZ5dZ6eGe=jnj dEdZ8dZ9eGe=jtjdZ;eGe=jxe- dEdZ=eGe=j|e-dիdZ?eGe=jjdZAeGe=jjdZCeGe=jjdZEeGe=je-dիdZGded}exfdZHeGe=je-dd'dZJeGe=je-dիdZLeGe=je-dd'dZNeGe=je-dիdZPeGe=jj$d4d ZReGe=jje=jjge-d!ZTeGe=jje=jjge-d"d"d#e[fd$ZUeGe jxjzjje jxjzjjge-d%ZVeGe=jj$e=jj$gd&ZYeGe=jjgd'Z[eGe=jje=jjge-dU˫dJdJdd(Z]eGe=jj$gd)Z_eGe=jje=jjgddd*d+ZbeGe=jjgddd*d,ZdeGe=jge-d-ZfeGe=jgd.ZheGe=jgd/ZjeGe=jgd0ZleGe=jgd1ZneGe=jgd2Zod3e[d4e[de[fd5Zpd6ZqeGe=jgdBedzfd7ZseGe=jgd8ZueGe=jgd9ZweGe=jjd:ZyeGe=je-d;Z{eGe=jj dFd<Z}eGe=jjd=Zd7d>ZeGe=jje=jjge-dGdd?d@ZeGe=jje=jjgdAZeGe=jje=jj e=jje=jj e=jje=jjge-d&d'dHdBZeGe=jjdCZeGe=jjdDZeGe=jjdEZeGe=j j"e=j$j"e=j j$e=j$j$e=j&je=j(je=j*jgdFZeGe=j.j"e=j0j"e=j.j$e=j0j$gd dGZeGe=j4j"e=j6j"gd dHZeGe=j:je=j:j<gdIZdJZeGe=jBj$e=jBj"gdKZeGe=jFj$e=jFj"gdLZeGe=jJjdMZeGe=jNj$e=jNj"gdNZeGe=jRj$e=jRj"gdOZeGe=jVjdPZeGe=jZj$e-d defdQZeGe=j^ge- dIdRZeGe=jbg dIdSZeGe=jfg dIdTZeGe=jjje=jljgd2dUZeGe=jpj"dVZeGe=jtjdWZeGe=jxdXZeGe=j|e-dYZeGe=jdZZeGe=jjd2d[ZÐd6d\ZeGe=jjd]ZeGe=jjd^Zǐd_ZȐd`ZɐdaZʐdbZ d2d?edce[dde[dee[dfe[dge[dhe[die[dje[dke[dle[dme[dne[doe[dpe[dqe[dre[dse[dte[due[dexdve\f,dwZ̐dxZd?ededce[dde[dee[dfe[dge[dhe[die[dje[dke[dle[dpe[dqe[dre[dse[dte[due[dexf&dyZΐdzZeGe=jjd{ZeGe=jj dDd|ZeGe=jjd}ZeGe=je-dd' dDd~ZeGe=je-dիdZd?edefdZGddeZd?edede[fdZeGe=jjdZeGe=je-dZeGe=je-dՐddZeGe=jjgdZeGe=jj dJdZeGe=jje=jjge-dddddddZeGe=jje=jjge-dddddddZeGe=jjdZeGe=jjdKdZd7d,e[de[de\fdZdZdZeGe=jjd2dZd2dZd4dZdZd4dZdLdZeGe=jjdZeGe=jdZeGe=jje=jje=jje=jjge-d4dZeGe=jje=jje=jje=jjgd4dZeGe=jg dMdedededede\de\dedzfdZdedee[dffdZeGe=jg dMdededededzde\dede\de\dedzfdZeGe=j g dNdededededzdede\de\dedzfdZeGe=jg d4dedededededededede[de[dede\dedededzfdZ eGe=jg dOdedededede\dedzdedzfdZ eGe=jg d6dededededededede\dedzdedzfdZ eGe=jg dPdededededzdede\dedzdedzdeeeffdZeGe=j g dQdededededzde\de\dedzfdZeGe=j$g dRdedededededzdedededededeZe\de\dedzfdZeGe=j(g d4dededededededededededede[de[dede\dedzf dZeGe=j,g dJdededededzdedzde[de[dede\de\dedzde[dzde[dzdedzdedzfdńZeGe=j0g d3dedededededededede[de[dede\dedededzde[dzde[dzf"dƄZeGe=j4g dSdedededBedzdedzdedzde[dzde[dzdede[de\dedzdedzdedzde[dzfdτZeGe=j8g dLdededededBedzdedzdedzde j:de j:dedededede[de\dedzde[dzde\f$dӄZ dTd+e j$dHe j$de j$de j$dBe j$dzde j$dzdDe jdzde\fd؄ZeGe=j@jg dTd+e j$dHe j$de j$de j$dBe j$dzde j$dzdDe jdzde\fdلZ! dUd+e j$dHe j$deZe j$deZe5deZe j$deZe5dBe j$dzdDe jdzdeZe6dzdeZe6dzde\fdބZ"eGe=jFjg dTd+e j$dHe j$deZe j$deZe5deZe6deZe j$deZe5deZe6dBe j$dzde jdzdeZe[dzde\fdZ$eGe=jJjLe=jJjNge-d7dZ(eGe=jRjLd7dZ*eGe=jVje=jVjge-d2dd2dZ,dZ-dZ.eGe=j^je=j`jgd4dZ/eGe=jbje=jdjgd6dZ1eGe=jfje=jhjg d6dedee[e j:zdee[e j:zdedzdedzf dZ3eGe=jjje=jljgd3dZ5eGe=jnje=jnjpe=jnje=jnjrgdVdZ:dZ;eGe=jxj d6dZ=eGe=j|jdZ>eGe=j~jdZ?dZ@dZAeGe=jje=jjgdGdZDeGe=jjdWdZEeGe=jjdXdZGeGe=je- dYdZIeGe=jje=jjge-d&d'dHdZKe jZMdZNeGe=jjdZOeGe=jjdZPeGe=jjdZReGe=jjdZSeGe=jj$e=jjge-ddddZVeGe=jge-dZdZXeGe=jje=jjg d6dZ[eGe=jjg d6dZ]eGe=jjdZ^eGe=jje=jjge-d3dZ_eGe jxjzjdZ`eGe jxjzjdZaeGe=je-ddddd d Zcd ZdeGe=jd ZfeGe=j d[d ZheGe=j d[dZjeGe=j d[dZleGe=je-ddddZneGe=je-de[d+edefdZpeGe=jd+efdZreGe=je-dU˫d+edefdZseGe=je-d+edefdZtdZu dUdedede j$dzde j$dzdedzdBedzde j$dzdDe jdzde\fdZveGe=je- d3dedededzdBedzdDe jdzdef dZxeGe=jg dUde j$de j$de j$de j$de j$dzdBe j$dzde j$dzdDe jdzde\fdZzeGe=je-ded,e[de\defd Z|eGe=je-d d!Z~eGe=je- d\d@ed'ed"e[d#e\d$e\def d%ZeGe=jj d]d&edaeZed&eZe[d'efd(Zd)ZeGe=jj d^d*Zd+Zee=j ee=jee=jee=jee=jee=jee=jee=jee=jee=jee=j ee=j"ee=j$ee=j&ee=j(ee=j*ee=j,ee=j.ee=j0ee=j2ee=j4d,ZeGe=j8e-d-ZeGe=j:e-dJd.d/ZeGe=j<e-dJd.d0Zee=j8Zee=j:Zee=j<Zddl.Z ddlZ ddlZ d1Zey(_N)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 S/var/www/html/engine/venv/lib/python3.12/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)NcVtfd}tj|S)Nc(tt|yN)r r)opfns r6registerz0register_meta..wrapper..registerDs  B 3r8)r pytree tree_map_)rDrErCs` r6wrapperzregister_meta..wrapperAs)  $ 4 2& r8r2)rCrHs` r6 register_metarI@s Nr8type_promotionctj|d|i\}}|Dcgc]}t||}}t|}t |dt j iScc}w)Ntype_promotion_kindrJ)utilsrrr$rrDEFAULT)rJargs_ result_dtyper;s r6elementwise_metarRMss .. *OA|?C C #A| 4 CD C T "D "  BJJ  DsActjtjtjtjtj tj i}|j||SrB)torch complex32halfcfloatfloatcdoubledoubleget)dtype from_complexs r6toRealValueTyper^asE  ekk u||L   E5 ))r8clttg|tjk(fdy)NcddS)Nzoutput with shape z# doesn't match the broadcast shape r2)broadcasted_shape self_shapesr6z)check_inplace_broadcast..ns$ZL0STeSfgr8)tupler#rT_check)rb args_shaperas` @r6check_inplace_broadcastrgjs0/ HZHI LLZ'gr8Fc < ttjr(tjj dk(dttjr(tjj dk(dt dfDrZt jtj  nFtjt j fdnxstjttjsJtjttfdttsJtjdk\dtjf|d|| S) NrcyNz:linspace only supports 0-dimensional start and end tensorsr2r2r8r6rcz(meta_linspace_logspace..r8cyrjr2r2r8r6rcz(meta_linspace_logspace..rkr8c3<K|]}t|tywrB) isinstancecomplex).0args r6 z)meta_linspace_logspace..s C:c7 # CscddS)Nzlinspace(): inferred dtype z& can't be safely cast to passed dtype r2)default_complex_dtyper\sr6rcz(meta_linspace_logspace..s56K5LLrsxryzr8cdtjdtjdtjdS)Nz4received an invalid combination of arguments - got (, ))type__name__)endstartstepssr6rcz(meta_linspace_logspace..sDu+r$s),,-RU 0D0D/EQHr8cy)Nz$number of steps must be non-negativer2r2r8r6rcz(meta_linspace_logspace..rkr8metar\layoutdevice pin_memory requires_grad)rnrTr redimanyrMrget_default_dtypeis_complex_dtyper\ _check_typerempty) r{rzr|baser\rrrrrts ``` ` @r6meta_linspace_logspacerrsL%& IIK1  P #u||$ GGIN P  CsE/B CC % A A  # # %!  =)E LL&&u-z  2002 eU[[ )) ) 5'" H eW %% % LL!KL ;; #  r8ctjjtjk(fdtj|j dk(xrj dk7 d|j jS)Nc"djS)Nz2take(): Expected a long tensor for index, but got r\indexsr6rczmeta_take..sDU[[MRr8rcy)Nz*take(): tried to take from an empty tensorr2r2r8r6rczmeta_take..rkr8)rTrer\long _check_indexnumel new_emptyshape)selfrs `r6 meta_takersm LL uzz!R   ZZ\Q  55;;=A#56< >>%++ &&r8rcTj}j}tj||k(dtjjdk(xrjdk(fdt j j }j |S)Ncy)Nz=linalg.cross: inputs must have the same number of dimensions.r2r2r8r6rczlinalg_cross..rkr8r.cVddjdjS)Nzlinalg.cross: inputs dimension z must have length 3. Got  and size)rotherrsr6rczlinalg_cross..s6-cU399S>"% 3'8 :r8)ndimrTrerr#rr)rrrx_dy_d out_shapes``` r6 linalg_crossrs ))C **C LL s O LL #!4 31 4 "$**ekk:I >>) $$r8c|t|dt|dtj|tjS)Nzlinalg.matrix_exp memory_format)squareCheckInputscheckFloatingOrComplexrT empty_likecontiguous_formatrs r6linalg_matrix_exprs3d/04!45   D0G0G HHr8valuesindicescZtj|j|j|j}tj|j|jtj }|j dk7r%|jdk7rt||j||fS)Nrr\r) rTrrrr\int64rrmaybe_wrap_dim)rrrrs r6 cummaxminrsp [[DKKtzz JFkk$**T[[ LG zz|qTYY!^sDII& 7?r8cxt||jtj|tjSNr)rrrTrr)rrs r6 logcumsumexprs+3 "   D0G0G HHr8cD|j}t|}||z }tt|}t|D cgc]} d} } |D]} d| | < gg} } |D]*} | | s| j | | j | ,| | z}t| }|j |d|}|j fdd|||dz}|j|}dgt|j|dz}|j|}|jd}||d<t|}tt|D]}|||||dz<|j|tjt|D cgc]} d}} d}|dz }|dk\r0||j dz|||<||||z}|dz}|dk\r0t||D]}|j d||z z|||<!|j|||j|Scc} wcc} w) NFTc|SrBr2)r; self_stridess r6rcz_exec_fft.. s <?r8keyreverserrr1r)rlenlistrangeappendstridesortpermuterreshaperresize_rTr as_strided_storage_offset)outr out_sizesrforwardr signal_ndim batch_dims dim_permuterPis_transformed_dimdleftright batch_endtmpinput batched_sizes batch_sizebatched_out_sizesi out_strides batch_numelrs @r6 _exec_fftrsZ 99Dc(K #JuT{#K).t5A%55 % $1%b%D !!$ KKN LLO  ,KD I;;=L jy !CHH*DH9 IJ//K LL %ED4 JK 899M MM- (EAJ!M!]+ 3s8_5#,SV#4!a% 5KK!1H1HKI$Dk*1*K*KQA q&&1CJJqM&A KN#yQ00  Q q&:t $G&)jja*n1E&F KN#GOOI{C,>,>,@A JW6@+s H Hrr exclude_lastct|}|j|dt|t|z j fd|S)Nc|SrBr2)rrs r6rcz_sort_dims..,s l1or8)r)rrrintr)rrr sorted_dimsrs @r6 _sort_dimsr(sLs)K;;=L6#k"S%667<< %= r8c tj|jj|s|j St ||}|j |j}t|||j||S)Nr) rTrer\ is_complexclonerrrr)rr normalizationrrrs r6 meta_fft_c2cr3sb LL&&' zz|T3'K .. %C S$ [' JJr8cft|tkDst|dk\r|ddk(r |ddk(ryy)Nrr1FT)rcufft_max_ndimrs r6use_optimized_cufft_pathrBs3 3x. SX]s1v{s1vQR{r8cztj|jjt |j }t |}|d}||dzdz}t |}|||<|r|||<t |dk(st |dk(rz|j|tj|j} |} t |dk(rt|rt| | ||dnt|dk(r|n|} t| | | |gdt|dkDr0|j|tj|j} |dd} | rx| | } } | j| jfd d tt t| } | t| | z d}t| | ||d| dt| | z } | rx|s:| j |||k7r#| j#|tj$ | } | S|j|tj|jS) Nrrr1cudaxpurTrc|SrBr2)rstridess r6rczmeta_fft_r2c..ps '!*r8rr)rTrer\is_floating_pointrr device_hintrrMrrrrrrminrrr)rrronesided input_sizesrlast_dimlast_dim_halfsizeonesided_sizesoutputworking_tensor target_sizesrmax_dims last_dimsrs @r6 meta_fft_r2crIs> LL--.tyy{#K[!I2wH#H-2Q6+&N0N8/ (4F"k$&75&@ U>>tzzJ   t  &+CC+H fnid K),CA 9>L fnlXJPT U3x!|!%U%F%Ftzz%R"0" cr(K)7(//1  ,d!~s;/?@'K(88(C(EF NNIt**GC ,>tzzJ  r8) generatorcBt|tj|gSrB)rrTSize)nrrs r6 meta_randpermrs S%**aS/ 22r8r\rrrc6tj|||||SNrrTr)rr\rrrs r6meta_randperm_defaultr s  ;; vf r8cxdtjkDfdtj|||||S)NrcddSNz:random_ expects 'from' to be less than 'to', but got from=z >= to=r2highlowsr6rczmeta_randint..LSEQXY]X^_r8rrTrer)rrr\rrrrs` @r6 meta_randintrs> C LL s _ ;; E&J r8cttjkDfdtj|||||S)NcddSrr2rsr6rcz"meta_randint_low..rr8rr)rrrr\rrrs`` r6meta_randint_lowrs9 LL s _ ;; E&J r8c6tj|||||Sr r )rr\rrrs r6meta_rand_defaultrs  ;; E&J r8rlastdimctj|jjt |dk(rt |j }|||d<|j|t|j}t|r.t||jtj||dSt|dkDrt||ddd|}n |jtj}t||||dgdS|}t|dkDr|dd}t|||d}|dd}t |j }|||d<|j|t|j} t| |||dS) NrrrrFrr1r)rTrer\rrrrrr^rrrrrr) rrrrrrtemprc2c_dimsrs r6 meta_fft_c2rrsk LL&&'4F"% $ #b' 1LM #C ( )@)@ A  3x!|#D#cr(Aw?zz0G0GzHVT9s2wiO O s8a<3BxH xNEbc(C& $ #b'nnYodjj.InJeYUCCr8cJddlm}||s#tj|dk(r t dt |t ra|j||}|j|jk7r.tjj||j|S)Nr)free_unbacked_symbolsr1zQmore than one element of the written-to tensor refers to a single memory location) %torch.fx.experimental.symbolic_shapesr rT_debug_has_internal_overlap RuntimeErrorrnr torr+ expand_copydefault)rsrc non_blockingr  intermediates r6 meta_copy_r*s L "$ 'E,M,Md,SWX,X _  #vvvdL1 99;,++- -    $ $\499; ? Kr8ct|j}t|j}||jk\rdn ||||z}|j |d|j ||||fSr0)rrrrinsert)tensorr result_sizesresult_strides new_strides r6inferUnsqueezeGeometryr1 sq &L&--/*NVZZ\)|C/@>RUCV/VJQ#z*  ''r8ct||jdz}t||\}}|j|||Sr0)rrr1r)rrg_sizes g_stridess r6meta_unsqueeze_r5s> dhhj1n -C/c:GYWi( Kr8rweight_metabias_activation_opt out_dtypec8t|j}|*|jd|jdk(sJd|jd|jddz k(sJ|jd|d<t|jdk(sJdd|jdf}|7|jt j k(r|t jk(sJd|j|| |jn|j||}|S) Nrzoutput size mismatchr1rrz*we can only handle the squashed input case9out_dtype is only supported for i8i8->i32 linear operatorr) rrrrr\rTint8int32r as_strided) rr6r7r8r9r: output_sizestransposed_stridesrs r6meta_sparse_structured_linearrB!s $L {{1~1-E/EE- ;;q>UZZ^a/ // /{{1~L u{{ q N"NN UZZ]+{{ejj(Y%++-E G E__&.ekkIj12 Mr8mat1 mat1_metamat2ct|jdk(sJt|jdk(sJt|jdk(sJ|jd|jddz k(sJ|jd|jdg}|7|jtj k(r|tj k(sJd|j|| |jn|}|S)Nrr1rr<rrrrr\rTr=r>r)rCrDrEr:r@rs r6meta_sparse_structured_mmrHCs tzz?a   y 1 $$ $ tzz?a   99Q<499Q z matmul!r) r\rTfloat32float16bfloat16r= float8_e4m3fnrr is_contiguousrr>r) rMrNr8rIr:rOrPrQrRis_8bit_input_typerm output_shapes r6meta__cslt_sparse_mmr\s ==       E E E     .Q0QQ . w}}  "M$MM "%++ E>+ &&r8)lengthsroffsetsaxisunsafeinitialdatarjrkrlrmc| tdfd}|||jS|+|jdd|jddz fz} || Std)Nz?segment_reduce(): indices based reduction is not supported yet.ctj|jdzdzjdtjS)Nr1r~r\rr)rTrrr\r) lengths_shaperlros r6segment_reduce_lengths_tensorz:meta_segment_reduce..segment_reduce_lengths_tensors>{{ DJJtaxz2 2**11   r8rr1z>" r8ctj|j|f}t|||}|j ||j |t j fSNrrMreduction_dimsr_compute_reduction_shaperrTrrrkeepdimr[s r6 meta_max_dimrR   tzzC6 2C+D#w?L |$ |5::6 r8c$|jdSrxryrs r6meta_minrr{r8ctj|j|f}t|||}|j ||j |t j fSr}r~rs r6 meta_min_dimr rr8c|jrt|j}nt|tj \}}t j||SNrLr)rrr\rr INT_TO_FLOATrTr)rrQrPs r6 meta_anglersI / ; ,  ? L L <   D 55r8ctj||j|j|j tj |SrB)rT_resize_output_rrcopy_angle)rrs r6meta_angle_outr"s6 #tyy{DKK8 99U[[& ''r8cyrBr2)vals r6 assert_asyncr( r8cyrBr2)r assert_msgs r6assert_async_metar-rr8cyrBr2)ss r6 print_metar2rr8r\rrrrc0tjddS)Nrr~rr rs r6make_dep_tokenr7s ;;q ((r8chddlm}t|ttfr t d||||y)Nr)constrain_range'Constraining SymFloat or Symbool is nyirmax)r!rrnr r ValueError)rrrrs r6sym_constrain_rangerCs/F$7+,BCCDcs+r8c6tj||||SNr)r+rrrr dep_tokens r6functional_sym_constrain_rangerMsTs4 r8c.ddlm}||tj|dk\yt |t t fr tdt|tur5|tj||k\|tj||ky||||y)Nr)_constrain_range_for_sizerr) r!rrTrernr r rrxr)rrrrs r6sym_constrain_range_for_sizerSsP {s{ TQY$7+,BCC DzS ? LL % ? LL %d5r8c6tj||||Sr)r+rrs r6'functional_sym_constrain_range_for_sizergs%%d%= r8c|SrBr2)rrrs r6functional_assert_async_metarms r8f_namec|jdk\s J|d|jd|jdk(s.J|d|jdd|jddy)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:? (EF? 99R=DIIbM ) (G RT VZ[_[d[deg[hZiirs )r8Anamectjjjk(fdtjjjk(fdtjj dj dk(fdtjj dj dk(fdy)Nc>djdjdS)Nz:Expected b and A to be on the same device, but found b on z and A on instead.rrrsr6rcz(linearSolveCheckInputs..s%H{{m:ahhZy :r8c>djdjdS)Nz=Expected b and A to have the same dtype, but found b of type z and A of type rrrsr6rcz(linearSolveCheckInputs..s%Kzzl/!'') =r8rrcRdjddjddS)Nz3A must be batches of square matrices, but they are rrrrrrsr6rcz(linearSolveCheckInputs..s0FF2J.sR,TF3 D $TYYr]O4 "  Hr8)rTrerr\r)rrrs```r6linearSolveCheckInputsrs LL qxx  LL agg  LL r affRj   LL r diim# r8tallow_low_precision_dtypescJ|jtj|jxs|j fd|sYtjtj tj tjtjfvfdyy)NcdS)Nz<: Expected a floating point or complex tensor as input. Got r2r\rsr6rcz(checkFloatingOrComplex..s6(VW\V]^r8cdS)Nz*: Low precision dtypes not supported. Got r2rsr6rcz(checkFloatingOrComplex..svhHPr8) r\rTrerrrXrZrWrY)rrrr\s ` @r6rrsn GGE LL /^ & ekk5<<u}}M M P  &r8arg_namec^tj|jdk\fdy)NrcddS)Nz: The input tensor z! must have at least 2 dimensions.r2)rrsr6rczcheckIsMatrix..s6(-hZ7XYr8)rTrer)rrrs ``r6 checkIsMatrixrs LL 1 Yr8Brctttjr#j dj dk(n"j dj dk(fdy)Nrrc drdnddjddjddjddjdd S) Nz2: Incompatible shapes of A and B for the equation zAX = BzXA = Bz (rr;rrrwr)rrrrsr6rcz#checkInputsSolver..s[hHxX. AaffRj\qvvbzl!AFF2J.s5i{;-/dm4 nU\\N Lr8)rTrer)rrrrs````r6checkSameDevicers&  LL % r8UPLOcj}tjtdk(xr |dk(xs|dk(fdy)Nr1ULcdS)Nz1Expected UPLO argument to be 'L' or 'U', but got r2)rsr6rczcheckUplo..sCD6Jr8)upperrTrer)rUPLO_uppercases` r6 checkUplors<ZZ\N LL D QKNc1J^s5JJr8 eigenvalues eigenvectorsr compute_vcTt|dt|t|j}|r/|j |}|j |t |dn|j dg}|j|j |t|j}||fS)Nz linalg.eighF row_majorrr) rrrrrrrpopr^r\)rrrrvecsvalss r6meta__linalg_eighrsa' dO ME{{5!  ;EU ST{{A3 IIK ;;uOAGG$<; =D :r8ct|dtj|jr |jntj|j}|j |j dd|S)Nzlinalg.eigvalsrr)rrMrr\rrr)r complex_dtypes r6meta__linalg_eigvalsrsce-.  ! !%++ .   . .u{{ ; ??5;;s+=? AAr8ct|dtj|jr |jntj|j}|j |j dd|}|j |j |}t|dk(}|j|j t|j |||fS)Nz linalg.eigrrrr) rrMrr\rrrrrr)rrrvectorsis_cudas r6meta_linalg_eigrse\*  ! !%++ .   . .u{{ ; __U[["-]_ CFooekko?G% F*G  0P 7?r8r'cv|jjtjj ddS)Nrrr)mTrrTr transpose)r's r6cloneBatchedColumnMajorrs* 66< H HR PPr8rct|SrB)r)rrrs r6_cholesky_solve_helperrs #4 ((r8ctjjdk\fdtjjdk\fdtd\}}t |||S)Nrc$djdS)Nz-b should have at least 2 dimensions, but has  dimensions insteadrrsr6rcz cholesky_solve.. s? {J]^r8c$djdS)Nz-u should have at least 2 dimensions, but has rrrsr6rcz cholesky_solve..$s?xGZ[r8cholesky_solve)rTrer!_linalg_broadcast_batch_dims_namer)rrrself_broadcasted A_broadcasteds`` r6rrsh LL Q^ LL ! ['H a!'#m ""2M5 IIr8c|jdk(r%tj|tjSt |dt |S)Nrrcholesky)rrTrlegacy_contiguous_formatrrrrs r6rr,s@ zz|qE4R4RSSdJ' "4 ((r8c0t|dt|S)Ncholesky_inverse)rrrs r6rr5sd./ "4 ((r8 check_errorsct|dt|d|j}t|}t |d}|j |}|j |||j |d|dz tj}||fS)Nzlinalg.choleskyFrrr) rrrrrrrrTr>)rrrA_shaper L_stridesrinfoss r6linalg_cholesky_exr =sa*+1/0ggG w.Vrkr8rrcy)Nzbtorch.linalg.householder_product: input.shape[-2] must be greater than or equal to input.shape[-1]r2r2r8r6rcz,linalg_householder_product..Zrkr8cy)Nz`torch.linalg.householder_product: input.shape[-1] must be greater than or equal to tau.shape[-1]r2r2r8r6rcz,linalg_householder_product..^rkr8r1c<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 rrr sr6rcz,linalg_householder_product..c')), 2Nuzzl \r8cdS)Nzltorch.linalg.householder_product: Expected batch dimensions of tau to be equal to input.shape[:-2], but got r2actual_batch_tau_shapesr6rcz,linalg_householder_product..m66L5MOr8c<djdjS)Nz,torch.linalg.householder_product: tau dtype z does not match input dtype rrsr6rcz,linalg_householder_product..us#:399+*5;;- 9r8z torch.linalg.householder_productr Frrrr\r) rTrerrrr\r empty_stridedrr)rr expected_batch_tau_shapers`` @r6linalg_householder_productrOsK  LL aZ LL 2%**R.(t LL 2#((2,&r  LL SXX"  zzA~#(;;s#3 !$3B "&> >   LL U[[  6UEJ    [[*5;;%Hkk||  r8c2t|dt|dd|j|j}|j |jt |jd|j|jddt j}||fS)Nz linalg.inv_exF)rrrrrrrrrrrTr>)rrrr s r6linalg_inv_ex_metar sqa)1o%P AGGAMM!''6qww%PQ KK EKKK 8E e8Or8LDpivotsinfo) hermitianrr$ct|dt|dtj|jt |jd|j |j}|j|jddtj}|j|jddtj}|||fS)Nztorch.linalg.ldl_factor_exFrrrrr) rrrTrrrr\rrr)rr$rr!r"r#s r6linalg_ldl_factor_ex_metar&sd894!=>    ZZ*4::Gjj{{  B ^^DJJsO599^ =F >>$**Sb/> ;D vt r8)r$cdtdtdtdtjj dk\fdj dd}tj|j k(fdtjtjjfdtjjjk(fdt\}}tj|t|d jj S) Nztorch.linalg.ldl_solverc$djdS)NzMtorch.linalg.ldl_solve: Expected B to have at least 2 dimensions, but it has rr)rsr6rcz'linalg_ldl_solve_meta..&&!4 6r8rc$djdS)Nzjtorch.linalg.ldl_solve: Expected LD.shape[:-1] and pivots.shape to be the same, but got pivots with shape  insteadrr"sr6rcz'linalg_ldl_solve_meta..))/h @r8c"djS)Nz.sNv||n]r8c<djdjS)Nz!torch.linalg.ldl_solve: LD dtype z does not match b dtype r)rr!sr6rcz'linalg_ldl_solve_meta..s"3BHH:=UVWV]V]U^_r8Frr)rrrrTrerrrMis_integer_dtyper\_linalg_broadcast_batch_dimsrrr)r!r"rr$expected_pivots_shapeB_broadcast_sizerPs``` r6linalg_ldl_solve_metar5sb232781b":; LL !  HHSbM LL-  LL v||,] LL AGG_7q"=a    *+;uMggxx  r8Pr)pivotr7chtjjdk\fdtj}|d}|d}t ||}||d<|rj |}nj dg}||d<j |}||d<||d<j |}|||fS)Nrc$djdS)Nz@linalg.lu: Expected tensor with 2 or more dimensions. Got size: r+r,rsr6rcz linalg_lu_meta..sRSTSZSZR[[cdr8rrr)rTrerrrrr) rr7sizesrZrkr6rrs ` r6linalg_lu_metar<s LL ! d ME b A b A Aq AE"I KK  KK E"I EAE"IE"I EA a7Nr8LU)r7rctjjdk\fdtj}|d}|d}tj |t |djj}|jt|||d<j|tj}|jj|tj}|||fS) Nrc$djdS)NzFtorch.lu_factor: Expected tensor with 2 or more dimensions. Got size: r+r,rsr6rcz*linalg_lu_factor_ex_meta..sXYZY`Y`Xaaijr8rrFrrr) rTrerrrrrr\rrrrr) rr7rr:rZrr=r"r#s ` r6linalg_lu_factor_ex_metar@s LL ! j ME b A b A    *5EBggxx  B IIKAq E"I [[eii[ 0F IIK ;;uEII; .D vt r8)radjointrActdtjjjk(fdtjjtjk(dt dt |dtjjdjdk(dtjjddjk(fdt\}}tj|t|| jj }|jd k7r"|s |jr|j}|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 r+r)rr=sr6rcz&linalg_lu_solve_meta.. s($$&HH:_QWWIX Or8cy)NzElinalg.lu_solve: pivots should be a Tensor of scalar type torch.int32r2r2r8r6rcz&linalg_lu_solve_meta..'rkr8zlinalg.lu_solvercy)NzYlinalg.lu_solve: Number of pivots per batch should be same as the dimension of the matrixr2r2r8r6rcz&linalg_lu_solve_meta../rkr8c$djdS)Nzclinalg.lu_solve: Expected LU.shape[:-1] and pivots.shape to be the same, but got pivots with shape r+r,r-sr6rcz&linalg_lu_solve_meta..5r.r8rrr)rrTrer\rrrrrr2rrrrrconj)r=r"rrrAr4rPrs``` r6linalg_lu_solve_metarIs.267 LL AGG  LL  !W b12b!T#45 LL  v{{2&k  LL " % 7q"=a  *+;4xPggxx F||~4    [[]F Mr8 unpack_data unpack_pivotsctjjdk\fd|r2tj|jtjk(dt j }|d}|d}t||}||d<|rj|}njdg}|r2||d<j|} ||d<||d<j|} n$jdg} jdg} || | fS)Nrc$djdS)NzFtorch.lu_unpack: Expected tensor with 2 or more dimensions. Got size: r+r,)r=sr6rcz lu_unpack_meta..UsXY[YaYaXbbjkr8c y)Nztorch.lu_unpack: LU_pivots is expected to be a contiguous tensor of torch.int32 dtype. Note: this function is intended to be used with the output produced by torch.linalg.lu_factorr2r2r8r6rcz lu_unpack_meta..Zs pr8rrr) rTrerr\r>rrrr) r=r"rJrKr:rZrr;r6rrs ` r6lu_unpack_metarOKs LL 1 k LLEKK '   NE b A b A Aq AE"I LL  LL! b LL b b LL  LL!  LL!  a7Nr8modecdk(rd}d}||fSdk(rd}d}||fSdk(rd}d}||fStjdfdfS)NreducedTcompleteFrcddS)Nzqr received unrecognized mode 'z=' but expected one of 'reduced' (default), 'r', or 'complete'r2)rPsr6rcz _parse_qr_mode..s1$8NOr8rTre)rP compute_qrRs` r6_parse_qr_moderXus y  g     g    g      g r8QRct|dt|dt|\}}|jd}|jd}t ||}|rMt |j}|r|n||d<|j |}|j|t|dn|j dg}t |j} |s|s|n|| d<|j | } | j| t| d|| fS)Nz linalg.qrrrFrr) rrrXrrrrrr) rrPrW reduced_moderZrr;Q_shaperYR_shaperZs r6linalg_qr_metar_s![!1k*,T2I|  A  A Aq Aqww-'aQ KK  g:7eTU KK 177mG#9!!GBK GAMM'6w%PQ a4Kr8sign logabsdetct|dt|dd|j}|j|dd}|j|ddt |j }t j|t|d|j |j}|j|ddt j}||||fS)Nzlinalg.slogdetFrrrr) rrrrr^r\rTrrrr>)rrr`rar=r"s r6_linalg_slogdetrcsa)*1.6 GGE ;;uSbz "D E#2Joagg.F GI    *5%8ggxx  B [[s5;;[ 7F B &&r8 full_matrices compute_uvdrivercbt|dt|dt|jdd}|jd}|jd}t ||}|r|||r|n|gz}|j |} | j |t|d||r|n||gz} |j | } t|dk(} | j | t| | n$|j dg} |j dg} |j ||gzt|j} | | | fS) Nz linalg.svdrrFrrrr) rrrrrrrrrr^r\)rrdrerfrrZrr;U_shaperV_shapeVrSs r6_linalg_svd_metarls#!\"1l+aggcrl#J  A  A Aq A 11== KK  g:7eTU]1== KK  a.F* g:7gVW KK  KK  J!$OAGG,D EA a7Nr8arg1arg2c|jdd}|jdd}t||}t|}||jd|jdgz }t|}||jd|jdgz }||fS)Nrr)rr#rr)rmrnarg1_batch_sizesarg2_batch_sizesexpand_batch_portionarg1_expand_sizearg2_expand_sizes r6r2r2s zz#2zz#2,-=?OP012 " 66012 " 66 - --r8c|r t|||t||\}}||jk(r|n|j|}||jk(r|n|j|}||fSrB)rr2rexpand)rmrnrrsrtarg1_broadcastedarg2_broadcasteds r6rrsv  tT40)EdD)Q&&!DJJ.DKK@P4Q!DJJ.DKK@P4Q - --r8rc|jdd}|jdk(xs-|jdz |jk(xr|j|k(}|S)Nrr1)rr)rrexpected_batched_rhs_shape vector_cases r6linalg_solve_is_vector_rhsr|sS!&Sb!1**/ Q%**$R8R)R r8)rrrr=r"r#ctdtjjjk(fdt }|rj dn} t | |dt| \} } tj|xs| d|r| ddn| } tj| t| | jj} j}tj|t|djj}j|ddtj}j|ddtj}||||f}| |||f}td |Drbt||D]S\}}t!||j|j#|j|j%t'||d U|S) Nz linalg.solvec>djdjdS)NzKlinalg.solve: Expected A and B to have the same dtype, but found A of type rDr+r)rrsr6rcz"_linalg_solve_ex..s%Ywwiqwwix 9r8rc y)Nzlinalg.solve: Vector broadcasting of the left hand side is not supported for left=False. In this case linalg.solve is equivalent to B / A.squeeze(-1)r2r2r8r6rcz"_linalg_solve_ex..'s  Kr8rFrrc3$K|]}|du ywrBr2)rpr;s r6rrz#_linalg_solve_ex..>s &Q1D= &s) copy_fromcopy_to exact_dtype)rrTrer\r| unsqueezerr2rrrrrr>allziprrrr!)rrrrrr=r"r#r{B_ B_broad_shaperP result_shaperesult_rLU_pivots_info_rresrTos`` r6_linalg_solve_exr s1n- LL 177 -Q2K'RQBaT>23B:M1 LL K *5="%-L!! *.ms))$7 9r8c$djdS)NzMtorch.triangular_solve: Expected A to have at least 2 dimensions, but it has rrrsr6rcz'triangular_solve_meta..tr)r8triangular_solveFrrrcy)Nz+triangular_solve: Got an unexpected layout.r2r2r8r6rcz'triangular_solve_meta..rkr8)rTrerrrstridedr2rrr\r sparse_csr sparse_bsrrr) rrrrrself_broadcast_sizeA_broadcast_sizesolutioncloned_coefficients `` r6triangular_solve_metarbsL LL Q  LL !  4$67xx5== 0LTST0U--&&$./BeT**;;   #00!./?5Q''88   ' '' U%% %U5E5E)E##D)!^^QC0 ' ''  UQR ' ''r8clt|dt|d|j|jdd}|j|j}|j |jt |jd|j|jddt j}|||fS)Nz linalg.detrFrrrr)rdetr=r"s r6_linalg_det_metarsa&1l+ ++aggcrl #C QWW BNN17775QR [["U[[[ 9F F?r8c0tjjdk\dtjjdk\d|rdndtjjjdk\fdtjjjdk(fdtjjdjdkdtjjjz d k(fd tjjjk(fd jdkDrejdd}jddtj|k(fd jddtj|k(fd tjjjk(fdtjjjk(fdt ddt ddtj jtjdjjS)Nrcy)Nz3torch.ormqr: input must have at least 2 dimensions.r2r2r8r6rczormqr..rkr8cy)Nz3torch.ormqr: other must have at least 2 dimensions.r2r2r8r6rczormqr..rkr8rrcddS)Ntorch.ormqr: other.shape[z0] must be greater than or equal to tau.shape[-1]r2left_size_conditionsr6rczormqr..s+,?+@@pqr8cddS)Nrz"] must be equal to input.shape[-2]r2rsr6rczormqr..s+,?+@@bcr8cy)NzHtorch.ormqr: tau.shape[-1] must be less than or equal to input.shape[-1]r2r2r8r6rczormqr..rkr8r1c<djdjS)Nz[torch.ormqr: Expected tau to have one dimension less than input, but got tau.ndim equal to rrrsr6rczormqr..rr8c<djdjS)Nzhtorch.ormqr: Expected other to have the same number of dimensions as input, but got other.ndim equal to rrrrsr6rczormqr..s+++0::,6RSXS]S]R^ `r8cdS)NzWtorch.ormqr: Expected batch dimensions of tau to be equal to input.shape[:-2], but got r2rsr6rczormqr..rr8cdS)NzYtorch.ormqr: Expected batch dimensions of other to be equal to input.shape[:-2], but got r2)actual_batch_other_shapesr6rczormqr..s66N5OQr8c<djdjS)NzPtorch.ormqr: Expected input and tau to have the same dtype, but input has dtype z and tau has dtype rrsr6rczormqr..s'##(;;-/B399+ Or8c<djdjS)NzRtorch.ormqr: Expected input and other to have the same dtype, but input has dtype z and other has dtype rrsr6rczormqr..s'##(;;-/DU[[M Sr8z torch.ormqrr rFrr) rTrerrr\rrrr) rr rrrexpected_batch_shaperrrs ``` @@@r6ormqrrs LL aV LL aV!%"" LL '(CIIbM9q LL '(EKKO;c  LL " R(Z  LL SXX"  LL ejj   zzA~${{3B/!$3B "&: :  $);;s#3  $(< <   LL U[[   LL u{{" M3u5M5%9    [[*5;;%Hkk||  r8crtjtdzk(fdj}|dzk(}|}| }|r*t d|D]}|xrj |dk7}n(t |D]}|xrj |dk7}tj|xs|fdy)Nrc,ddzdtS)Nzpadding size is expected to be rz , but got: r)rpaddingsr6rcz,_padding_check_valid_input..s1!c'+c'l^Tr8r1rc:ddzddzdjS)N Expected r1zD or rzcD (batch mode) tensor with possibly 0 batch size and other non-zero dimensions for input, but got: r,)rrsr6rcz,_padding_check_valid_input..s2aycAgY/AAF  Or8)rTrerrrr)rrr input_dim is_batch_modevalid_batch_modevalid_non_batch_moders``` r6_padding_check_valid_inputrs LL G CT  I#'*M$,,q)$ GA/FEJJqMQ4F  Gy! OA#7#NEJJqMQ.12%%*G2eWOE7*UZU`U`Tacr8cddS)Nz input (W: z%) is too small. Calculated output W: r2)input_woutput_wsr6rcz_pad1d_common..9s*WI%J8*Ur8r)rrrrTrer) rr is_reflection dim_planenbatchnplanerrrrrs ` @@@@@r6 _pad1d_commonrsI E F zzQA  Q ug15LE5 ZZ "FjjG&H GO /   LLA U  zzQ1229::r8ct||dSNTr)rrrs r6meta_reflection_pad1drB t <.M =ekk>Q>Q>S=TTUXr8Fr)rTrer\boolrrs` r6meta_replication_pad1drH5 LL uzz!X u ==r8cd|s#tjt|dk(djdk(rdz |\j }|zz|r&tj|kxr|kfdtjj k(fdj j S)Nr1rcy)Nz padding size is expected to be 2r2r2r8r6rcz(_pad1d_backward_common..Urkr8r.c4ddddjSrr,rsr6rcz(_pad1d_backward_common..brr8c2ddjSNz(grad_output width unexpected. Expected: , Got: rr grad_outputrsr6rcz(_pad1d_backward_common..j":8*GKL\L\]bLcKder8rTrerrrrr) rrrrrrrrrs `` @@@@r6_pad1d_backward_commonrRs E  S\Q&(RS zzQ  LE5jjG&H GO /   LLK$$U++e ??5;; ''r8 grad_inputc t|||dSrrrrrs r6meta_reflection_pad1d_backwardrps "+ugT RRr8c t|||dS)NFrrrs r6meta_replication_pad1d_backwardrvs "+ugU SSr8c ddd}d}t|dj}|dk(r jd}dz dz |dz }|\ j|}j j z z zz |rLtj kxr kfdtj kxr k fdtj dk\xs dk\ fdjd k(rj | fSj || fS) Nrr1rrc4ddddjSrr,rsr6rcz_pad2d_common..rr8c4ddddjSNzcArgument #6: Padding size should be less than the corresponding input dimension, but got: padding (rvrrr,dim_hrpad_bpad_tsr6rcz_pad2d_common..rr8c ddddS)Nz input (H:  W: z%) is too small. Calculated output H: r2)input_hroutput_hrsr6rcz_pad2d_common..s* gY/$$,:T( =r8r.rrrrTrer)rrr dim_slicesrrrrrrrrrrrrrs` @@@@@@@@@@r6 _pad2d_commonr |sU E EJ Fug15 ::D qyA    a !(E5% ZZ #FjjGjjG&H&H GO /    GO /   LLA &Q  zzQ(;<<(CDDr8ct||dSr)r rs r6meta_reflection_pad2dr rr8ctjjtjk7fdt |dS)Nc@djjdS)Nz)"replication_pad2d" not implemented for 'rrrsr6rcz(meta_replication_pad2d..rr8Fr)rTrer\rr rs` r6meta_replication_pad2drrr8c^tj|}tj|}||fSrBrTr)grad_wsaved_vsaved_g saved_normsrgrad_vgrad_gs r6meta_weight_norm_backwardrs.  g &F   g &F 6>r8c d d d}|j}|jdk(r dz dz |dz }|\}}}}| } | } | |z|z | |z|ztjj k( fdtj j k( fd|j |jS)Nrr1rrc2ddjSrrrsr6rcz%meta_pad2d_backward..rr8c2ddjSNz)grad_output height unexpected. Expected: rrrrrsr6rcz%meta_pad2d_backward..";H:W[M]M]^cMdLefr8)rrrTrerr)rrrrrbrrrrrrrrrrs` @@@@r6meta_pad2d_backwardr s E EIJ xxzQ    Q !(E5%GG&H&H LLK$$U++e LLK$$U++f >>$** %%r8c $ d ddd}t|djdk(}|r%jd} dz dz dz |dz }|\j|}j j j  zz zz zz|rrtj kxr k fdtj kxr kfdtj kxr kfd tjdk\xs dk\xs dk\ fd |rj | fSj | fS) Nr.rr1rrc4ddddjSrr,rsr6rcz_pad3d_common.. rr8c4ddddjSrr,rsr6rcz_pad3d_common..rr8c4ddddjS)NzcArgument #8: Padding size should be less than the corresponding input dimension, but got: padding (rvrrr,)dim_drpad_bkpad_fsr6rcz_pad3d_common..s2%%*G2fX_UG:V[VaVaUbdr8c ,dddddd S)Nz input (D:  H: rz%) is too small. Calculated output D: r2)input_drroutput_drrsr6rcz_pad3d_common..#s7 gYd7)<$$,:T(4z Kr8r )rrrr batch_moderrr&rrr+rrr,rrrr'r(rrrs` @@@@@@@@@@@@@@@r6 _pad3d_commonr.s E E EIug15qJA      Q 07-E5%v ZZ "FjjGjjGjjG'H&H&H GO /    GO /    GO 0 0   LLA 7Q7(a-  (HMNN(HEFFr8ct||dSr)r.rs r6meta_reflection_pad3dr0/rr8ctjjtjk7fdt |dS)Nc@djjdS)Nz)"replication_pad3d" not implemented for 'rrrsr6rcz(meta_replication_pad3d..:rr8Fr)rTrer\rr.rs` r6meta_replication_pad3dr35rr8c tjt|dk(d|jdkDsJj|jk(sJdd d |jdk(rdz  dz dz |\}}}}}}|j } |j } |j } | |z|z| |z|z| |z|ztjj k(fdtjj k( fdtjj k( fd |j |j S) Ncy)Nz padding size is expected to be 6r2r2r8r6rcz%meta_pad3d_backward..Irkr8r.rr1r"c2ddjSrrrsr6rcz%meta_pad3d_backward..arr8c2ddjSrrrsr6rcz%meta_pad3d_backward..err8c2ddjS)Nz(grad_output depth unexpected. Expected: rr)r&rr,sr6rcz%meta_pad3d_backward..irr8r)rrrrrrrr(r'r+rrr&rrr,rrs` @@@@@@r6meta_pad3d_backwardr:?s_ LLW"$NO ::>>   uzz )) ) E E E zzQ      07-E5%vjjGjjGjjG'H&H&H LLK$$U++e LLK$$U++f LLK$$U++e ??5;; ''r8pcJtj|jd|jd}|dkr0|j dgj tj S|j ||dz zdzfj tj S)Ncy)Nz(_pdist_forward requires contiguous inputr2r2r8r6rcz%meta__pdist_forward..srkr8rr1rr)rTrerXrrr$r)rr;rs r6meta__pdist_forwardr>os LL P ! AAv~~qc"%%E4R4R%SS~~qAE{a/1255886  r8gradpdistctj|jdtj|jdtj|tjS)Ncy)Nz._pdist_backward requires self to be contiguousr2r2r8r6rcz&meta__pdist_backward..rkr8cy)Nz/_pdist_backward requires pdist to be contiguousr2r2r8r6rcz&meta__pdist_backward..rkr8r)rTrerXrr)r?rr;r@s r6meta__pdist_backwardrD~sW LL V LL X   D0N0N OOr8)rJrIc 0 ddlm}m}jd}jd}jd} |t j |j ||| frj||| ft jjdk(dt jjdk(dtjsGt jjjcxk(xrjk(ncfdj } j | d | d t j d k(xr d k( fd jjS) Nr) guard_or_truesym_eqr1rr.cyNzbatch1 must be a 3D tensorr2r2r8r6rczmeta_baddbmm..rkr8cyNzbatch2 must be a 3D tensorr2r2r8r6rczmeta_baddbmm..rkr8cVdjdjdjS)Nz+Input dtypes must be the same, got: input: z , batch1: z , batch2: r)batch1batch2rsr6rczmeta_baddbmm..s0A$**ZX^XdXdWeeopvp|p|o}~r8c .ddddddd SNz@Expected size for first two dimensions of batch2 tensor to be: [rvz ] but got: [rr1].r2 batch2_sizesbscontraction_sizesr6rczmeta_baddbmm..s:t2&'|LO3DB|TUFWWY [r8)r!rFrGrrTsym_notrrvrer exp_config&skip_dtype_check_in_meta_registrationsr\r)rrMrNrJrIrFrGdim1dim2dim3 batch1_sizesrSrTrUs``` @@@r6 meta_baddbmmr]s7L ;;q>D ;;q>D ;;q>DU]]6$**tT46H#IJK{{D$-. LL"$HI LL"$HI  < < JJ&,, 6&,, 6 ~ <>$))+ &&r8cLtj|tjSrrarrs r6meta_bernoullir`s   D0G0G HHr8c|SrBr2rr;rs r6meta_bernoulli_rc Kr8cLtj|tjSrrarbs r6meta_bernoulli_prfs   D0G0G HHr8c,tj|SrBrr_s r6 meta_poissonrh   D !!r8c tj| |jkdtj|tj} tj|| fS)Ncy)NzJError in fused_moving_avg_obs_fake_quant_cpu: ch_axis must be < self.dim()r2r2r8r6rcz6meta__fused_moving_avg_obs_fq_helper..rkr8r)rTrerrr)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_helperrysO  LL$((*\   D 3D   T "D ))r8c@tj|jdk(dtj|jdk(d|j\|j\tjk(fd|gtj||jk(xsA|tj k(xr,|jtj tjfvd| |jn|}|jf|S)Nrcy)Nz a must be 2Dr2r2r8r6rczmeta_mm..rkr8cy)Nz b must be 2Dr2r2r8r6rczmeta_mm..rkr8c "ddddd S)Nz/a and b must have same reduction dim, but got [rvz] X [rQr2)M1M2Nr6sr6rczmeta_mm..s(A!Brd%PRtSUVWUXXZ[r8cy)NzFout_dtype must be the same as input dtype or fp32 for fp16/bf16 inputsr2r2r8r6rczmeta_mm..rkr8r) rTrerrr\rTrUrVr)r4r5r:rQr~rrr6s @@@@r6meta_mmrs LLA56 LLA56 GGEAr GGEB LL b[   U]]*?GG u~~>> \  (/177YL ;;1v\; 22r8c|r(tfdtjDStjj S)Nc3HK|]}|vrj|ndyw)r1Nr,)rprdimsrs r6rrz+_compute_reduction_shape..s$UqatmTZZ]:Us")rdrrrMcompute_reduction_output_shaper)rrrs`` r6rrs7UE$))DTUUU  / / D AAr8ct|tjjr|jj St |drEt |jdr/|jj dk7r|jj Sy)Nrrxr~r)rnrT _subclasses FakeTensor fake_devicerxhasattrr)r-s r6rrsg&%++667!!&&&! FMM6 * MM  & (}}!!!r8 input_tensorrrdilation is_transposedgroupsoutput_paddingc$dtdtdtdtdtdtf d}dtdtdtdtdtdtdtfd } |jd d} |jd d|r||jd z} n<|jd } |jd |z|jd k7r td |jd | gt|tr|gt z}n t |d k(r|d gt z}t|tr|gt z}n t |d k(r|d gt z}t|tr|gt z}n t |d k(r|d gt z}d} |rCt|tr|gt z} n#t |d k(r|d gt z} n|} t t D]]} | r/j| | || || | | || | | 4j|| || || | | || _d dlm }tj|d dDcgc]}|d kD c}fdScc}w)Nlnr;rr;rr>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)rr;rr;rs 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)rr;rr;rrCs r6_formula_transposedz6calc_conv_nd_return_shape.._formula_transposed$ s2"Q!|a!e#a1q5k1B6::r8rr1rzInvalid channel dimensions)sym_orc.dtddddS)NzGiven input size per channel: z&. Calculated output size per channel: rz. Output size is too small)r)r ret_shapesr6rcz+calc_conv_nd_return_shape..v s*0d =//8}o>#$r8) rrr#rnrrrrr!rrTre)rr6rrrrrrrr kernel_size out_channelsoutput_padding_listrrr;rrs @@r6calc_conv_nd_return_shaper s7S7S7S7S7S7S7";;;;;;;QT;&,,qr"K   ab !D Q/ ||A <<?V #|'9'9!'< <;< <##A& 5I&'"CI% V )s4y('7#)c$i' W 1:,T*(G$:D ) X! QK=3t9,,0 ng .#1"2SY">   A %#1!#4"5D "A "0  3t9     #GAJQKN1I'*     a'!*hqk;q>6RS9U #*= LL !" .1Q./ $  /s2 J cbtjj|tjk(SrBrT _prims_commonr channels_lasttens r6is_channels_lastr~ s$    4 4S 9U=P=P PPr8 running_mean running_vartrainingexponential_average_factorepsiloncrj}| |jn |j} | |jn |j} fd} j|j| } |r#j| } j| }n"jd} jd}| | |fS)NctrtjSjtjrtjStjSr)rrTrrXr)rsr6pick_memory_formatz2meta_miopen_batch_norm..pick_memory_format sF L )&& &  % %E4K4K % L** *&&&r8rr)rrr$)rr6r8rrrrrrsave_mean_shapesave_var_shaperr save_meansave_vars` r6meta_miopen_batch_normr s""I-9,Dl((&,,O*5*A[&&v||N'   + . .=O=Q . RC **?; )).9 **40 ))$/  8 ##r8c t||||||||r|nd} d} d} |j| dk(rd| | <|j| } | SNr1r)rrr) rr6r8rrrrrr shape_outinput_channels_dimoutput_channels_dimrs r6 meta_convr sl*'T I+,1)* %&   +C Jr8mkldnnc t|||||d|g} |j| } tj} |j dk(rtj } | j | } | S)NFr"r)rrrTrrchannels_last_3dr$) rr6r8rrrrattrscalars algorithmrrout_memory_formats r6meta_mkldnn_convolution_defaultr sp. &&'8UFB $$Y/!//     " % 6 6 ff#4f5 r8cb|jg|jdd|jdSNrrrr)rr6r8rrrs r6meta_linear_pointwise_defaultr s5%%&Q (:(:3B(?&Qa&QRRr8mklcb|jg|jdd|jdSrr)r packed_weight orig_weightr8rs r6meta_mkl_linearr s:))@,$$Sb)@;+<+. -H QR+Sr8c,djdS)Nzw must be a 2D tensor, got rrrsr6rcz/meta_int4mm_packed_weight_cpu.. rr8c"djSNz#expected x to be f32/f16/bf16, got rrsr6rcz/meta_int4mm_packed_weight_cpu.. s9!''Cr8c"djSNzexpected w to be uint8, got rrsr6rcz/meta_int4mm_packed_weight_cpu.. s.J177),Tr8c"djS)Nz q_group_size must be int64, got r) q_group_sizesr6rcz/meta_int4mm_packed_weight_cpu.. s6|7I7I6JKr8c"djS)Nz5q_scale_and_zeros must have the same dtype as x, got r)q_scale_and_zerossr6rcz/meta_int4mm_packed_weight_cpu.. sKL]LcLcKder8rr) rTrerr\rTrUrVrrrrr;rrrs````r6meta_int4mm_packed_weight_cpur  s QUUW\#ST QUUW\#ST GG u}}ennE E C   GGu{{ "$T     %++ - K    # #qww . e {{166!9affQiqww{??r8ctjjk(xrjk(fdy)NcjdddddjddjzS)NzExpected a tensor of dimension z and tensor.size[z] == rvzbut got : dimension z] = rr)rdim_sizerr-sr6rcz check_dim_size.. sP1#6GzQVW[V\\^_ .?zfll[cNdMe fgr8)rTrerr)r-rr rs````r6check_dim_sizer s6 LL > X 6$ > gr8cd}|d|\}} tjt|dvdtjjtjtj tj tjfvfdt|dk(r|| } } n%t|dk(r |d|d} } n |d|\} } |d |\} } tj|duxs|dk7d jd k(rjd nd}jd }jd}jd}t||| | d|}t|| | | d|}tj}t|| | | | | dd||||||jdk(r|||g}n||||g}tj|jj|S)Nctjt|dvfd|d}t|dk(r|n|d}||fS)Nr1rcddS)Nz avg_pool2d: 4 must either be a single int, or a tuple of two intsr2rsr6rcz1meta_avg_pool2d..unpack.. l4&(\]r8rr1rTrerrrHWs` r6unpackzmeta_avg_pool2d..unpack G H  ]  FSQACF!t r8rrr1rcyNzOavg_pool2d: stride must either be omitted, a single int, or a tuple of two intsr2r2r8r6rcz!meta_avg_pool2d.. rkr8c@djjdS)Nz""avg_pool2d" not implemented for 'rrrsr6rcz!meta_avg_pool2d.. 6u{{7J7J7L6MQQr8rr1rrcyNzdivisor must be not zeror2r2r8r6rcz!meta_avg_pool2d.. rkr8rrrrr.rr)rTrerr\ruint16uint32uint64rrpooling_output_shaperMrpool2d_shape_checkrr)rrrrrcount_include_paddivisor_overriderkHkWdHdWpadHpadWrr inputHeight inputWidthrrrrs` r6meta_avg_pool2dr3 sM; /FB LL F y a LL EKKu||U\\RRQ 6{aRB V F1IB&)B 7+JD$ LLD 9$4$9* %yy{a/UZZ^QF**R.K**R.KBJ' Rr1iPL&z2tRINK//6M         $ yy{a\;7 \;? ;; kk||#  r8ct|||||||dd| | | | | ||j}| }t|||dz |t|||dz | t|||dz | y)Nr1r.r)r(rr)r gradOutputrr+r,r-r.r/r0rr1r2rr mem_formatr nOutputPlanes r6avg_pool2d_backward_shape_checkr8- s"         $ 99;DL:tTAX|<:tTAX|<:tTAX{;r8ctjt|dk(xst|dk(d|d}t|dk(r|n|d} tjt|dk(xst|dk(xst|dk(dt|dk(r|n|d} t|dk(r| nt|dk(r| n|d} tjt|dk(xst|dk(d|d} t|dk(r| n|d} tj|duxs|dk7d|j}|j dk(r|d nd}|d }|d }|d }t ||| | d|}t || | | d|}t j|}t||||| | | | | ||||||tj||j|j| S)Nr1rcy)NzKavg_pool2d: kernel_size must either be a single int, or a tuple of two intsr2r2r8r6rcz*meta_avg_pool2d_backward..g rkr8rcyrr2r2r8r6rcz*meta_avg_pool2d_backward..m rkr8cy)NzGavg_pool2d: padding must either be a single int, or a tuple of two intsr2r2r8r6rcz*meta_avg_pool2d_backward..s rkr8cyr"r2r2r8r6rcz*meta_avg_pool2d_backward..z rkr8rrr#rrrr) rTrerrrr'rMrr8rr\r) gradOutput_rrrrrr)r*r+r,r-r.r/r0 input_sizerrr1r2rrr6s r6meta_avg_pool2d_backwardr@Y s LL KA6[!1Q!6] QB;1$+a.B LL F q@CK1,@F q0@a6{aVAYB6{a3v;!+;RB LL G .S\Q.Y 1:Dw<1$4'!*D LLD 9$4$9* J$yy{a/Z^QFR.KR.KBJ' Rr1iPL&z2tRINK,,U3J#       $ ;;kk||  r8ctjt|dvd|d}t|dk(r|n|d}t|dk(r|n|d} tj| xs t|dvdtjjtjtj tj tjfvfd|s|n|d} |s|nt|dk(r| n|d} |s| nt|dk(r| n|d} tjt|dvd|d} t|dk(r| n|d}t|dk(r| n|d}tjjd vd tj| xs|dk7d jd}jd }jd }jd}jd}t||| | d|}t|||| d|}t|| || d|}t|||| | | | | ||ddd||||||ddjdk(rj||||fSj|||||fS)Nr1r.cyNzFavg_pool3d: kernel_size must be a single int, or a tuple of three intsr2r2r8r6rcz!meta_avg_pool3d.. rkr8rr1rcyNzJavg_pool3d: stride must be omitted, a single int, or a tuple of three intsr2r2r8r6rcz!meta_avg_pool3d.. rkr8c@djjdS)Nz""avg_pool3d" not implemented for 'rrrsr6rcz!meta_avg_pool3d.. r r8cyNzBavg_pool3d: padding must be a single int, or a tuple of three intsr2r2r8r6rcz!meta_avg_pool3d.. rkr8rr"cyNz9non-empty 4D or 5D (batch mode) tensor expected for inputr2r2r8r6rcz!meta_avg_pool3d.. rkr8cyr"r2r2r8r6rcz!meta_avg_pool3d.. rkr8rr#rrz avg_pool3d()T)check_input_sizer) rTrerr\rr$r%r&rrr'pool3d_shape_checkr)rrrrrr)r*kTr+r,dTr-r.padTr/r0rnslicesitimeiheightiwidthotimeoheightowidths` r6meta_avg_pool3drZ s LL KF"X QB;1$+a.B;1$+a.B LL  +c&kV+\ LL EKKu||U\\RRQvayBc&kQ&6F1IBc&kQ&6F1IB LL G T 1:Dw<1$4'!*Dw<1$4'!*D LL fK  LL 5 0A 5* ZZ]FjjnG JJrNEjjnG ZZ^F D"a CE"7Bb!YGG !&"dB9 EF               -2 zzQ@AAHIIr8ctjt|dvd|d}t|dk(r|n|d} t|dk(r|n|d} tj| xs t|dvd|s|n|d} |s| nt|dk(r| n|d} |s| nt|dk(r| n|d} tjt|dvd|d}t|dk(r|n|d}t|dk(r|n|d}tj|jdvd tj| xs|dk7d |j d }|j d }|j d }|j d}t |||| d|}t || || d|}t || || d|}t ||||| | | | | |||||||||d|j|jS)NrBcyrDr2r2r8r6rcz*meta_avg_pool3d_backward.. rkr8rr1rcyrFr2r2r8r6rcz*meta_avg_pool3d_backward.. rkr8cyrIr2r2r8r6rcz*meta_avg_pool3d_backward.. rkr8rJcyrLr2r2r8r6rcz*meta_avg_pool3d_backward.." rkr8cyr"r2r2r8r6rcz*meta_avg_pool3d_backward..' rkr8rr#rrzavg_pool3d_backward()) rTrerrrr'avg_pool3d_backward_shape_checkrr)rrrrrrr)r*rPr+r,rQr-r.rRr/r0rSrTrUrVotime_for_shape_checkoheight_for_shape_checkowidth_for_shape_checks r6meta_avg_pool3d_backwardre s LL KF"X QB;1$+a.B;1$+a.B LL  +c&kV+\vayBc&kQ&6F1IBc&kQ&6F1IB LL G T 1:Dw<1$4'!*Dw<1$4'!*D LL fK  LL 5 0A 5* jjnG JJrNEjjnG ZZ^F0D"aS27Bb!YW1&"dB9U#           ', ??5;; ''r8c,tjjdk(xsjdk(fdjddt |z}t j }tj|jj|S)Nr.rc"djS)Nz"Expected 3D or 4D tensor, but got r,rsr6rcz*meta_adaptive_avg_pool2d..P 4TZZLAr8rrr) rTrerrrdrMrrr\r)r output_sizer[rs` r6meta_adaptive_avg_pool2drjL s| LL Q($))q.A::cr?U;%77L//5M ;;jj{{#  r8ctjjdk(xsjdk(fdjjddt |zS)Nrr"c"djS)Nz"Expected 4D or 5D tensor, but got r,rsr6rcz*meta_adaptive_avg_pool3d..b rhr8r#)rTrerrrrd)rris` r6meta_adaptive_avg_pool3drm^ sO LL Q($))q.A >>$**Sb/E+,>> ??r8cj}td|D].tjj dkDfd0tj|dk(xs|dk(fdtjj j k(fdtj }trtj}jjj|S) Nr1rc*djddS)Nz{adaptive_avg_pool2d_backward(): Expected grad_output to have non-zero size for non-batch dimensions,  with dimension being emptyr,)grad_outrsr6rcz4meta__adaptive_avg_pool2d_backward..m s&66>nn5EEUVWUXXdfr8r.rc"djS)NzBadaptive_avg_pool2d_backward(): Expected 3D or 4D tensor, but got r,rsr6rcz4meta__adaptive_avg_pool2d_backward..r sTUYU_U_T`ar8c<djdjSNexpected dtype z! for `grad_output` but got dtype r)rrrsr6rcz4meta__adaptive_avg_pool2d_backward..v s/$**-Nx~~N^_r8r) rrrTrerr\rrrrrr$)rrrrrrs`` @r6"meta__adaptive_avg_pool2d_backwardrwg s ==D 1d^  MM! q  f   LL  TQYa LL hnn$_++M++ >>$** % ( (} ( EEr8cdt|dtj|tjS)Nadaptive_avg_pool3d_backwardr)!_adaptive_pool_empty_output_checkrTrrrrs r6"meta__adaptive_avg_pool3d_backwardr|~ s(&k3QR   D0N0N OOr8rcj}td|D]/tjj dkDfd1y)Nr1rc.djddS)Nzc(): Expected grad_output to have non-zero size for non-batch dimensions, but grad_output has sizes rprqr,)rrrsr6rcz3_adaptive_pool_empty_output_check.. s/*--8->->,??OPQsR^`r8)rrrTrer)rrrrs`` @r6rzrz sG   D 1d^    Q ! #   r8c j}tj|dvfdtd|D]. tjj dkD fd0tjt |dk(dd}d}d}jdk(rj d}|dz }j |dz }|\}}jd k(r;|||f}j |} j |tj } | | fS||||f}tj} j |j| } j |tj j| } | | fS) Nr.rc"djS)Nz:adaptive_max_pool2d(): Expected 3D or 4D tensor, but got: r,rsr6rcz*meta_adaptive_max_pool2d.. LU[[MZr8r1rc*djddS)Nzjadaptive_max_pool2d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes rprqr,rrsr6rcz*meta_adaptive_max_pool2d.. %'',{{m3CA3lTr8rcy)NzCadaptive_max_pool2d(): internal error: output_size.size() must be 2r2r2r8r6rcz*meta_adaptive_max_pool2d.. rkr8rr.rr) rrTrerrrrrrMrr$) rrirdimHsizeBsizeDosizeHosizeWrrrrrs ` @r6meta_adaptive_max_pool2dr s| ::D LL Z1d^  JJqMA     LL KAU D E E zzQ 1    JJtax E NFF zzQFF+ ooi(//)5;;/?G|E662 33E: ooi(++-+H//)5;;/?BB'C G|r8cNj}tj|dvfdtdtjjjk(fdt j }jjj|S)Nrc"djS)NzKadaptive_max_pooling2d_backward(): Expected 3D or 4D grad_output, but got: r,rsr6rcz3meta_adaptive_max_pool2d_backward.. s]^i^o^o]pqr8adaptive_max_pool2d_backwardc<djdjSrur)rrsr6rcz3meta_adaptive_max_pool2d_backward.. s!/%++.OP[PaPaObcr8r) rrTrerzr\rMrrrr$)rrrrrs`` r6!meta_adaptive_max_pool2d_backwardr s   D LL q &k3QR LL {(((c //6M ??5;; ' * * * GGr8c j}tj|dvfdtd|D]. tjj dkD fd0tjt |dk(dd}d}d}|dk(rj d}|dz }j |}|\}}}|d k(r||||f} n|||||f} j | } j | tj } | | fS) NrJc"djS)Nz:adaptive_max_pool3d(): Expected 4D or 5D tensor, but got: r,rsr6rcz*meta_adaptive_max_pool3d.. rr8r1rc*djddS)Nzjadaptive_max_pool3d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes rprqr,rsr6rcz*meta_adaptive_max_pool3d.. rr8r.cy)NzCadaptive_max_pool3d(): internal error: output_size.size() must be 3r2r2r8r6rcz*meta_adaptive_max_pool3d.. rkr8r"rr)rrTrerrrrr) rrirdimDrrosizeTrrrrrrs ` @r6meta_adaptive_max_pool3dr s ::D LL Z1d^  JJqMA     LL KAU D E E qy 1    JJt E(FFF qyFFF3 E666: //) $Cooiu{{o;G <r8cPt|d|j|jS)Nadaptive_max_pool3d_backward)rzrr)rrrs r6!meta_adaptive_max_pool3d_backwardr s"&k3QR ??5;; ''r8c>| td|j|S)Nz:cannot repeat_interleave a meta tensor without output_size)r#r)repeatsris r6meta_repeat_interleave_Tensorr s%WXX   [ ))r8c&|jjsJ|jjsJt|jt |j|jt |jt j }|SNrJ)r\rrRr$rrrN)realimagrs r6 meta_complexr sp :: ' '' ' :: ' '' '  +DJJ78 +DJJ786>>F Mr8) fill_valuercd|j||jftjSr})rrrTr)rrrs r6nonzero_staticr s& >>4,EJJ> ??r8c tjtjdtj|j |j fd|j ftj|jS)Ncy)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.r2r2r8r6rcznonzero..' rkr8r1r\r) rT_check_not_implementedrWmeta_nonzero_assume_all_nonzerorrrrrrs r6nonzeror" sf   22 S    txxz" DJJLjj{{  r8c tjtdg}tD]\ftjjtj tj tjtjfvdjtjtjfvrȉj}t|tjjzjkfdtjD]`tjjjzk(fd|j|jdb]|jp|j|tjtjkfdddlm}t%|j&tjkr*jdtjkr*d}d}D]|dk(r d}|dk(rd }nd }|sg}g}tD]*\ |j|j,tD]*\ |j|j,j)||gggtD]\\} @rjj| )jj| Ht%j^fd } j+zz} dd lm} | j1dk(r| S| } t3j4| \}}t%|t%tt|k7r~t3j6| j|}t3j8|}t3j6|t3j:|}| j=| j?|} | S) Ncy)Nz#at least one index must be providedr2r2r8r6rcz#meta_index_Tensor..7 rkr8cy)Nz?tensors used as indices must be long, int, byte or bool tensorsr2r2r8r6rcz#meta_index_Tensor..? rkr8c"djS)N)too many indices for tensor of dimension rrsr6rcz#meta_index_Tensor..F sG {Sr8c NdjddjdzS)NzThe shape of the mask at index z0 does not match the shape of the indexed tensor r,)rrjr;rsr6rcz#meta_index_Tensor..K s<"8 ZPQsSJJN**U_`ade`e_f!hr8r1c<djdtdS)Nrz (got rw)rr)rrsr6rcz#meta_index_Tensor..V s!;DII;fSQX\NZ[\r8rFrTczz}t|j}dgtz|tt|jtz |j ||S)zI This follows restride_src in TensorAdvancedIndexing.cpp r)rrrrr?)rrr after_shape before_shapereplacement_shapes r6 _restride_srcz(meta_index_Tensor.._restride_src so00;>t{{}%KL#PS Q K L!C Oc+6F$FGug..r8guard_or_false) rTrer enumerater\rrr=rrrrrrrselect torch._refs_refsrr$rrr!rrrM3compute_elementwise_output_logical_to_physical_perm apply_permr invert_permr?r)rrrrrefsstatehas_contiguous_subspacertransposed_indicesrrrrrestrided_selfpermrP perm_shaper0rrrrrr;rs`` @@@@@@@r6meta_index_Tensorr5 s LLg MN#%Fg&!5   LL  EIIuzz5::NNY {{uzz5::66--/K"" Ndii/Suzz*8A&& A$**QU*;;h MM'..A"67 8 e$ MM% /!0G LL G  !\ (4(('23G g, "t g, " E# ' A:  aZ}  '#' #!'* 1HAu  A"))%0 1"'* 1HAu} A"))%0 1||D!$!LK#%(2 U = ""4::c?3##DJJsO4 $U[[ 1 2 / ..(99KG HCDdjjla'( #4(NGGWGD! DzT%D *++%%cii6 66zB %%j%2C2CD2IJ nnSXXZ4 Jr8c d} d} d} | dr|j|j} | dr|j|j} | dr|j|} | | | fSNrr1rrr) grad_output_input_weight_bias_sizes_optrrr transposedrr output_maskbackend_grad_inputbackend_grad_weightbackend_grad_biass r6meta_convolution_backwardr sy 1~)33FKKMB1~*44W\\^D1~(22>B  35F GGr8cjd}jd}|j||f}tjj dk(dtjj dk(dtjjdjdk(fdtjjdjdk(fdtj|jd|k(xr|jd|k(d |j |jS) Nr1rr.cyrIr2r2r8r6rczmeta_addbmm.. rkr8cyrKr2r2r8r6rczmeta_addbmm.. rkr8rcPdjddjdS)Nz8batch1 and batch2 must have same number of batches, got rrrrMrNsr6rczmeta_addbmm.. s.J6;;WX>JZZ_`f`k`klm`n_opr8c djddjddjddjdd S)Nz#Incompatible matrix sizes for bmm (r1r;rrrwrrsr6rczmeta_addbmm.. sQ1&++a.1A6;;q>BRS;;q>"!FKKN#31 6r8cy)Nz.self tensor does not match matmul output shaper2r2r8r6rczmeta_addbmm.. rkr8)rrvrTrerr)rrMrNrJrIrYrZs `` r6 meta_addbmmr s ;;q>D ;;q>D ;;d| $D LL"$HI LL"$HI LL A&++a.(p LL A&++a.(  LL ! 51!5@ >>$))+ &&r8c @|j|jSrBr)rrkwargss r6meta_randint_liker s >>$))+ &&r8) grad_scale found_infc n||||||fD])tjttfd+y)Nc dtSNz'exponent must be a tensor list but got rxlsr6rcz#meta__fused_adam_.. =d1gYGr8rTrernr)rgradsexp_avgs exp_avg_sqsmax_exp_avg_sqs state_stepslrbeta1beta2 weight_decayepsamsgradmaximizerrrs @r6meta__fused_adam_r s:&E8[/; O  q$  G  r8c ||||||fD])tjttfd+d}||||||||||fS)Nc dtSrrrsr6rcz"meta__fused_adam..&rr8cR|Dcgc]}tj|c}Scc}wrBr) tensor_listrs r6empty_like_listz)meta__fused_adam..empty_like_list)s!-89  #999s$r)rrrrrrrrrrrrrrrrrs @r6meta__fused_adamrss&E8[/; O  q$  G  : ! $(  r8cjtjjdk(dtjjdk(dtjjtjufdtjjtjufdtjj dj dk(fdj j dj dftj S) Nrcy)Nza must be a 2D tensorr2r2r8r6rczmeta__int_mm..8rkr8cy)Nzb must be a 2D tensorr2r2r8r6rczmeta__int_mm..9rkr8c"djS)Nzexpected self to be int8, got r)r4sr6rczmeta__int_mm..<0 :r8c"djS)Nzexpected mat2 to be int8, got r)r5sr6rczmeta__int_mm..@rr8r1rc djddjddjddjdd S)Nz'Incompatible matrix sizes for _int_mm (rr;r1rrwrr3sr6rczmeta__int_mm..DsM5affQi[!&&)M66!9+Qqvvayk ,r8r)rTrerr\r=rrr>r3s``r6 meta__int_mmr 5s LLA>? LLA>? LL 5::: LL 5::: LL q QVVAY  ;;q 166!9-U[[; AAr8cftjjdk(dtjjtjufdj d}j ddz}j |dz||dzzd|dzftj S) NrcyNzw must be a 2D tensorr2r2r8r6rcz2meta__convert_weight_to_int4pack..Nrkr8c"djSrrrsr6rcz2meta__convert_weight_to_int4pack..Q.qwwi8r8rr1 r)rTrerr\rrrr>r inner_k_tilesrr;s` r6 meta__convert_weight_to_int4packrLs LLA>? LL 5;;8 q A q A A ;; F -"$ %  Q   kk  r8cJtjjdk(dtjjtjufdj d}j d}j ||dzftjS)Nrcyr r2r2r8r6rcz:meta__convert_weight_to_int4pack_for_cpu..brkr8c"djSNzexpected w to be int32, got rrsr6rcz:meta__convert_weight_to_int4pack_for_cpu..err8rr1r)rTrerr\r>rrrrs` r6(meta__convert_weight_to_int4pack_for_cpur`s LLA>? LL 5;;8 q A q A ;; AF kk  r8c.tjjdk(dtjjdk(dtjjtjtj tj fvfdtjjtjufdjjdjddzj S) NrcyNzx must be a 2D tensorr2r2r8r6rcz*meta__weight_int4pack_mm..qrkr8rcy)Nzw must be a 4D tensorr2r2r8r6rcz*meta__weight_int4pack_mm..rrkr8c"djSrrrsr6rcz*meta__weight_int4pack_mm..u5aggY?r8c"djSrrrsr6rcz*meta__weight_int4pack_mm..yrr8rrr rTrerr\rTrUrVr>rrrs`` r6meta__weight_int4pack_mmr#os LLA>? LLA>? LL EMM5==%..AA? LL 5;;8 ;;qvvay!&&)a-qww; ??r8c(tjjdk(dtjjdk(dtjjtjtj tj fvfdtjjtjufdjjdjdjS)Nrcyrr2r2r8r6rcz2meta__weight_int4pack_mm_for_cpu..rkr8cyr r2r2r8r6rcz2meta__weight_int4pack_mm_for_cpu..rkr8c"djSrrrsr6rcz2meta__weight_int4pack_mm_for_cpu..r r8c"djSrrrsr6rcz2meta__weight_int4pack_mm_for_cpu..rr8rr) rTrerr\rTrUrVrrrrs`` r6 meta__weight_int4pack_mm_for_cpur)~ LLA>? LLA>? LL EMM5==%..AA? LL 5;;8 ;;qvvay!&&)177; ;;r8c(tjjdk(dtjjdk(dtjjtjtj tj fvfdtjjtjufdjjdjdjS)Nrcyrr2r2r8r6rcz;_weight_int4pack_mm_with_scales_and_zeros..rkr8cyr r2r2r8r6rcz;_weight_int4pack_mm_with_scales_and_zeros..rkr8c"djSrrrsr6rcz;_weight_int4pack_mm_with_scales_and_zeros..r r8c"djSrrrsr6rcz;_weight_int4pack_mm_with_scales_and_zeros..rr8rrr")r;rrqScaleqZeross`` r6)_weight_int4pack_mm_with_scales_and_zerosr2r*r8r4r5c||zdz |z|zSr0r2r3s r6 kai_roundupr4s UQY1  !!r8c |dk(ry||k(r(d}d}d}dddd fd fd}||||||S|dzd k(rC||zd k(r:d}d}d}dddd fd } fd  d fd  fd|||||||Syyy)Nrrrrc8t||zd}t||S)Nrr4)r;krsrkr_sr_roundedup4s r6kai_k_roundedupz3get_kai_packed_weight_size..kai_k_roundedups#$/rBw#: "1&677r8cX|||}|dzdk(sJd||dzzzzzS)Nrrzk_internal must be evenr2) r;nrr8r9 k_internalr;kai_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_qsu4cxs1s0sY-QB7 "Q1,G.GG,1_23+,))r8c>t|||z}|||||zSrBr7)rr;r=r8r9num_rowsrBs r67kai_get_rhs_packed_size_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0z[get_kai_packed_weight_size..kai_get_rhs_packed_size_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0s6'q"-3O2r2r8rrc|||zdk(sJ| zdk(sJ|zdk(sJt|||z}||||||zSrgr7) rr;r=r8r9blrDkai_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_qsu4c32s1s0spRA~%~//A555//A555&q"-3Q2r2rr8c||zdk(sJ| zdk(sJ|zdk(sJ } ||} ||}|||zz zzSrgr2)r;r=r8r9rGnum_bytes_multiplier_rhsnum_blocks_per_rownum_bytes_per_blockrH#kai_get_bf16_datatype_size_in_bytesrJkai_num_blocks_per_rowr?kai_num_bytes_per_blockrAs r6rIz_get_kai_packed_weight_size..kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0sRA~%~//A555//A555,O+P(%;Ar%B"&=0'#(+==+,()r8cy)Nrr2r2r8r6rPzGget_kai_packed_weight_size..kai_get_bf16_datatype_size_in_bytessr8c6|zdk(sJt|||zSrgr7)r;rGrHs r6rQz:get_kai_packed_weight_size..kai_num_blocks_per_rows)//A555"1b)R//r8c(|zdk(sJ|dz|zS)Nrrr2)rGrMrHs r6rRz;get_kai_packed_weight_size..kai_num_bytes_per_blocks'//A555a#;;;r8r2)n_bitsrK groupsizekai_nrkai_krkai_srrErKrHrPrIrBr;rJrQr?r@rRrAs @@@@@@@@@@@r6get_kai_packed_weight_sizer\s { >FFF$% !+, (!"  8  K1fff ^q Q]a%7FFF$% !!" !" !#     ,  0 <M1fffi u&8 ]r8ctjjtjufdtjj j r||k(r|jtjk(s2||kre|dzdk(r]||zdk(rU|jtjk(r8td|||}jt|tjSj|jz}|||jz }j|tjS)Nc"djSrr)weightssr6rcz2meta__dyn_quant_pack_4bit_weight..s.w}}o>r8rrrr) rTrer\rbackendskleidiai is_availablerXrVr\rrr)r_ scales_zerosr8 block_size in_features out_featurespacked_weight_sizes` r6 meta__dyn_quant_pack_4bit_weightrhs  LL $> ~~++- { "|'9'9U[['H  $R1$j(A-""enn48 |[*   %7!8  LL <+=+=+?? djjl*   /u{{  CCr8cntjjdk(dtjjtjk(xs$jtj k(xrk(fdj d}j||jS)Nrcy)Nzinput must be a 2D tensorr2r2r8r6rcz-meta__dyn_quant_matmul_4bit..1rkr8c.djddS)NzPexpected input to be f32 or bf16 (bf16 requires block_size == in_features), got z with block_size=z and in_features=r)rdreinpsr6rcz-meta__dyn_quant_matmul_4bit..5s*99+.zl:KK= Zr8rr)rTrerr\rTrVrr)rlpacked_weightsrdrerfrs` `` r6meta__dyn_quant_matmul_4bitrn)s LLa!DE LL emm # G II ' EJ+,E   A ==L = ::r8c(tjjdk(dtjjtjtj tj fvfdtjjdk(dtjjtjufdjjdjdjS)Nrcyrr2r2r8r6rcz*meta__weight_int8pack_mm..@rkr8c"djSrrrsr6rcz*meta__weight_int8pack_mm..Cr r8cyr r2r2r8r6rcz*meta__weight_int8pack_mm..Erkr8c"djS)Nzexpected w to be int8, got rrsr6rcz*meta__weight_int8pack_mm..Hs-aggY7r8rr) rTrerr\rTrUrVr=rr)r;rq_scaless`` r6meta__weight_int8pack_mmru>s LLA>? LL EMM5==%..AA? LLA>? LL 5::7 ;;qvvay!&&)177; ;;r8cntjjdk\fdtjjdk\fdtjjdjdk(fdtjt j j fdtjt j j fdtj|dk\d tjd vfd jd }jd }jdd }jdd }ttj||}|j||gj|S) Nrc,djdS)Nz1cdist only supports at least 2D tensors, X1 got: rrx1sr6rcz$meta_cdist_forward..QCBFFH:QOr8c,djdS)Nz1cdist only supports at least 2D tensors, X2 got: rrx2sr6rcz$meta_cdist_forward..Urzr8rcPdjddjdS)Nz4X1 and X2 must have the same number of columns. X1: rz X2: r)ryr}sr6rcz$meta_cdist_forward..Ys,Frwwr{mSXY[Y`Y`acYdXefr8c"djS)Nz3cdist only supports floating-point dtypes, X1 got: rrxsr6rcz$meta_cdist_forward..]EbhhZPr8c"djS)Nz3cdist only supports floating-point dtypes, X2 got: rr|sr6rcz$meta_cdist_forward..arr8rcy)Nz)cdist only supports non-negative p valuesr2r2r8r6rcz$meta_cdist_forward..crkr8rcdS)Nz(possible modes: None, 0, 1, 2, but was: r2) compute_modesr6rcz$meta_cdist_forward..fs:<.Ir8r) rTrerrrMis_float_dtyper\rrbroadcast_shapesextendr) ryr}r;rr1r2 batch_tensor1 batch_tensor2r[s `` ` r6meta_cdist_forwardrMsJ LL A O LL A O LL  rwwr{"f LL RXX&P LL RXX&P LLaLM LL'I B BHHSbMMHHSbMM..}mLMLR! << %%r8c4|jd}|jd}|jd}|jdd}|jdd} ttj|| } | j } | j ||gt j| } |dk(s|dk(s |dk(s| dk(rtj|S| t|jk7r|j| }tj|tjS)Nrrrr) rrrTrcopyrmathprod zeros_likervrr) r?ryr}r;cdistc1rrrrrrtensor1_expand_size batch_products r6meta_cdist_backwardrqs "B "B "BHHSbMMHHSbMM 6 6}m TU.335Bx(II23M Qw"'R1W (:##d288n, YY* +   Be.E.E FFr8c tjjtjtjfvfdtjjtjtjfvfdtjt j jfdjd} |rtj| dk\d| dz} j| jd} }tj|tk(dtjjdk(fdtjjjk(fd fd d fd } td k7r|jjd} jj} |tk(r"j| jd}njd}n| | |}|ttfvs|s!jjd} njd} j| } jd}|tk(rA|rtj|dk\d|dz}j|jd}nj| j}| | | |fS)Nc"djS)Nz(expected indices to be long or int, got r)rsr6rcz$meta_embedding_bag..:7==/Jr8c"djS)Nz(expected offsets to be long or int, got r)rksr6rcz$meta_embedding_bag..rr8c"djS)Nz/expected weight to be floating point type, got r)r6sr6rcz$meta_embedding_bag..sA&,,Pr8rr1cyNz1include_last_offset: numBags should be at least 1r2r2r8r6rcz$meta_embedding_bag..rkr8cy)Nz@embedding_bag: per_sample_weights only supported with mode='sum'r2r2r8r6rcz$meta_embedding_bag..rkr8c$djdS)Nz1expected per_sample_weights to be 1D tensor, got rr)per_sample_weightssr6rcz$meta_embedding_bag..sGHZH_H_G``abr8cNdjdjdS)Nz%expected per_sample_weights.numel() (z$ to be the same as indices.numel() (rwr)rrsr6rcz$meta_embedding_bag..s478J8P8P8R7ST66=mmo5FaIr8cD|||xr|jddk(SNrr1r)r'rpr padding_idxis_fast_path_index_selects r6is_fast_path_index_select_scalez;meta_embedding_bag..is_fast_path_index_select_scales( %c6; ? XELLQROWXDX r8c|jtjk(xs|jtjk(xr1|j ddk(xr|j ddk(xr|dkSr)r\rTrXrVr)r'rrs r6rz5meta_embedding_bag..is_fast_path_index_selectsb YY%++ % @ejj)@  1 "  a A% a  r8c2| ||||S|||SrBr2)r'rprrrrs r6 is_fast_pathz(meta_embedding_bag..is_fast_paths)  23v{S S,S&+F Fr8cpucyrr2r2r8r6rcz$meta_embedding_bag..rkr8)rTrer\rrrMrrrMODE_SUMrrrMODE_MAX MODE_MEANr)r6rrkscale_grad_by_freqrPsparserinclude_last_offsetrnum_bagsrr offset2bagbag_size max_indices fast_path_sumnumBagsrrs``` ` @@r6meta_embedding_bagrs} LL %**eii00J LL %**eii00J LL V\\*P ||AH M G  A    h A 7F% H  V    # #q ( b    $ $ &'--/ 9    G 7u$&&w||A7 $$W\\^4 8 !++Hfkk!nEK!++A.K$V-?U Ix( (  **7<<?;J **1-J$$X.--" 8 " qLO1 !++GV\\!_EK!++HMMO(>{{ zz ;;r8rct||d}tj|j|}t |||}|j ||S)NT)rr)rrMrrrr)rrrr\rr[s r6 meta_nansumrsI(u$OL    T 2D+E4AL ??<|? <.3rkr8rcddS)Nz"Repeats cannot be negative, found rr2)rrepsr6rczmeta_repeat..8s8ZsKr8r) rTrerrrrdrrr)rrnum_new_dimensions padded_sizer target_sizers ` @r6 meta_repeatr/s LL G  "lG$ 3 1H K  W 2++eDJJ.??K8=c'l8KL1;q>GAJ.LKL >>+ &&Ms1Cc|SrBr2rs r6 meta_zero_rCrdr8czt|tjr t|j|j|SrB)rnrTr rgrrrs r6meta_binop_inplacerHs)%& EKK8 Kr8cd}d}d}||r||r td||r||s tdt|tjr t |j |j |S)a* Some checks for inplace ops. Checks for promotion rules for some dtypes. int.add/sub_(float) and bool.add/sub_(others) are rejected. Promoting in these in-place operations would require reallocating and copying over elements, hence not allowed. Checks for alpha param. ct|trtj|jSt|t SrB)rnrrMr1r\rrqs r6 is_integericz.meta_binop_inplace_alpha..is_integericks. c: &))#))4 4c7+ +r8ct|trtj|jSt|t SrB)rnrrMrr\rrs r6 is_floaticz,meta_binop_inplace_alpha..is_floaticqs. c: &'' 2 2c9- -r8ct|trtj|jSt|t SrB)rnrrMis_boolean_dtyper\rrs r6 is_booleanicz.meta_binop_inplace_alpha..is_booleanicws. c: &))#))4 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#rnrTr rgr)rrrIrrrs r6meta_binop_inplace_alpharYsz$, . -Dj/ k  D,u"5 m  %& EKK8 Kr8c:t||tjSrrRrrNrrrIs r6meta_binop_alphars  e$C$K$K r8c 8t|tjSrr)rrs r6 meta_roundrs  <DD r8cltjtjjfdt tj r8tjtjjfdytjt tfdy)Nc&djS)Nz7: Expected input tensor to have an integral dtype. Got r)rrsr6rcz#shift_dtype_check..s7)RSWS]S]R^_r8c&djS)Nz6: Expected shift value to have an integral dtype. Got rrrsr6rcz#shift_dtype_check..swiUVYV_V_U`ar8cdS)Nz): Expected shift value to be an int. Got r2rsr6rcz#shift_dtype_check..swiHNr8)rTrerMr1r\rnr r)rrrs```r6shift_dtype_checkrsp LL tzz*_#u||$  " "399 - a  sG $ N r8cTtd||t||tjS)NrshiftrrrRrrNrs r6 meta_rshiftsr)he,  e$C$K$K r8cTtd||t||tjS)Nlshiftrrrs r6 meta_lshiftsrrr8c8|j|jSrBrrs r6 meta_zerors >>$** %%r8c|SrBr2rrs r6 meta_fill_rrdr8c,tj|SrBrrs r6 meta_fillr   D !!r8c|SrBr2rs r6 meta_relu_rrdr8c:t||tjSrrrs r6meta__add_relurs  e$C$K$K r8c,tj|SrBrrnoiselowerrrrs r6meta_rrelu_with_noisers   D !!r8cVtj|tj|fSrBrrs r6 meta_rrelu_with_noise_functionalrs%   D !5#3#3E#: ::r8c|SrBr2)rrrrrs r6meta_rrelu_with_noise_rs  Kr8c,tj|SrBrrrr accumulates r6meta_index_putr rr8cFt|j|j|SrBrgr)rrxvalues r6meta_masked_fill_rsDJJ 3 Kr8c|j|jjtj|}|Sr)rrr$rMr)rrxrp masked_scales r6meta__masked_scalers<>>$))+.1111$72L r8ctj|jtjtjfvdtjjjk(fdS)Ncy)NzMask must be bool or uint8r2r2r8r6rcz&meta_masked_scatter_.. rkr8c<djdjS)NzEmasked_scatter: expected self and source to have same dtypes but got rr)rr^sr6rcz&meta_masked_scatter_..s"**U6<<.:r8)rTrer\rr)rrxr^s` `r6meta_masked_scatter_rsU LL uzz5;;//1U LL fll" : Kr8ct||\}}tj|tj}t |||Sr)r$rTrrr)rrxr^rs r6meta_masked_scatterrs;"$-JD$   d%2I2I JF f 55r8c$|j|SrBry)rrxr:s r6meta_masked_scatter_backwardrs >>%  r8c|SrBr2r s r6meta_index_put_r!rdr8c  ddlm}m}tj|j dk(dtj|j dk(d|j }|j |d |d|d} d} || ftj|| d | d fd|r|jtjk(xs|jtjk(xr|tjk(} tj||jk(xs| d |jj|} n|j} |sYWtjj dk(d tj|j fd | S) Nr)sym_andrGr.cyrIr2r2r8r6rcz)common_meta_baddbmm_bmm..)rkr8cyrKr2r2r8r6rcz)common_meta_baddbmm_bmm..*rkr8rr1c .ddddddd SrPr2rRsr6rcz)common_meta_baddbmm_bmm..7s5RSURV  l<?*;2l1o=NbRr8cy)Nzfout_dtype only supported for torch.float32 output with float16/bfloat16 inputs or same as input dtypesr2r2r8r6rcz)common_meta_baddbmm_bmm..@rkr8cy)Nzself must be a 3D tensorr2r2r8r6rcz)common_meta_baddbmm_bmm..Hrkr8c0ddjS)Nz*Expected an input tensor shape with shape z but got shape: r)ri self_baddbmmsr6rcz)common_meta_baddbmm_bmm..Ks!@ M]^j^o^o^q]rsr8) r!rrGrTrerrr\rUrVrTrr$)rMrNis_bmmr&r:rrGr\res_rowsres_colssupported_out_dtyperrSrTrUris ` @@@@r6common_meta_baddbmm_bmmr+&sE LL"$HI LL"$HI;;=L;;=L aB#AAHAHx*K LL|A+VLOEU-VW R  LLEMM ) KV\\U^^-K)5==(    % <)< | !!+.11)<!!+. l. \%%'1,.PQ <$$& 4 s Mr8ct||dS)NTr+)rrEs r6meta_bmmr.Qs "4t 44r8c t||d|S)NT)r:r-)rrEr:s r6meta_bmm_dtyper0Vs "4ty IIr8ch||z}||z}|dk7r"t|dkt|dkk7r|dz}|Sr)r)r;r<qrTs r6div_rtnr3[sB QA AA Av4A;$q1u+- Q Hr8ct||z|z||dz zz dz |r|dz ndz|dz}|r|dz |z||zk\r|dz}|Sr)r3) inputSize kernelSizerrrrr outputSizes r6pooling_output_shape_pad_lrr8es    *q.) *   'vzA  /     Nf $ E(9 9 !OJ r8c tj|dk7dtjdk\fdtjdz zdzdzkfdt|||S)Nrcy)Nzstride should not be zeror2r2r8r6rcz&pooling_output_shape..rkr8cdS)Nz'pad must be non-negative, but got pad: r2padsr6rcz&pooling_output_shape..s%LSE#Rr8r1rcdddS)NzApad should be at most half of effective kernel size, but got pad=z, kernel_size=z and dilation=r2)rr6r=sr6rcz&pooling_output_shape..s'OPSuU%,nXJ @r8)rTrer8)r5r6r=rrrs `` ` r6r'r'ss LL1AB LLRS LL a8+a/A55  ':sC9 r8c \ j} tjdkDxrdkDfdtjdkDxrdkDfdtjdkDxrdkDfdjddk7xrjddk7}|tjk(r7tj|dk(xr|xrjddk7fd nWtj|dk(xrjddk7xr|xs|dk(xr|xrjddk7fd tjdzk\xrdzk\fd tj dk\xr dk\  fd y) NrcddS)Nz5kernel size should be greater than zero, but got kH: , kW: r2)r+r,sr6rcz$pool2d_shape_check..sGt6RTQUVr8cddS)Nz0stride should be greater than zero, but got dH: , dW: r2)r-r.sr6rcz$pool2d_shape_check..sB2$fRDQr8cddS)Nz9dilation should be greater than zero, but got dilationH: , dilationW: r2) dilationH dilationWsr6rcz$pool2d_shape_check..sKI;Vcdmcnor8r1rrr.c*djS)NzExpected 4D (batch mode) tensor expected for input with channels_last layout with optional 0 dim batch size for input, but got: rrsr6rcz$pool2d_shape_check..sCCH::<.Rr8c*djS)NzYExpected 3D or 4D (batch mode) tensor with optional 0 dim batch size for input, but got: rrsr6rcz$pool2d_shape_check..sopupzpzp|o}~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,r/r0sr6rcz$pool2d_shape_check..s&ygbT>r8c .ddddddd SNzGiven input size: (r;z). Calculated output size: (z). Output size is too smallr2)r1r2rr7rrsr6rcz$pool2d_shape_check..s:%k]!K=*N$$0><.+O##r8)rrTrerr)rr+r,r-r.r/r0rFrGrr1r2rrrr valid_dimsr7s`````````````` @r6r(r(s" 99;DL LL Q26V LL Q26Q LLA ')a-o A!#: 1 (:J+++ AI ;* ;A!); R  QY <5::a=A- <* A ?j?UZZ]a-? ~  LL a4+B!GtO >  LLq.\Q. # #r8rSrPr+r,rQr-r.pTpHpW dilationTrFrGrTrUrVrWrXrYrNc J j}tjdkDxr dkDxrdkDfdtjdkDxr dkDxrdkDfdtj dkDxr  dkDxr dkD fdtj|dvfdt|D]:|dk(rdk(rtjj dkDfd<|r/tjk\xr k\xrk\fd tjd z k\xrd z k\xrd z k\  fd tjd k\xr d k\xrd k\fd y)NrcdddS)Nz5kernel size should be greater than zero, but got kT: z, kH: rAr2)r+rPr,sr6rcz$pool3d_shape_check..s#$fRDrd ,r8cdddS)Nz0stride should be greater than zero, but got dT: z, dH: rCr2)r-rQr.sr6rcz$pool3d_shape_check..s >rd&FSURV Wr8cdddS)Nz9dilation should be greater than zero, but got dilationT: z , dilationH: rEr2)rFrQrGsr6rcz$pool3d_shape_check..s$#M)M) Vr8rJc&djS)Nz/: Expected 4D or 5D tensor for input, but got: r,)rrsr6rcz$pool3d_shape_check..s7)J5;;-Xr8r"cLdjdjdS)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)rrrsr6rcz$pool3d_shape_check..s3)--2[[M+EJJqM?:KMr8c .ddddddd S)Nzinput image (T: r*rz ) smaller than kernel size (kT:  kH:  kW: rwr2)rUrTrVr+rPr,sr6rcz$pool3d_shape_check..s9"5'gYd6(C$$&4uRDbT<r8rc ,dddddd S)NzHpad should be smaller than or equal to half of kernel size, but got kT: rZrYz padT: z padW: z padH: r2)r+rPr,rOrNrPsr6rcz$pool3d_shape_check..s6$eB4uRDt72$gbT Kr8r1c:dddddddddSrLr2)rUrTrVrSrXrWrYsr6rcz$pool3d_shape_check..!sI!'!E7!G9AfXF((/y%'!F8L' (r8)rrTrerr)rrSrPr+r,rQr-r.rNrOrPrQrFrGrTrUrVrWrXrYrrNrrs````````````````````` @r6rOrOs0 ::D LL Q$26$b1f  LL Q$26$b1f   LLA 9)a-9IM  LL X 4[   19a  JJqMA     RK :GrM :fl    LL Q" 6a26"q&B,   LL  3v{3w!|  r8cj|j}t|||||||| | | | | |||||||||t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |y)Nrr.rr1rrOr)rrrrSrPr+r,rQr-r.rNrOrPrQrFrGrTrUrVrWrXrYrrs r6max_pool3d_backward_shape_checkr_)s2 ::D            +0;dQh8;dQh6;dQh8;dQh77D$(G47D$(E27D$(G47D$(F3r8c|j}t||||||||| | | ddd| | |||||dt|||dz |t|||dz |t|||dz |t|||dz |y)Nr1Trr.rr^)rrrSrPr+r,rQr-r.rNrOrPrTrUrVrWrXrYrrs r6rarags* ::D                -2;dQh8;dQh6;dQh8;dQh7r8cd}|d|\}}tjt|dvdt|dk(r||} } n |d|\} } |d|\} } |d|\} }|jd }|jd }|jd }t j |}|tj k(r)tj|jd k(d nR|tjk(r(tj|jdvdntjddt||| | | |}t||| | ||}t|||| | | | | ||||||||||fS)Nctjt|dvfd|d}t|dk(r|n|d}||fS)NrcddS)Nz max_pool2d: rr2rsr6rczEmax_pool2d_checks_and_compute_shape..unpack..rr8rr1rrs` r6rz3max_pool2d_checks_and_compute_shape..unpackrr8rrcy)NzOmax_pool2d: stride must either be omitted, a single int, or a tuple of two intsr2r2r8r6rcz5max_pool2d_checks_and_compute_shape..rkr8rrrrr#rrrcy)NzMnon-empty 4D (batch mode) tensor expected for input with channels_last layoutr2r2r8r6rcz5max_pool2d_checks_and_compute_shape..rkr8rcy)Nz9non-empty 3D or 4D (batch mode) tensor expected for inputr2r2r8r6rcz5max_pool2d_checks_and_compute_shape..rkr8Fcy)NzAUnsupported memory format. Supports only ChannelsLast, Contiguousr2r2r8r6rcz5max_pool2d_checks_and_compute_shape..rkr8) rTrerrrMrrrrr'r()rrrrrrrr+r,r-r.r/r0rFrGrr1r2rrrs r6rrsM; /FB LL F y a 6{aRB&)B 7+JD$!*h7Iy**R.K**R.KBJ//6M+++ IIK1  c  %11 1 IIK6 ! O   W ( Rr9iXL&z2tRIVK       $  k 11r8c| t|||||\} tjjjk(fd| j fd} | | |t j } tjjjj| S)Nc<djdjS)NzExpected dtype z for `gradOutput` but got dtype rr{sr6rcz7meta_max_pool2d_with_indices_backward..s /$**-MkN_N_M`ar8clt|dz t|dz t|dz y)Nr.rr1)r)rr7rrrs r6_check_dim_sizez>meta_max_pool2d_with_indices_backward.._check_dim_sizes9q$q,7q$q,7q$q+6r8rr) rrTrer\rrMrrrr)rrrrrrrrrrkrr7rrrs`` @@@@r6%meta_max_pool2d_with_indices_backwardrls , k67Hi    LL k'''a L 99D7 K G//5M ;; jj{{#  r8ct||||||\}}}|jdk(r|jdnd} tj|} |jdk(r|||g} n| |||g} t j | |j|j| t j | t j|j| fSr) rrrrMrrTrr\rrrs r6meta_max_pool2d_with_indicesrns , {FGXy   %yy{a/UZZ^QF//6M yy{a\;7 \;? ++<<'   ++<<'    r8ct tjjdvfdj}t|dz |D]. tjj dkD fd0tjt dk(dtjt |dk(dj d }j d  j d  |d k(rj d}nd }tjj j k(dtjjdk(fdj d}j d }j d tj||k\dtj||k(dtj dk( fdtj|ddzd z k fdtj|d d zd z k fdjd k(r |||d|d g} n ||d|d g} tj| j jtj| tjjfS)Nrc"djS)Nz:fractional_max_pool2d: Expected 3D or 4D tensor, but got: rrsr6rcz,meta_fractional_max_pool2d..@sLTYYKXr8r.rc2djddS)Nz_fractional_max_pool2d: Expected input to have non-zero size for non-batch dimensions, but got rpz emptyr)rrsr6rcz,meta_fractional_max_pool2d..Gs'77;yy{mCSTUSVV\^r8rcy)NzNfractional_max_pool2d: kernel_size musteither be a single int or tuple of Intsr2r2r8r6rcz,meta_fractional_max_pool2d..Nrkr8cy)NzOfractional_max_pool2d: output_size must either be a single int or tuple of Intsr2r2r8r6rcz,meta_fractional_max_pool2d..Srkr8r#rrrr1cy)Nz6Expect _random_samples to have the same dtype as inputr2r2r8r6rcz,meta_fractional_max_pool2d..arkr8c"djS)Nz1Expect _random samples to have 3 dimensions got, r)random_samplessr6rcz,meta_fractional_max_pool2d..esCNDWDWCXYr8cy)Nz=Expect _random_samples.size(0) no less then input batch size.r2r2r8r6rcz,meta_fractional_max_pool2d..mrkr8cy)Nz.qrkr8cddS)Nz/Expect _random_samples.size(2) equals to 2 got .r2)rsr6rcz,meta_fractional_max_pool2d..ss#RSTRUUV!Wr8cdddS)Nz%fractional_max_pool2d: kernel height rz' is too large relative to input height r2) input_heightrsr6rcz,meta_fractional_max_pool2d..ws7 A7GGno{n|}r8cdddS)Nz$fractional_max_pool2d: kernel width r1z& is too large relative to input width r2) input_widthrsr6rcz,meta_fractional_max_pool2d..{s6{1~6FFlmxlyzr8r) rTrerrrrr\rrrr) rrrirvrinput_channels input_batchrcrrr|r~s `` ` @@@r6meta_fractional_max_pool2dr<ss LL VX 99D 4!8T "  IIaL1  ^   LL KA 2  LL KA 2 YYr]N99R=L))B-K qyiil   LL n***H LLq Y AAAAAA LL [O LL ^N LLaWX LLAQ'!+|;} LLAQ'!+{:z  xxzQ^[^[^L A A?  **;;  ++;;  r8c|tjt|dvd|d}t|dk(r|n|d}t|dk(r|n|d}tj| xs t|dvd|s|n|d} |s|nt|dk(r| n|d} |s|nt|dk(r| n|d} tjt|dvd|d} t|dk(r| n|d} t|dk(r| n|d}tjt|dvd|d}t|dk(r|n|d}t|dk(r|n|d}tj|jd vd |jd k(r|j d nd}|j d }|j d}|j d}|j d}t ||| | ||}t ||| | ||}t |||| ||}t |||||| | | | | ||||||||||d|jd k(xr&tj|tjk(}|jdk(rK|jd}|j xr |jtj}||||f}n|||||f}|j|}|j|tj}|r@|jtj}|jtj}||fS)NrBcyNzMmax_pool3d: kernel_size must either be a single int, or a tuple of three intsr2r2r8r6rcz.meta_max_pool3d_with_indices..rkr8rr1rcyNzQmax_pool3d: stride must either be omitted, a single int, or a tuple of three intsr2r2r8r6rcz.meta_max_pool3d_with_indices..rkr8cyNzImax_pool3d: padding must either be a single int, or a tuple of three intsr2r2r8r6rcz.meta_max_pool3d_with_indices..rkr8cyNzJmax_pool3d: dilation must be either a single int, or a tuple of three intsr2r2r8r6rcz.meta_max_pool3d_with_indices..rkr8rJcyrLr2r2r8r6rcz.meta_max_pool3d_with_indices..rkr8r"rr#rrzmax_pool3d_with_indices()rrr)rTrerrrr'rOrMrrrrXrrr$)rrrrrrrPr+r,rQr-r.rNrOrPrQrFrGrrSrTrUrVrWrXrYrinput_channels_last_checkrrrs r6meta_max_pool3d_with_indicesrsJ LL KF"_ QB;1$+a.B;1$+a.B LL  +c&kV+cvayBc&kQ&6F1IBc&kQ&6F1IB LL G [ B7|q gajB7|q gajB LL H \ I ]a/ Xa[I ]a/ Xa[I LL fK %zzQUZZ^AFjjnG JJrNEjjnG ZZ^F BIy IE"7BB 9MG !&"b"i KF            #+2  aXE77>%BXBXX zzQ$)OOA$6!)779 9 '55006   eWf5 WeWf= //) $Cooiu{{o;Gff5#9#9f:**5+A+A*B <r8ctjt|dvd|d}t|dk(r|n|d} t|dk(r|n|d} tj| xs t|dvd|s|n|d} |s| nt|dk(r| n|d} |s| nt|dk(r| n|d} tjt|dvd|d}t|dk(r|n|d}t|dk(r|n|d}tjt|dvd|d}t|dk(r|n|d}t|dk(r|n|d}tj|jd vd |j d }|j d }|j d }|j d}|j d }|j d }|j d}t |||||| | | | | ||||||||||||d|jdk(xr&t j|tjk(}|jdk(rD|jd}|j xr |jtj}|j|j}|r |jtj}|S)NrBcyrr2r2r8r6rcz7meta_max_pool3d_with_indices_backward..rkr8rr1rcyrr2r2r8r6rcz7meta_max_pool3d_with_indices_backward..rkr8cyrr2r2r8r6rcz7meta_max_pool3d_with_indices_backward..rkr8cyrr2r2r8r6rcz7meta_max_pool3d_with_indices_backward.. rkr8rJcyrLr2r2r8r6rcz7meta_max_pool3d_with_indices_backward..(rkr8rr#rrz"max_pool3d_with_indices_backward()r"rr)rTrerrrr_rMrrrrXrrr$)rrrrrrrrrPr+r,rQr-r.rNrOrPrQrFrGrSrTrUrVrWrXrYrrrs r6%meta_max_pool3d_with_indices_backwardrs LL KF"_ QB;1$+a.B;1$+a.B LL  +c&kV+cvayBc&kQ&6F1IBc&kQ&6F1IB LL G [ B7|q gajB7|q gajB LL H \ I ]a/ Xa[I ]a/ Xa[I LL fK jjnG JJrNEjjnG ZZ^F   R Er"G   b !F#            ,/6  aXE77>%BXBXX zzQ$)OOA$6!)779 9 '55006   -J]]1G1G]H r8gridcztjjjk(fdtjjtjk(xrjtjk(fdtjj dj dk(fdtjj dj dz k(fdtdj D],tjj dkDfd.y) Nc<djdjS)NzNgrid_sampler(): expected input and grid to be on same device, but input is on z and grid is on rrrsr6rcz+check_grid_sampler_common..ds'\\N"24;;- Ar8c<djdjS)NzTgrid_sampler(): expected input and grid to have torch.strided layout, but input has z and grid has )rrsr6rcz+check_grid_sampler_common..ks&nT[[M Cr8rc<djdjS)NzZgrid_sampler(): expected grid and input to have same batch size, but got input with sizes  and grid with sizes r,rsr6rcz+check_grid_sampler_common..rs' % },A$** Or8rrcBdjdz djS)Nz+grid_sampler(): expected grid to have size rz, in last dimension, but got grid with sizes )rrrsr6rcz+check_grid_sampler_common..ys,9%**q.9IJ226** ?r8c*djddS)NzYgrid_sampler(): expected input to have non-empty spatial dimensions, but input has sizes rprqr,rsr6rcz+check_grid_sampler_common..rr8)rTrerrrrrr)rrrs``@r6check_grid_sampler_commonras LL  #  LL  %F$++*F  LL A$**Q-'  LL 2%**q.( 1ejj !  KKNQ    r8ceZdZdZdZdZy)GridSamplerInterpolationrr1rN)ry __module__ __qualname__BILINEARNEARESTBICUBICr2r8r6rrsHGGr8rinterpolation_modectjjdk(xrjjk(fdtjjdk(xr|tjj k( dy)Nr"c<djdjS)Nzdgrid_sampler(): expected 5D input and grid with same number of dimensions, but got input with sizes rr,rsr6rcz'check_grid_sampler_3d..s&449KK=#DJJ< 1r8cy)Nz.rkr8)rTrerrrr)rrrs`` r6check_grid_sampler_3drsp LL a3EJJ$))3  LL JJ!O M"&>&F&F&L&LL  O r8c|d}|r&tj|tj}nd}tj|tj} || fSNrr)rTrrr rrrr padding_mode align_cornersrinput_requires_gradr grad_grids r6grid_sampler_2d_backward_metarsQ&a.%%e5;R;RS    U5L5LMI  ""r8ct||t||||jd}|jd}|jd}|jd}|jd} |j||||| fS)Nrr1rr.)rrrr) rrrrrrCout_Dout_Hout_Ws r6grid_sampler_3drsveT*%'9: AA AA JJqME JJqME JJqME ??Aq%6 77r8rct||t||||d}|r&tj|tj}nd}tj |tj} || fSr)rrrTrrrrs r6grid_sampler_3d_backwardrsmeT*%'9:%a.%% !?!?    U5S5STI y  r8c|jd}|stj|}||d<tj|g|i|S)Nr\)r[rM get_dtyperTr)rrrOrr\s r6fullrsC JJw E  +F7O ;;t -d -f --r8cN|tjk(rtj|dudtjd| |jn||| |j n||}|j r>|j|j|j|jn/|j|j|jd|jd|Stjj||||||}|j!d|S)Ncy)Nz9memory format option is only supported by strided tensorsr2r2r8r6rczzeros_like..rkr8rrTr)rT sparse_coorerr\r is_sparsesparse_resize_and_clear_r sparse_dim dense_dimr _coalesced_r+rr&fill_)rr\rrrrrs r6rrs!!! T ! O kk % $**5"(.4;;f!   >>  ( ( T__.0@   ( (dhhj! D  // ! ! # " CIIaL Jr8rc|tj}|tj}|tj}tj|||||Sr rTrget_default_devicerrrr\rrrrs r6 meta_onesrT }'') ~))+ ~ ;; E&J r8c|tj}|tj}|tj}tj|||||Sr rrs r6 meta_zerosr-rr8c,tj|SrBrMclone_preserve_strides)rr'rrs r6meta_select_scatterrC  ' ' --r8c,tj|SrBr)rr'rr{rzsteps r6meta_slice_scatterrHrr8 dim_post_expr wrap_scalarcv|dkr|sJd}| }|dz }||ks||kDrJd|d|d|d|dkr||z }|S)Nrr1zdim z out of bounds (rvrwr2)rrrrrs r6rrNsm{ .C ! Cc S3YR4u4DSEC5PQ)RR ' Qw } Jr8cJ|jdk(rdS|j|Srr )rrs r6ensure_nonempty_sizerZs!1 1.!''#,.r8c:tjd}tjd}tj||k(dt |D];k7s tjt t kfd=y)Nr1cy)NzDIndex tensor must have the same number of dimensions as input tensorr2r2r8r6rcz$gather_shape_check..drkr8cNddjdjdzS)Nz!Size does not match at dimension z expected index  to be no larger than self  apart from dimension r,)rrrrsr6rcz$gather_shape_check..js7;A3>Nu{{m\/ |;QRUQVWXr8)rrrTrerr)rrr self_dims index_dimsrs``` @r6gather_shape_checkr_sDHHJ"IUYY[!$J LLZV9  8 LL$UA.2FtQ2OOX r8cpddlm}t||j}|j dk(}|s`t j jt jk(xsjt jk(fdt|||jjS)Nrrc"djS)Nz8gather(): Expected dtype int32/int64 for index, but got rrsr6rczmeta_gather..xsNu{{m\r8) r!rrrrrTrer\rrrrr)rrr sparse_gradr wrapped_dimis_index_emptys ` r6 meta_gatherrosD dhhj1K#EKKMQ$67N  KK5:: % A )A \  4e4 >>%++ &&r8c|r6|dk(ry|dk(ry|dk(ry|dk(ry|d k(ry tjd d y|d k(ry|dk(rytjd dy)Nr REDUCE_ADDrREDUCE_MULTIPLYmean REDUCE_MEANamaxREDUCE_MAXIMUMaminREDUCE_MINIMUMFcy)Nz=reduce argument must be either sum, prod, mean, amax or amin.r2r2r8r6rcz#get_operator_enum..rkr8addmultiplycy)Nz/reduce argument must be either add or multiply.r2r2r8r6rcz#get_operator_enum..rkr8rV)reduce_use_new_optionss r6get_operator_enumrs{ e   $    #  #  S   e   "$ UUVr8cPddlm}||jdk7rStj|j tj k(xs|j tjk(fd|1tj|j |j k(fdyy)Nr)rFcdS)Nz((): Expected dtype int32/int64 for indexr2 method_namesr6rcz,scatter_gather_dtype_check..s{m#KLr8cdS)Nz0(): Expected self.dtype to be equal to src.dtyper2rsr6rcz,scatter_gather_dtype_check..s{m#STr8)r!rFrrTrer\rr)rrrsrc_optrFs` r6scatter_gather_dtype_checkrsvCU[[]a'( KK5:: % A )A L  JJ'-- ' T r8ct|dSr0)rrs r6ensure_nonempty_dimrs sA;r8cddlm}|jdk(rytjt j t j k(dd}t j }t|D]'}t|}|k(r|t|kDs%d}n|s1/t|D]!}t|}|t|kDsd}nftjt j t j k(dtj| fdytj| fdy) NrrcyNzCIndex tensor must have the same number of dimensions as self tensorr2r2r8r6rcz%scatter_shape_check..rkr8FTcyrr2r2r8r6rcz%scatter_shape_check..rkr8cbdjdjddjzS)NExpected index rrz and to be no larger than src r,)rrrrsr6rcz%scatter_shape_check..s8oekk]2Mdjj\Z&se+I'--YZr8cHdjdjdzS)Nr rrr,)rrrsr6rcz%scatter_shape_check..s,oekk]2Mdjj\Z&se,-r8) r!rrrTrerrrr) rrrrris_wrong_shaperr index_d_sizes ```` r6scatter_shape_checkrsPDekkmq() LLDHHJ'+>uyy{+KKU N#DHHJ/I9 +E15 8  .tQ7 7!N   g1y! A/q9L27A>>!%     +/B599;/O O Y     Z    - r8ct||j}td|||t||||| t ||yy)Nscatter)rrrrr)rrrr'rrrs r6scatter_meta_implrsE dhhj1Ky$s;k5#6'?3r8cVt||||d|j|jSNrrrrrrrr's r6meta_scatter_addrs%dCU3 >>$** %%r8c$t||||d|Srrrs r6meta_scatter_add_rsdCU3 Kr8ct|tjr|nd}t||||||j |j SrB)rnrTr rrrrrr src_or_valuer_r's r6 meta_scatterrs;%\5<<@,dCdCV4 >>$** %%r8c`t|tjr|nd}t||||||SrB)rnrTr rrs r6 meta_scatter_rs-%\5<<@,dCdCV4 Kr8queryrr dropout_p is_causalreturn_debug_maskrpc .|jd}|jd}|jd} |jd} |jd} tj|} tj||| ftj|j } |ra| dkDrdnd}t j| |z }| dkrd}n| dkrd}tj||| |f|j|j }n,tjd|j|j }tjjrtjjst|d k(rMtjd tjd }tjd tjd }nLtjdtjd }tjd tjd }| | dd| | |||f S) Nrr1rr.r@rr2r~)rrTrrrXrrceilr\versionhiprrbrrr&)r rrr!r"r#rpr num_headsmax_seqlen_batch_qhead_dimmax_seqlen_batch_k attention logsumexp blocksize_c max_seqlen_k debug_maskseedoffsets r6(meta__scaled_dot_product_flash_attentionr6sAJ 1 IAzz!}H!  'I Y 23kk||I %]c yy!3k!AB  $L 3 &L[[ $6 E++<< [[%++ellK  }}UZZ446+e:LPU:U{{2UZZ?Ruzz&A{{Aell6BRu||FC      r8 res_shape.ctj|k(rtj}|St gdfdd}|Dcgc]}|| }}t t |Dcgc]}|j|}}tj|jjj|}|Scc}wcc}w)N)rr1rr.c*j|SrBr)idxr s r6rcz,alloc_with_matching_layout..Zs%,,.*=r8Trr) rdrrTrsortedrrrrr\rr)r r7r dim_orderr:permuted_shaper final_permutes` r6alloc_with_matching_layoutr?Rs U[[Y&u% J =t 5>>S)C.>>5:3y>5JK+K Kkk %++ell '-   J ?Ks  C,C attn_biascompute_log_sumexpc |jd} |jd} |jd} |jd} |jd} | | | | f}t||}tj| | | dftj|j }tjdtj d}tjdtj d}||dd| | ||df SNrr1rrrr2r~rr?rTrrXrr)r rrr@rAr!r"r#rprrS_QS_KVD_Vr7r logsum_expr4r5s r6(meta__scaled_dot_product_cudnn_attentionrIes  1 A 1 A **Q-C 88A;D **R.CAsC I $UI 6C AsAkk||J ;;rF ;D [[5::f =F         r8c |jd}|jd} |jd} |jd} |jd} || | | f} t|| }tj|| | ftj|j }tjdtj d}tjdtj d}||dd| | ||df SrCrD)r rrr@r!r"r#rprH_QrErFrGr7rrHr4r5s r65meta__scaled_dot_product_fused_attention_overrideablerLs  1 A **Q-C **Q-C 88A;D **R.CCc"I $UI 6C C kk||J ;;rF ;D [[5::f =F         r8rrr0 cum_seq_q cum_seq_kmax_qmax_k philox_seed philox_offsetctj|}tj|}tj|}|||fSrBr)rrr rrrr0rMrNrOrPr!r"rQrRrpgrad_qgrad_krs r6'meta__scaled_dot_product_flash_backwardrVsA,  e $F   c "F   e $F 66 !!r8 attn_maskc |jd}|jd}|jd} tj|} tj|| |ftj|j j dd} | | fS)Nrr1rr)rrTrrrXrr) r rrr!r"rWrprr+r,r/r0s r60meta__scaled_dot_product_flash_attention_for_cpurYsAJ 1 IA  'I    kk||i1o  r8c ntj|jd|j|j} tj|jd|j|j} tj|jd|j|j} | | | fS)Nrrr1r.r)rTempty_permutedrr\r) rrr rrrr0r!r"rWrprTrUrs r69meta__scaled_dot_product_flash_attention_for_cpu_backwardr]s& ! ! kk|| F  ! !  iizz F  ! ! kk|| F 66 !!r8 dropout_maskcd}|\||\} } ||\} } j\| j\} } } fd} fd}dk\s | kr dk\r|S| S)Ncl|jdk(r|jddfS|jdkDrxd}t|jdz D]}||j|z}|j ||j d|j d|j ddfS|d fS) Nr.rTrr1r#rrF)rrrrviewr)r;rrs r6 ensure_4dzBmeta__scaled_dot_product_attention_math_for_mps..ensure_4d.s 557a<;;q>4' ' UUWq[J1557Q;' )aggaj(  )66*affRj!&&*affRjI4O Oe8Or8cXjj} r|j}jf} rajdk(r|j d}||fSt jdd|jddz}|j |}||fS)Nr.rr#r1r)rrview_asrsqueezerra) rattnrrmax_seq_lengthnum_headq_q_sizer  unsqueezeds r6sdpa_vector_fast_mpszMmeta__scaled_dot_product_attention_math_for_mps..sdpa_vector_fast_mps@sll288$ ++e$C||Z6>JK yy{a||ADyU[["-.Aa@yy'Dyr8crd}jj}j|f}||fS)Nrr)blocksrr)r head_sizerhrirjs r6sdpa_vector_2pass_mpszNmeta__scaled_dot_product_attention_math_for_mps..sdpa_vector_2pass_mpsNs>ll288$||Z669$UV L  r8iir,)r rrrWr!r"r^rprbk_rPv_k_sizerlrprrorgrhrirjrks` @@@@@@@r6/meta__scaled_dot_product_attention_math_for_mpsrt#s u%NB cNEB e EB.0hh+J&)#%88 Av~q  !! $FVO$8N$&&#%%r8c@|jdd}|jdd}|jdd}|jd}|jd} |jd} |jd} tj|| | | |j|j } tj jr&tjjr |r| nd} n|rtj| dz dznd} tj|| | ftj|j }| jdd} tjdtjd }tjdtjd }| |||fS) Nr1rrrrrrr2r~)rrrTrr\rr)r*rrbrr(rXr)r rrr@rAr!r"rprrr+Kvr logsumexp_dimrHr4r5s r6,meta__scaled_dot_product_efficient_attentionrxZsD OOAq !E --1 C OOAq !E 1 A 1 A 2I BB ++aIrU\\ RC }}UZZ446 0Q 2D !b&)B.!  I}%kk||J --1 C ;;rF ;D [[5::f =F  D& ((r8grad_input_maskc |jd} |jd}|jd}|jd}|jd}|jd}tj| |||fd|j|j}tj| |||fd|j|j}tj| |||fd|j|j}d}|~| dry|jd}|dzdk(r|n |dz|dzz }t |j}||d<tj ||j|j}|d d|f}||||fS) Nrr1rr.r[rrr.)rrTr\r\rrr)rrr rrr@rr0rQrRr!ryr"rprr+rOr- head_dim_vrPrTrUr grad_biaslastDimlastDimAligned new_sizess r6+meta__scaled_dot_product_efficient_backwardrs{(AJ 1 I JJqMEzz!}HAJ HHQKE  ! ! Yx0kk|| F  ! ! Yx0iizz F  ! ! Yz2kk|| F I!3..$$+bLA$57R<'TV,;V)* & " KK Y__Y5E5E c8G8m, 669 ,,r8ctj|}tj|}tj|}|||fSrBr)rrr rrrr0rQrRr@rMrNrOrPr!r"rprTrUrs r6'meta__scaled_dot_product_cudnn_backwardrsA.  e $F   c "F   e $F 66 !!r8window_size_leftwindow_size_right seqused_k alibi_slopesc||jdn|jdz }||jdn|}||jdn|}|jd}|jd}tj|}|4tj|||ftj |j }nC|jd}tj||ftj |j }| ra|dkDrdnd}tj||z }|dkrd}n|dkrd}tj||||f|j|j }n,tjd|j|j }d \}}tjjrktjjrMtjd tjd }tjd tjd }nLtjd tjd }tjd tjd }|||||fS) Nrr1rrrr%r&r'NNr2r~r)rrrTrrrXrrr(r\r)r*rrbrr&)r rrrMrNrOrPr!r"r#rprrrrrr,r.r+r-r/r0total_qr1r2r3r4r5s r6meta__flash_attention_forwardrs4#,"3A9JQ9NJ*3*;A(1(9!u 2Izz"~H  'IKK $6 7++<< **Q-KK   ELL %]c yy!3k!AB  $L 3 &L[[ $6 E++<< [[%++ellK LD& }}UZZ446{{2UZZ?Ruzz&A{{Aell6BRu||FC   r8ctj|}tj|}tj|}|||fSrBr)rrr rrrr0rMrNrOrPr!r"rQrRrprr grad_querygrad_key grad_values r6meta__flash_attention_backwardr0sA0!!%(J$H!!%(J x ++r8 cu_seqlens_q cu_seqlens_k max_seqlen_qr2custom_mask_typecausal_diagonalseqlen_k window_sizec|jd}|jd}|jd}|jd}|jd}tj|||||j|j}||jddz n|}|}||J|}||n|}| rt j |dz dznd}tj|||ftj|j}tjdtjd}tjdtjd}||||||fS) Nrr1rrrrr2r~) rrTrr\rrr(rXr)r rrr8rrrr2r!rrArprrrrrrr+rvrlogsumexp_batch_dimactual_max_seqlen_qactual_max_seqlen_krwrHr4r5s r6!meta__efficient_attention_forwardrOs9,  1 A 1 A  A 2I BB ++aIrU\\ RC7C7O,++A.2VW'''**6*B,4F %*+b0A i7kk||J ;;rF ;D [[5::f =F  D&*=?R RRr8bias_requires_gradnum_splits_keyshared_storage_dqdkdvc|rtj|jd|jdk(dtj|jd|jdk(dtjg|jddd|jd|jd|j|j }|j d d}|j d d}|j d d }n?tj|}tj|}tj|}|z|jd}|d zdk(r|n |d z|d zz }t|j}||d<tj||j|j }|d d|f}n!tjd |j }||||fS)Nr1cy)Nz,seqlen must match for `shared_storage_dqdkdvr2r2r8r6rcz4meta__efficient_attention_backward..rkr8r.cy)Nz3embedding dim must match for `shared_storage_dqdkdvr2r2r8r6rcz4meta__efficient_attention_backward..rkr8rrrrr#rr.r2r) rTrerrr\rrrrr)rrr rrr8rrrr2r0r!rQrRrrrprrchunkrrrr}r~rr|s r6"meta__efficient_attention_backwardrs2 KKNciil * B   KKNciil * I  Eekk!B E EEKKO EU[[_ E++<<  \\"a( <<A&\\"a( %%e, ##C(%%e,  ))B-$+bLA$57R<'TV,;V% & " KK DKKP c8G8m, KK5<<8 xY 66r8scale_ascale_b scale_resultuse_fast_accumc d}tjjdk(xrjdk(fdtj|jxr|jfdt dk(st dk(rd} d} d } tj| j xs| fd tj| j xs| fd tjj d d zdk(fdtjj dd zdk(xrj d d zdk(fdj\j d jtjk(xrjtjk(xs<jtjk(xrjtjk(} jd k(rfjd k(rStjjtjk(xrjtjk(dn| rjtjk(rd } dznd} d}t| }t|ddz}|t|z|z|t|z|zjk(r_jk(rLtjjdtjjdntjdfdntjjtjk(xrjtjk(dtjjdk(xrjdk(fdj dk(rtj d d k(r`j dd k(rLj d k(r8tjjxrjdnj dk(rUj d j dcxk(rtdk(r"nnj d tdk(rn|j dk(rKj d j dcxk(rtdk(rnnj d k(rntjdfd||n j}tjj dj d |jS)Nc|tjtjtjtjtj fvSrBrTrW float8_e5m2float8_e4m3fnuzfloat8_e5m2fnuzfloat4_e2m1fn_x2rs r6is_fp8_or_fp4_typez2_check_scaled_mm_sizes..is_fp8_or_fp4_typeA        ! !  ! !  " "    r8rcLdjdjSNz%Inputs must be 2D but got self.dim()=z and mat2.dim()=rrErsr6rcz(_check_scaled_mm_sizes..'7 |CSTXT\T\T^S_`r8c<djdjSNz?Expected both inputs to be fp8 or fp4 types but got self.dtype=z and mat2.dtype=rrsr6rcz(_check_scaled_mm_sizes..&QRVR\R\Q]]mnrnxnxmyzr8rrc,|d|dkDxr|ddk(Srr2rs r6 is_row_majorz,_check_scaled_mm_sizes..is_row_major"!9vay(;VAY!^ ;r8c&|ddk(xr|ddkDSrr2rs r6 is_col_majorz,_check_scaled_mm_sizes..is_col_major!9>3fQi!m 3r8cV|jddk(xs|jddk(Srr tensor_2ds r6 has_zero_dimz,_check_scaled_mm_sizes..has_zero_dim)>>!$)CY^^A->!-C Cr8c*djSNz#self must be row_major, got stride rrsr6rcz(_check_scaled_mm_sizes..9$++-Ir8c*djSNz#mat2 must be col_major, got stride rrEsr6rcz(_check_scaled_mm_sizes..rr8r1rrc,djdSNzBExpected self.size(1) to be divisible by 16, but got self.size(1)=r1rrsr6rcz(_check_scaled_mm_sizes..XY]YbYbcdYeXfgr8c"djSNz?Expected both dimensions of mat2 to be divisible by 16 but got r,rsr6rcz(_check_scaled_mm_sizes..UVZV`V`Uabr8cy)NzNFor tensorwise scaling, both scale_a and scale_b must be float (fp32) tensors.r2r2r8r6rcz(_check_scaled_mm_sizes.. rkr8rr&rcy)Nzscale_a must be contiguousr2r2r8r6rcz(_check_scaled_mm_sizes..+rkr8cy)Nzscale_b must be contiguousr2r2r8r6rcz(_check_scaled_mm_sizes../rkr8Fc Zddjddjd S)NzTInvalid blockwise scaling configuration. For blockwise scaling, scale_a should have  elements, got z, scale_b should have rzr)expected_a_sizeexpected_b_sizerrsr6rcz(_check_scaled_mm_sizes..4sHFFUEVVefmfsfsfuevw//>.?w}}N__`br8cy)NzKFor rowwise scaling, both scale_a and scale_b must be float (fp32) tensors.r2r2r8r6rcz(_check_scaled_mm_sizes..=rkr8cLdjdjS)NzLFor non-tensorwise scaling, scale tensors must be 2D, but got scale_a.dim()=z and scale_b.dim()=rrrsr6rcz(_check_scaled_mm_sizes..Bs,gY`YdYdYfXhh|nunynyn{m}~r8cy)Nz@Both scale_a and scale_b must be contiguous for rowwise scaling.r2r2r8r6rcz(_check_scaled_mm_sizes..Nrkr8cJddddtdd dtddtdddtdd zdtddd jd djd d jd djd 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 (rvr&z), zscale_b should be (z<). For (BlockWise1x128, BlockWise1x128), scale_a should be (z). Got scale_a.size()=(rr1z) and scale_b.size()=(rwr7r)_krZrrrsr6rcz(_check_scaled_mm_sizes..bsCCD#EabcadeVVWUXXZ[cdfhk[lZmmpr0S0A/B"XaQTEUDVWTTUSVVXYabdfiYjXkknp p 0S0A/B"QCH//6||A.?r',,q/ARS//6||A.?r',,q/ARRSU Ur8r)rTrerr\rrrrfloat8_e8m0fnurWrrTr7rXrr)rrErrr8rr:rrrrris_blockwise_scaling block_size_k block_size_mn num_k_blockspadded_num_k_blocks _out_dtyperrrrZrs```` @@@@@r6_check_scaled_mm_sizesrs  LL  a+DHHJ!O` LL4::&I+=djj+Iz 4F"k$&75&@ < 4 D   ' =<+= I    ' =<+= I   IIaL2  " g   IIaL2  " =tyy|b'8A'= b  2 IIaL !5!55:MMU%9%99   !4!449MMU%8%88  ==?a GMMOq$8 LL .Q7==EMM3Qh "}} 3 33 " !V! M#B 5L"*<";a"? M ::=PP M ::=PP   ?2MMO6 ))+8 ))+8   LL .Q7==EMM3Qe  LL "9w{{}'9~   Q1$LLOq(LLOq(LLOq( ))+G0E0E0G^  Q1$LLOw||AK(2s:KKLLOx3'77 Q1$LLOw||AK(2s:KKLLOq(   (3J ;;tyy|TYYq\DKK XXr8c &t||||||||SrB)r)rrErrr8rr:rs r6meta_scaled_mmrss" " dGWdL)^ r8scale_recipe_ascale_recipe_b swizzle_a swizzle_bc h  !"#$%d} d} tjjdk(xrjdk(fdtj| jxr| jfdjdjdjd | jr| jrd} | z|Dcgc] }t |}}|Dcgc] }t |}}rDcgc] }t |c}nt jg r Dcgc] }t |c} nt jg tdk(std k( rd }d }d }tj|jxs|fd tj|jxs|fdtjjddzdk(fdtjjddzdk(xrjddzdk(fddtt dtt fd}dtt dtt fd}dtt dtt fd}dtt dtt fd}dtt dtt fd}dtt dtt fd}dtt dtt fd}|||rtjdjdk(xrZdjdk(xrBdjtjk(xr djtjk(dn|||rtjdjdk(xrrdjk(xrZdjtjk(xr8dj k(xr djtjk(fdn|||rqdjtjk(xr djtjk(}d$$jdk(xrk$jddzk(xrT$jddk(xr>$jdk(xs($jddk(xr$jddk(}d%%jd k(xrk%jddzk(xrT%jddk(xr>%jd k(xs(%jddk(xr%jddk(}tj|xr|xr| $%fdna|||rdjtjk(xr djtjk(}tdz dd$$jdk(xrk$jddzk(xrT$jddk(xr>$jdk(xs($jddk(xr$jddk(}d%%jd k(xrk%jddzk(xrT%jddk(xr>%jd k(xs(%jddk(xr%jddk(}tj|xr|xr| $%fd n|||rdjtjk(xr djtjk(}tdz dd$$jdk(xrk$jddzk(xrT$jddk(xr>$jdk(xs($jddk(xr$jddk(}d%%jdk(xrk%jd dzk(xrT%jddk(xr>%jdk(xs(%jddk(xr%jddk(}tj|xr|xr| $%fd!nK|||rtjj rct#jdd"jdz!t#jdd"jdz"t j#ntjddtt#jdd"dz!tjddtt#jdd"dz"t j$#tjdj!k(xrndjtj&k(xrLdj"k(xr4djtj&k(xrd#k(xr d#k(!"# fd#n|||rntdtt#ddz!t dtt#ddz"t j$#tjdj!k(xrdjtj(k(xrdjdk(xrdjtjk(xrdj"k(xrndjtj(k(xrLdjdk(xr4djtjk(xrd#k(xr d#k(!"# fd$ntjd% fd&||n j}tj* |j,'Scc}wcc}wcc}wcc}w)(Nc|tjtjtjtjtj fvSrBrrs r6rz5_check_scaled_mm_sizes_v2..is_fp8_or_fp4_typerr8c(|tjfvSrB)rTrrs r6 is_fp4_typez._check_scaled_mm_sizes_v2..is_fp4_types//111r8rcLdjdjSrrrsr6rcz+_check_scaled_mm_sizes_v2..rr8c<djdjSrrrsr6rcz+_check_scaled_mm_sizes_v2..rr8rr1rrc,|d|dkDxr|ddk(Srr2rs r6rz/_check_scaled_mm_sizes_v2..is_row_majorrr8c&|ddk(xr|ddkDSrr2rs r6rz/_check_scaled_mm_sizes_v2..is_col_majorrr8cV|jddk(xs|jddk(Srrrs r6rz/_check_scaled_mm_sizes_v2..has_zero_dimrr8c*djSrrrsr6rcz+_check_scaled_mm_sizes_v2..rr8c*djSrrrsr6rcz+_check_scaled_mm_sizes_v2..rr8rc,djdSrrrsr6rcz+_check_scaled_mm_sizes_v2..rr8c"djSrr,rsr6rcz+_check_scaled_mm_sizes_v2..rr8recipe_arecipe_bct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr& TensorWiserrs r6 is_tensorwisez0_check_scaled_mm_sizes_v2..is_tensorwisesTH ":MQ&:QK;#9#99:QK;#9#99  r8ct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr&RowWisers r6 is_rowwisez-_check_scaled_mm_sizes_v2..is_rowwisesTH "7MQ&7QK;#6#667QK;#6#66  r8ct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr& BlockWise1x32rs r6is_mxz(_check_scaled_mm_sizes_v2..is_mxsTH "=MQ&=QK;#<#<<=QK;#<#<<  r8ct|dk(xrnt|dk(xr^|dtjk(xrF|dtjk(xr.|dtjk(xr|dtjk(S)Nrrr1)rr& BlockWise1x16rrs r6is_nvz(_check_scaled_mm_sizes_v2..is_nvsH ":MQ&:QK;#<#<<:QK;#9#99:QK;#<#<< : QK;#9#99  r8ct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr&BlockWise1x128rs r6is_1x128_1x128z1_check_scaled_mm_sizes_v2..is_1x128_1x128sTH ">MQ&>QK;#=#==>QK;#=#==  r8ct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr&r BlockWise128x128rs r6is_1x128_128x128z3_check_scaled_mm_sizes_v2..is_1x128_128x128sWH "@MQ&@QK;#=#==@QK;#?#??  r8ct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr&r r rs r6is_128x128_1x128z3_check_scaled_mm_sizes_v2..is_128x128_1x128sTH ">MQ&>QK;#?#??>QK;#=#==  r8cy)Nz\For Tensorwise scaling, both scale_a and scale_b must be single element float (fp32) tensorsr2r2r8r6rcz+_check_scaled_mm_sizes_v2.. rkr8c djdddjdjdddjd S)Nz'For Rowwise scaling, scale_a must have rz elements (got: z), and scale_b must have r1rw)rr)rErrrsr6rcz+_check_scaled_mm_sizes_v2..)sa=djjm_L\]def]g]m]m]o\pq//3zz!}o=MgVWjN^N^N`Maabdr8r&cdddzdjddjdddzdjddjdS)Nz>For 1x128 x 1x128 blockwise scaling, scale a must have shape [rvr&] (got: ) and stride [1, )scale b must have shape [rwrr)rWrrsasbsr6rcz+_check_scaled_mm_sizes_v2..Fs~001s"Q#XJhrxxjPabcaddlmomvmvlwx001s"Q#XJhrxxjPabcaddlmomvmvlwwxzr8rcdddzdjddjdddzdjddjdS)Nz]For 128x128 x 1x128 blockwise scaling, L4 = {round_up(K / 128, 4)}, scale a must have shape [rvr&rrrrwrrWL4rrrrsr6rcz+_check_scaled_mm_sizes_v2..es~002t2a3hZxzQbcebffnoqoxoxnyz001s"Q#XJhrxxjPabcaddlmomvmvlwwxzr8cdddzdjddjdddzdjddjdS)Nz]For 1x128 x 128x128 blockwise scaling, L4 = {round_up(K / 128, 4)}, scale a must have shape [rvr&rrrrwrrsr6rcz+_check_scaled_mm_sizes_v2..s~001s"Q#XJhrxxjPabcaddlmomvmvlwx002t2a3hZxzQbcebffnoqoxoxnyyz|r8rcdddjdddjdtjddjddjdd 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: rw)rrTrr\expected_scale_a_elemsexpected_scale_b_elemsexpected_swizzlerrrrsr6rcz+_check_scaled_mm_sizes_v2..s78N7OwW^_`WaWgWgWiVjk--C,DGGTUJL\L\L^K_`""'"6"6!7|GAJDTDTCUU^_fgh_i_o_o^pq..>-?|IVWL>Ybclmncobppqsr8cdddjdddjddddddS) Nz!for NV scaling scale_a must have rrrz). Must have swizzle type r!r rwrr"sr6rcz+_check_scaled_mm_sizes_v2..sv78N7OwW^_`WaWgWgWiVjk--C,DGGTUJL\L\L^K_`$$4#5\)A,yYbcdYeXffgir8Fcbddddtdd dtddtdddtdd zdtddd d jd dd jd d d jd dd jd d zSrr)rWrrrrsr6rcz+_check_scaled_mm_sizes_v2..s??@cA]^_]`aRRSQTTVW_`acfWgVhhkm,HQ,<+=RC@P?QRPPQsRTU]^_adUeTffik k ,HQ,<+=RsC++21:??1+=*>bQRAS@TU++21:??1+=*>bQRAS@TTUW Wr8r)rTrerr\rr&r' NO_SWIZZLErrrrrrTr=r)r*r7SWIZZLE_32_4_4rrWrr)&rrErrrrr8r:rrrrrK_packed_multipliersirrrrrrrr r rtypes_ok scale_a_ok scale_b_okrrWrrrr#r$r%rrs&``` ` `` @@@@@@@@@r6_check_scaled_mm_sizes_v2r/s  2 LL  a+DHHJ!O` LL4::&I+=djj+Iz 1 A 1 A 1 A4::;tzz#:  0>?"k"o?N?0>?"k"o?N?/89[_9   " " /89[_9   " " 4F"k$&75&@ < 4 D   ' =<+= I    ' =<+= I   IIaL2  " g   IIaL2  " =tyy|b'8A'= b  D$5 kAR  k!2 d;>O  D- k9J  D- k9J  T+%6 ${BS  tK'8 DDU  tK'8 DDU   8 LL   "a'6AJ$$&!+6AJ$$ 56AJ$$ 5v   7 LL   #q(6AJ$$&!+6AJ$$ 56AJ$$&!+6AJ$$ 5 NN ;    EMM1Wgaj6F6F%--6W B q THHQK18+TIIaLA%TYYq\Q&R288A;!+;+Q ! PQ@Q  B q THHQK18+TIIaLA%TYYq\Q&R288A;!+;+Q ! PQ@Q  LL6Z6J nn =    EMM1Wgaj6F6F%--6W !c'1%BB r!UHHQK18+UIIaLA%UYYq\R'SBHHQK1,<,R1QRAR  B q THHQK18+TIIaLA%TYYq\Q&R288A;!+;+Q ! PQ@Q  LL6Z6J nn =    EMM1Wgaj6F6F%--6W !c'1%BB q THHQK18+TIIaLA%TYYq\Q&R288A;!+;+Q ! PQ@Q  B r!UHHQK18+UIIaLA%UYYq\R'SBHHQK1,<,R1QRAR  LL6Z6J >> 2}}  )1$**Q-)DtzzRS})T&)1$**Q-)DtzzRS})T&#.#9#9 )1$**Q-)ETZZ]B/I*&*2$**Q-)ETZZ]B/I*&$/#=#= LL   "&<<5AJ$$(<(<<5AJ$$&*@@5AJ$$(<(<<5aL$44 5 aL$44 >> 2%-a%5!RRS8T%T "%-a%5!RRS8T%T "*99  LL   "&<< 5AJ$$(;(;; 5AJ$$&!+ 5AJ$$ 5 5AJ$$&*@@ 5 AJ$$(;(;; 5 AJ$$&!+ 5AJ$$ 5 5aL$44 5aL$44 $ LL  (3J ;;q!:dkk BBw@?: :s0r r%"r*r/rcontraction_dimsc .t|||||||| |||  S)N)r8r:rrr)r/) rrErrrrrrr8rr0rs r6meta_scaled_mm_v2r2s2 %   %  r8cZt|||||d|j|jSNT)rrrrrr'r_r]s r6meta_scatter_reduce_twor6s)dCVTJ >>$** %%r8c(t|||||d|Sr4rr5s r6meta_scatter_reduce__twor8sdCVTJ Kr8cttjdjcxkxrdkncfdjdk(r0tj|tjj Stjj d|tjj S)Nrrc*djS)NzAThe probability distributions dimensions must be 1 or 2, but got rrsr6rcz"meta_multinomial..sSTYT]T]T_S`ar8r1r)rTrerrrrr)r num_samples replacementrs` r6meta_multinomialr=s| LL EIIK1a yy{a{{;ejjNN ;; 1 {%**U\\ r8c"d}|D]}||z} |Sr0r2)vsrTvs r6multiply_integersrAs$ A  Q Hr8cLtjtk(fddztjtk(fdtjtdddDxrtdDfddd\}}||gS)Nc&ddtS)Nz%It is expected output_size equals to , but got size r)num_spatial_dimsrisr6rcz'upsample_common_check..s78H7IY\]hYiXjkr8rc&ddtS)Nz$It is expected input_size equals to rDr)expected_input_dimsr?sr6rcz'upsample_common_check..s67J6K?[^_i[jZklr8c3&K|] }|dkD ywrNr2rprs r6rrz(upsample_common_check..s *aAE *c3&K|] }|dkD ywrIr2rJs r6rrz(upsample_common_check..s2NQ1q52NrKcddS)NzDInput and output sizes should be greater than 0, but got input size z and output size r2)r?risr6rcz'upsample_common_check..s \!2;-Ar8)rTrerr)r?rirErchannelsrGs``` @r6upsample_common_checkrOs LL K,,k+Q. LL J..l  LL *:ab> **Ns2N+2N/N A ""1~FH H +{ ++r8c4tjjdk7xstj ddfdt j |d}j |jtjS)Nrr1c*djS)Nz>Non-empty 3D data tensor expected but got a tensor with sizes rrsr6rcz$upsample_nearest1d..*PQVQ[Q[Q]P^_r8rEr rTrerrArrOrr$rMr)rriscalesfull_output_sizes` r6upsample_nearest1drW$ LL A/ QR0@A_- kA ??+ , / /11%8 0 r8ctjjdk7xstj ddfdt j |d}j |}tj}j\}}}}jjdk(r|dkrtj}|j|}|S) Nrr1c*djSNz>Non-empty 4D data tensor expected but got a tensor with sizes rrsr6rcz$upsample_nearest2d..:rRr8rrSrrr)rTrerrArrOrrMrrrrxr contiguous) rriscales_hscales_wrVrrrP n_channelss ` r6upsample_nearest2dr`4s LL A/ QR0@A_- kA__- .F//6M ++Az1a ||F"zA~//   ]  ;F Mr8rir?r]r^cXt||dtjjdk(fdt dD]2tjj k(fd4j |jtjS)NrrSrc"djSNzFExpected grad_output to be a tensor of dimension 4 but got: dimension rrsr6rcz-upsample_nearest2d_backward..`XYdYiYiXjkr8c DddddjS)NzCExpected grad_output to have the same shape as output; output.size() = z but got grad_output.size(rrVrrsr6rcz-upsample_nearest2d_backward..es> !s$'7':&;,QCtK4D4DQ4G3HJr8r) rOrTrerrrrr$rMr)rrir?r]r^rVrs` @@r6upsample_nearest2d_backwardrhNs-K! LLAk1X    Q #3A#6 6      , / /11+> 0 r8c4tjjdk7xstj ddfdt j |d}j |jtjS)Nrr1c*djS)Nz>Non-empty 5D data tensor expected but got a tensor with sizes rrsr6rcz$upsample_nearest3d..wrRr8r.rSrrT)rriscales_dr]r^rVs` r6upsample_nearest3drlqrXr8ctj|tj|tj}}||t|tsJt|tsJ|j }|j } t||}t||}|j|| |j|| t||t||||fS||fS)Nr)rr) rTrrrnrrrrrr!) rstabler descendingrrr@rr out_strides r6 meta_sortrqs   D !5#3#3D #LqA g1&*---':...GG XXZ "695#GY79j1Iz2F3G4w a4Kr8c tjjdk(fdtjjjk(fdj dtjjdk(fdtjj k(fdtjjjk(fdtjjdk(fdj d zztjj k(fd tjt fd fDd y) Nrc"jdSNz != 2r) input_gatessr6rcz%rnn_cell_checkSizes..;3C3C2DE0Jr8c:jdjSN != r,) hidden_gatesrusr6rcz%rnn_cell_checkSizes..s ;$$%T,*<*<)=>r8r1c"jdS)Nz != 1r) input_biassr6rcz%rnn_cell_checkSizes..sjoo5Fe3Lr8c.jdSrxr) gates_sizer|sr6rcz%rnn_cell_checkSizes..sz'')*$zl;r8c:jdjSrxr,) hidden_biasr|sr6rcz%rnn_cell_checkSizes..s z''([->->,?@r8c"jdSrtr) prev_hiddensr6rcz%rnn_cell_checkSizes..rvr8rc `jdjddddd S)Nryrz * z // z (aka rw)rr)expected_prev_hidden_numelfactorr~rursr6rcz%rnn_cell_checkSizes..sB;$$&'tK,<,.s*  HH ** * s#&cy)Nz%expected all inputs to be same devicer2r2r8r6rcz%rnn_cell_checkSizes..rkr8)rTrerrrrr)rurzr|rrrrr~s``````@@r6rnn_cell_checkSizesrsB LL!!Q&(JK LL\///>!!!$J Z__)+LM    * , ;      1 1 1 @  LL!!Q&(JK!,!1!1!!4z!AV!K LL99 LL  #J [I  8 r8c t||||d|tj|tj}tj|tj}tj|tj}|||fS)Nrr)rrTrr)rurzcxr|r workspacehycys r6_thnn_fused_lstm_cell_metarsk \:{ArR  E ?j(^*CD  __Y 'F~-z;GJ z [[5<< 0 \\* % zN2JI JBAAMoom5;;o?G 2r7J ..r8c(|r|jdn|jd}|r|jdn|jd}| }|r|||gn|||g}|j|}|"tjd|j}n|j|j}|"tjd|j}n|j|j}tjd|jtj }||||fS)Nr1rrr)rrrTrrr)rw0w1w2w3hx_cx_rrrPrr has_biasesrrrrroutput_chanelsrrrrrs r6mkldnn_rnn_layerr s&$/QEKKNJ#.QEKKNJ N  Z0*n 5 __Y 'F { [[5<< 0 ]]399 % { [[5<< 0 ]]399 % Aell%++FI 2r9 $$r8c|jdk(r%tjdk(xsdk(fdytj|jdk7fdy)NrrcdS)Nz4: Expected reduction dim -1 or 0 for scalar but got r2rrsr6rcz'zero_numel_check_dims..7swiSTWSXYr8cddS)Nz: Expected reduction dim z to have non-zero size.r2rsr6rcz'zero_numel_check_dims..<swi8=TUr8)rrTrr)rrrs ``r6zero_numel_check_dimsr3sR yyA~  1H !r Y  IIcNa  U r8c|(t||j}t||ytj|j dk7fdy)NrcdS)Nz@: Expected reduction dim to be specified for input.numel() == 0.r2rsr6rcz%check_argmax_argmin..Hstf\]r8)rrrrTrer)rrrs` r6check_argmax_argminrAsC S$((*-dC. JJLA  ] r8ctd||tj|j||fnd}t |||}|j |t jS)Nargmaxr)rrMrrrrrTr)rrrrrs r6argmax_argmin_metarLsQ$,    coSF4 PD $T4 9E >>%u{{> 33r8c||tjk(rtj}tjd||||S)Nr2r)rTjaggedrr)rr\rrrs r6 scalar_tensorrTs5  ;; %v* r8ct||jd}|jdk(rdn|j|}tj|dk\tj||kdt |j }t|dkDr|||<|j||j|tjfS)NTrrr1cy)Nzk not in range for dimensionr2r2r8r6rcztopk_meta..frkr8r) rrrrTrerrrrr)rr;rlargestr; sliceSizetopKSizes r6 topk_metar`s dhhjd ;CXXZ1_$))C.I LLa LLi!GHDJJH 8}q >>( #T^^HEKK^%P PPr8c| |Jd|j}|j} tj|| j| j| j S)Nz;segment_reduce(): Either lengths or offsets must be defined)r\rr)r\rTrr\rr) r?rror_rjrkrlrn data_contig grad_contigs r6meta__segment_reduce_backwardrnsj  '"5E 5//#K//#K   !!!!  r8cddlm}t|jd|jdkDr|j nd}t j ||dk\||kfdt|jd|jdzdz}|r%|jdkDr|jd|j||j|t jfS)Nr)rTrr1cdS)Nz9kthvalue(): selected number k out of range for dimension r2rsr6rczkthvalue_meta..sKC5Qr8r) r!rrrrrTrerrr,rr)rr;rrrdimSizers ` r6 kthvalue_metars> dhhjd ;C $ QdiinAG LLQW %Q DS!DJJsQwy$99 :E488:> S! >>% $..ekk."J JJr8c||n|}tj|jdk(d|j}|(tj|j|k(d|(tj|j|k(dtj|j|k(dtj|j|k(dtj|jdk(dtj|j |d|d zd zk(d y) NrcyNr2r2r8r6rcz(checkLSTMBackwardSizes..rkr8cyrr2r2r8r6rcz(checkLSTMBackwardSizes..rkr8cyrr2r2r8r6rcz(checkLSTMBackwardSizes..rkr8cyrr2r2r8r6rcz(checkLSTMBackwardSizes..rkr8cyrr2r2r8r6rcz(checkLSTMBackwardSizes..rkr8cyrr2r2r8r6rcz(checkLSTMBackwardSizes..rkr8rr1rcyrr2r2r8r6rcz(checkLSTMBackwardSizes..rkr8)rTrerrr)grad_hygrad_cyrrr defined_gradexp_sizes r6checkLSTMBackwardSizesrs%17wL LL!!#q(*5  "H W\\^x/< W\\^x/< LLh& 3 LLh& 3 LLA%z2 LL"hqkHQK&?!&CCZPr8c||yt|||||tj|t}tj|t}|r|j ddnd}|||fS)NNNNrrF)r)rrTrlegacy_contiguous_memory_formatr) rrrrrhas_bias grad_gatesgrad_cxr|s r6#_thnn_fused_lstm_cell_backward_implrsl7?7GRY?!!!@Jr1PQG4< q%0$I w ))r8cd}d}d}|dr|j|j}|ds|drQ|j|jd|jdf}|j|jd}|||fS)Nrr1rrr)rrrrr grad_weightr|s r6linear_backwardrsJKI1~!++FKKM: 1~Q",,l.?.?.CV[[QS_-UV  **<+<+.is_channels_lasts$""88=ATATTTr8c2r.tdk(rtjStjSj tjrtjSj tj rtj Sy)Nrr)rrTrrrXpreserve_format)rrsr6rz.meta_pixel_shuffle..pick_memory_formats| D !4 F*...***   e.E.E  F** *   e.C.C  D(( (Er8rrr)rrrr$) rupscale_factorrrHrWrrrrs ` @r6meta_pixel_shufflers DJJ! 2.>2Q RVW W 2$**=TUcTdeW U ) 2>N:;A B. (B B. (B-$**Sb/-1-b-"-I .. #C &&13& 4C Jr8cX|j|j}|j|j}|j|j}|j|j}|j|j}|j|j}|||||||fSrBr)rweight0weight1weight2weight3rcx_tmprhy_cy_grad_output_r_opt grad_hy_r_opt grad_cy_r_optrrPrrrrrrrrdiff_xdiff_hxdiff_cxdiff_w1diff_w2diff_bs r6mkldnn_rnn_layer_backwardrs4__U[[ )FmmCII&Gv||,G .G .G   w}} -F 7GVVWg EEr8) out_int32rctj||rtjntjtjS)Nr\r)rTrr>rr)r boundariesr rs r6meta_bucketizer s2    &ekkEKK-- r8cdtdk(r'tjjfdtdk(r%jrt j dtjt tfdtjdkDfdtjt tfd tjt tfd tjk\fd tjjj S) Nzhistc()rc$djdS)Nz%"histogram_cpu" not implemented for 'rrrsr6rczmeta_histc..s=ekk]!Lr8rz%_histc_cuda with floating point inputc$dtS)Nz#: argument 'bins' must be int, not rbinsrsr6rczmeta_histc..s7)>tDzlKr8rcdS)Nz: bins must be > 0, but got r2r sr6rczmeta_histc..sgY.J4&#Qr8c$dtS)Nz%: argument 'min' must be Number, not r)rrsr6rczmeta_histc..7)@c Lr8c$dtS)Nz%: argument 'max' must be Number, not r)rrsr6rczmeta_histc..#r r8cdS)Nz: max must be larger than minr2)rsr6rczmeta_histc..%sy0M%Nr8r) rrTrerrMrrnrrrrr\)rr rrrs````@r6 meta_histcr sG5U"  # # % L 5V#(?(?(A %%&MN LL4!K LLQR LL3L LL3L LLNO ;;tELL DDr8cBtj|d}tjj dk7xs#t djddDfdj |jtjS)NrrSrc3&K|] }|dkD ywrIr2)rprs r6rrz,meta_upsample_bimode2d_aa..7s!Ht$(!HrKr1c*djSr[rrsr6rcz+meta_upsample_bimode2d_aa..8rRr8r) rOrrTrerrrr$rMr)rrirr]r^rVs` r6meta_upsample_bimode2d_aar )s- kA LL Hc!Huzz|AB7G!HH_ ??+ , / /11%8 0 r8cTt||dtjjdk(fdt dD]0tjj k(fd2j |jtjS)NrrSrc"djSrcrrsr6rcz4meta_upsample_bimode2d_aa_backward..Mrdr8c DddddjS)NzD Expected grad_output to have the same shape as output; output.size(rfz but got grad_output_size(rrgsr6rcz4meta_upsample_bimode2d_aa_backward..Rs@DDE3dK[\]K^J_`D!1!1!!4 59r8r) rOrTrerrrrr$rMr)rrir?rr]r^rVrs` @@r6"meta_upsample_bimode2d_aa_backwardr ?s-K! LLAk1X    a $4Q$7 7 9     , / /11+> 0 r8cPtj|jdk(dtj|jdk(dtj|jjdtj|jjdy)Nr1cy)Nz%found_inf must be a 1-element tensor.r2r2r8r6rcz<_amp_foreach_non_finite_check_and_unscale_.._rkr8cy)Nz%inv_scale must be a 1-element tensor.r2r2r8r6rcz<_amp_foreach_non_finite_check_and_unscale_..brkr8cy)Nz!found_inf must be a float tensor.r2r2r8r6rcz<_amp_foreach_non_finite_check_and_unscale_..frkr8cy)Nz!inv_scale must be a float tensor.r2r2r8r6rcz<_amp_foreach_non_finite_check_and_unscale_..jrkr8)rTrerr\r)rr inv_scales r6*_amp_foreach_non_finite_check_and_unscale_r \s| LLQ O LLQ O LL))3 LL))3r8c,tj|SrBr)rnanposinfneginfs r6 nan_to_numr# orir8c|jtjtjtjtj hvsJd|jd|j }t||}t||}||k(r|St|j}t|j}||||c||<||<||||c||<||<|j|||S)Nz>torch.transpose_: in-place transposition is not supported for z layout) rrTr sparse_cscr sparse_bscrrrrrr)rdim0rYndimsrrs r6rrus ;;       I U\]   IIE $ &D $ &D t|  D $++- F!'vd|F4L&,!$ZdDJT T6" Kr8c|j}|jr8|j}|j}|dkr|dk(s,Jd|d|d|j dks Jd|dt |d|dkrdSdS) 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 rr1)rrrrrr)rr( rrs r6t_r* s IIE ~~__& NN$ Q9> !l,yk9J L 1 xxzQ DUG1 M  dAEAIq 551 55r8)r rsidesorterc^tjtjdkxsjddjddk(fdtjduxsjjk(fdtj|dk7xs| d|rtjntj }t tjr&tj|tjStjd|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_sequencesr6rcz#meta_searchsorted..s8 3378M8M3N2OP""&tzz"2!3 5r8cldtjdtjSgS)Nz[torch.searchsorted(): boundary and sorter must have the same size, but got boundary tensor z and got sorter tensor r/ )r0 r, sr6rcz#meta_searchsorted..sO ##'(=(=#>"??V%+%7tFLL!@ B>@@ Br8rcy)Nzetorch.searchsorted(): side and right can't be set to opposites, got side of left while right was Truer2r2r8r6rcz#meta_searchsorted..rkr8r r2r) rTrerrr>rrnr rrrr)r0 rr rr+ r, r\s`` ` r6meta_searchsortedr3 s LL O ! !"a' 9  " %CR 8   LL$?///6<<?  LL #e) $ %EKK%++E$ % U-D-D  {{2U?3I3IJJr8ctjtjtjtjfvfdy)NcdS)Nz/Unsupported input type encountered for isin(): r2rsr6rcz3_check_for_unsupported_isin_dtype..sA%Ir8)rTrer complex128 complex64rs`r6!_check_for_unsupported_isin_dtyper8 s/ LL ejj%"2"2EOODDIr8cJ|j||jdf}|S)Nrr)rr num_weightsrrrs r6meta_embedding_dense_backwardr; s*''k6F6Fr6J(KLK r8c j| rtj||||||||| | St||||||||| | SrB)r+_embedding_bag_sparse_backward!meta_embedding_bag_dense_backward) r?rrkrrmaximum_indicesr: rrPrrrs r6meta_embedding_bag_backwardr@ se22            1            r8c Ntjjtjtjtj tj fvfd|tk(rtj|duj|jdf} | S)Nc"djS)Nz$Unsupported input type encountered: r)r?sr6rcz3meta_embedding_bag_dense_backward..$s6tzzlCr8r1) rTrer\rUrVrTfloat64rrr) r?rrrr? r: rrPrrindex_grad_weights ` r6r> r> sv LL u}}ennemmU]]SSC x _D01 TYYq\'BC r8c|jd}tj|tk(dtj|j dk(tj|j dk(|jd}tj|j dk(tj|jd|k(|j |f} | S)Nr1cy)NzHembedding_bag_backward: per_sample_weights only supported for mode='sum'r2r2r8r6rcz@meta_embedding_bag_per_sample_weights_backward..9rkr8rr)rrTrerrr) r?r6rrkrrPrembedding_featuresr;rs r6.meta_embedding_bag_per_sample_weights_backwardrH ,s1 LL Z LLq! LL!#$,,q/K LL"# LLQ#556 ^^[N +F Mr8) assume_uniqueinvertctjt|txst|tdt|ts!tj||j }t|ts!tj||j }t |jt |jtj|tjS)Ncy)Nz.Irkr8rr) rTrernr r-rr8 r\rr)elements test_elementsrI rJ s r6 meta_isinrO Ds LL8V$I =&(IN h '<<1E1EF mV , ]8??K %hnn5%m&9&9:   HEJJ 77r8rctj|dk\dt|tj\}}tj ||S)Nrcy)Nz,polygamma(n, x) does not support negative n.r2r2r8r6rcz meta_polygamma..Yrkr8rr)rTrerrrr)rrrPrQs r6meta_polygammarR VsF LLaOP( ;HHOA|   D 55r8ctd)Nz.Tensor.item() cannot be called on meta tensors)r#rs r6meta_local_scalar_denserT as G HHr8c,tj|SrBrrs r6silurV frir8clt|tj\}}tj||Sr)rrrrTr)rrPrQs r6sigmoidrX ls3) ;HHOA|   D 55r8cR|jdk(}|jdk(}|r|r4|jd|jd|jdg}n"tj|jd|jdk(d|jd|jdg}n|r[tj|jd|jdk(d|jd|jdg}njtj|jd|jdk(d|jd|jd|jdg}|xs |j}tj j rZd|jz}|d|zdz |z|z}||k(r |d|z|dg} n|dg} tj|| ||j } | Stj|||j } | S) Nrrr1cyNz matrix batch sizes have to matchr2r2r8r6rcz2_create_grouped_mm_output_tensor..rkr8rcyr[ r2r2r8r6rcz2_create_grouped_mm_output_tensor..rkr8cy)Nzbatched dimension has to matchr2r2r8r6rcz2_create_grouped_mm_output_tensor..rkr8rr) rrrTrer\r)ritemsizerrr) rCrEoffsr: mat1_is_2d mat2_is_2dr alignment size_paddedrprs r6 _create_grouped_mm_output_tensorrd vsqJqJ  ! diilDIIaLAH LL !  ! ,.X  ! diim4H  LL !  ! ,.X  ! diil3H LL !  ! ,.V  ! diilDIIbMBH'TZZI }}),,, |i/!3 AIM  #"1+ 3[!DJ%q)J!! j $++  Jkk()DKKH Jr8mat_amat_br_ c tjduduk(dduxrdu} | rstjjrtjntj } tjj | k(xrj | k(fdnTtjj tjk(xrj tjk(fdtjjdvxrjdvfdjdk(} jdk(} | r| s8tjjdjdk(d | rDd } d }tj| fd tj|fd d}|d|dtjj tjk(xrj tjk(xs<j tjk(xrj tjk(fdj tjk(xrj tjk(dfd }| r| rjdnd}|dd||dd|tj|dud| s| r}tjdufdxtjjdk(fdtjj tjk(fdntjdudtj|dudtj|duxs|tjk(dt|S)Ncy)Nz,Either both scale factors are given, or noner2r2r8r6rcz)_meta_grouped_mm_common..rkr8c>djdjdS)Nz5Expected inputs of E4M3 FP8 type but got mat_a.dtype= and mat_b.dtype=rzrre rf sr6rcz)_meta_grouped_mm_common..s%KEKK=Xijojujuivvwxr8c>djdjdS)Nz1Expected inputs of BF16 type but got mat_a.dtype=rj rzrrk sr6rcz)_meta_grouped_mm_common..s%G }Tefkfqfqerrstr8)rr.cLdjdjS)Nz3Multiplicands must be 2D or 3D but got mat_a.dim()=z and mat_b.dim()=rrk sr6rcz)_meta_grouped_mm_common..s'Eeiik]Rcdidmdmdocpqr8rrrcy)Nz3contraction dimension of mat_a and mat_b must matchr2r2r8r6rcz)_meta_grouped_mm_common..rkr8cF|j}|ddkDxr|ddk(SNrr1rrmat mat_strides r6rz-_meta_grouped_mm_common..is_row_majors*Jb>A%=*R.A*= =r8cF|j}|ddk(xr|ddkDSrp rrq s r6rz-_meta_grouped_mm_common..is_col_majors*Jb>Q&=:b>A+= =r8c0djddS)NzNExpected mat_a tensor to be row major in the last two dimensions, got strides rr)re sr6rcz)_meta_grouped_mm_common..s#dejeqeqestvtwexdyzr8c0djddS)NzQExpected mat_b tensor to be column major in the last two dimensions, got strides rr)rf sr6rcz)_meta_grouped_mm_common..s#ghmhththvwywzh{g|}r8cjdz djz}jdz dk(rGtdjdz k\r%t j |zdk(fdydk(rJdz tdjk\r(t j dz |zdk(fdyt j dfdy)Nr1rrc"ddddS)Nr stride along % dim to be multiple of 16 bytes, got rzr2end_dimmat_namers sr6rczF_meta_grouped_mm_common..check_valid_strides..s))H:^G9Dijtu|j}i~~Ar8c.dddz ddz dS)Nrry r1rz rzr2r{ sr6rczF_meta_grouped_mm_common..check_valid_strides..sK)H:^GaK=HmnxzADEzEoFnGGHIr8Fc*ddjdS)NzInvalid strides/sizes, got z for strides and z for sizes.r,rq sr6rczF_meta_grouped_mm_common..check_valid_strides..s!5j\ARSVS\S\R]]hir8)r element_sizerrrrTre)r} rr rb r| rs s`` @@r6check_valid_stridesz4_meta_grouped_mm_common..check_valid_stridess'')a-#**,, ZZ\ gk "a 'Jw,?3 syy1%D -  LL7#i/14A  A %*Wq[*AS syy!F +  LL7Q;')3q8I  LLi r8re rf 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=rzrrsr6rcz)_meta_grouped_mm_common..sT~@G@M@MNNabiboboappqrr8r1c jdk(rtjjfd r;tjjjk(fdytjjdk(fdtjjdjzk(fdytjj ddk(fd tjjdjdk(fd  rtjj j dz k(fd j\}}d }t||z d  t|d tjjdk(xrjd zk( fdytjjdk(fdtjjdjdzk(fdy)NrcddS)Nrz to be contiguous.r2 scale_namesr6rcz>_meta_grouped_mm_common..check_scale..si |3EFr8cBddjdjS)NzKFor MXFP8, scale must have same number of dimensions as target tensor, but  has mat.ndim= and scale.ndim=rrr rpr sr6rcz>_meta_grouped_mm_common..check_scale..sZ"mnxmyzHILIQIQHRRbchcmcmbn!or8r1c2ddjdS)Nrz to be 1D tensor, but got D tensor.rrpr sr6rcz>_meta_grouped_mm_common..check_scale..#)J<7QRWR[R[R]Q^^g hr8rcVddjzdjddS)Nrz to have rrz elements.r,)rr rpscale_multiplierr scaled_dimsr6rcz>_meta_grouped_mm_common..check_scale..!sW)J_meta_grouped_mm_common..check_scale..&si |3[\r8cPddjddjddS)Nrz batch dimension to be r, got rzr,r sr6rcz>_meta_grouped_mm_common..check_scale..*s6i |3J399UV<.X^_d_j_jkl_m^nnopr8cBddjdjS)Nz0For MXFP8, 3d tensor should have 2d scales, but r r rr sr6rcz>_meta_grouped_mm_common..check_scale..1sF"RS]R^^lmpmumulvwGHMHRHRGS!Tr8rrr&cNdjddzdjS)NzFor MXFP8, expected mat.shape=z to have scale shape of (,z ), but got r,)G blocked_K blocked_Nrr rpsr6rcz>_meta_grouped_mm_common..check_scale..;sQ"@ Kdefdgghiru~i~h@KLQLWLWKX!Yr8c2ddjdS)Nrz to be 2D tensor, but got r rr sr6rcz>_meta_grouped_mm_common..check_scale..@r r8cVddjdzdjddS)Nrz non-batch dimension to be r1r rzr,)rr rpr r sr6rcz>_meta_grouped_mm_common..check_scale..DsT)J<7RSVS\S\]^ak]kSlRmmstyttABuCtDDE!Fr8)rrTrerXrrrr=) r rprr r r rWrrdr r r is_mxfp8s ````` @@@r6 check_scalez,_meta_grouped_mm_common..check_scale swwyA~ '')FLL swwy0o LL q(hLL A#))J*?BR*RRU  LL$)\ KKNciil2p LL chhl2T "iiGAq!!#J (Z ;I (C 0ILL A!+W A)iBW0WY LL q(hLL A#))A N*CCFr8rrrcy)Nz:Scale result tensor provided, but it is not supported yet.r2r2r8r6rcz)_meta_grouped_mm_common..Orkr8cNdjdjdS)Nz/Offsets tensor not provided, but is needed for zD/zD multiplicand layouts.rrk sr6rcz)_meta_grouped_mm_common..Us*Eeiik]RTUZU^U^U`Taaxyr8c,djdS)Nz.Offsets tensor must be 1D, but got offs.dim()=rzrr_ sr6rcz)_meta_grouped_mm_common..ZsH TUVr8c$djdS)Nz7Offsets tensor must be integer (int32) tensor, but got rzrr sr6rcz)_meta_grouped_mm_common..^sQRVR\R\Q]]^_r8cy)NzJOffsets tensor provided, but is not needed for 3D/3D multiplicand layouts.r2r2r8r6rcz)_meta_grouped_mm_common..crkr8cy)Nz2Bias tensor provided, but it is not supported yet.r2r2r8r6rcz)_meta_grouped_mm_common..hrkr8cy)Nz4If output dtype provided, it must be torch.bfloat16.r2r2r8r6rcz)_meta_grouped_mm_common..mrkr8r)rTrer)r*rrWr\rVrrrTrrr>rd )re rf rrr_ r8rr:rscaled fp8_dtype mat_a_is_2d mat_b_is_2drrr r r r s````` @r6_meta_grouped_mm_commonr s\ LL Dgo.>D 8WD%8F -2]]->->E))EDWDW  KK9 $ A )A x  KK5>> ) KekkU^^.K t  LL v7%))+"7q ))+"K))+"K k JJrNejjn , I  > >    z     } 0''w2 ]]emm + N 0N  !5!55:MMU%9%99 r   MMU11 1 6 !5!55  : z"-++DJJqMST  Iwq2BCIwq2BC D P k   y    LL aV  LL ekk)_   DL `  LL  D  LLT8Y%..8F ,E5$ JJr8c (t||dd||d|S)N)rrr_ r8rr:)r )re rf r_ r8r:s r6meta_grouped_mmr ss) #      r8c R|xstj}t||||||||| S)N)rrr_ r8rr:r)rTrVr ) re rf rrr_ r8rr:rs r6meta_scaled_grouped_mmr s:+U^^I "    !%  r8r; half_to_floatc"|r.|jtjtjfvsJt j |tj j\}}|s|n|}tj||tj}|S)Nrr ) r\rTrVrVrMrrrNrr)r;rr computation_dtyperQrs r6softmaxr sxww5::u~~6666&+&>&> uDDLL'#|(5<:KL   1L@W@W XC Jr8c  tjtdzdk(fd|jt tdz} |z tj |k\ fdt dDr|}t D]td z dz z  dkr*|j   |j z} dzdksL|j d|j dzz}v|jStd }t |D]^tdzdzz  z z dzz}tj|dk\  fd|j|`tj||j|j|jt|S) Nrrc dtS)Nz1Length of pad must be even but instead it equals rr<sr6rcz'_constant_pad_nd_meta..sCCH:Nr8c(dtddS)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_inpr=sr6rcz'_constant_pad_nd_meta..s"225c(;P 'r8c3^K|]%}t|tjxr|dk'ywrI)rnrMIntWithoutSymInt)rpr;s r6rrz(_constant_pad_nd_meta..s) I:a// 0 ;Q!V ; Is+-r1c Fdzdddzdzd 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_diffr=pad_idxsr6rcz'_constant_pad_nd_meta..sGok&1*&=%>>V7|nE#gk"2!34117! OMr8)r\rrr)rTrerrrrnarrowrrrrr\rrr) rr=rl_padc_input new_shapenew_dimrrr r r s ` @@@@@r6_constant_pad_nd_metar s LL C1 N ++K  E HME U]F LL   IS IIvu% TA519q=)G7|a!..G }gmmA&6W&E7Q;!#!..Aw}}Q/?#gPQkBR/RS T}}[&)*I 5\ "c(q1uk*fqj)CL83w{;KK qL M ! " ;;kk||))+E2  r8rrrc|jdk(sJd|j}|j}|jdk(r|df}n$|jdk(r |d|df}n g||d}|j}|j ||S)Nrz'weight' must be 2-Drr1r)rrrr\r) r6rrrr weight_shape indices_shaperr:s r6 embeddingr s ::<1 444 <._f s =JJ  r8rIr"funcr s r6_create_unary_float_meta_funcr  *4]  Ir8c |js|js |jr td|jdk(r,|j|j|jdfS|j d} |j d}|j d}|j| ||}| r1| r|j| ||}||fS|j| |||}||fS|jd}||fS)NzP_native_multi_head_attention fake implementation does not support nested tensorsrr1) is_nestedrurrrr)r rr embed_dimrh qkv_weightqkv_bias proj_weight proj_biasrx 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 !??1a3L L !! !??1h1=L L !!q) L !!r8cBt|td}|S)Nc:t||tjSrr r:s r6r z*_create_binary_float_meta_func.._fR s q!@!M!M  r8r r s r6_create_binary_float_meta_funcr Q r r8ctfd}jd}||_ttt||}|S)Nc`|g|i|}t|j|j|SrBr)rrOrrrDs r6_fnz#_register_inplace_meta.._fnv s.''' CII6 r8rP)rryrIgetattrr+)rDr inplace_names` r6_register_inplace_metar u sO 2Y kk]!$LCL 4-l3 4S 9C Jr8cftjjjk(fdg}ttrQj dk7r1tjjjk(fd|j t|dtjiS)Nc<djdjS)Nrvz for `end`, but got dtype r)rzr{sr6rczlerp.. s/%++.H Tr8rc<djdjS)Nrvz for `weight`, but got dtype r)r{r6sr6rczlerp.. s!/%++6STZT`T`Sabr8rJ) rTrer\rnrrrrRrrN)r{rzr6rOs``` r6lerpr s LL syy T 3. rkr8r)rTrerMr1r\rRrrNr s r6addcdivr sb LL  " "7== 1 6&&w}}5    w0O0W0W r8ci}dD]}t|}|D]}||vs||||<!|jD]\}}t|tjj r,t|t sJ|jtjjj|tjj|jdr|tdvst|d|jr|jdvrd|jvrtj!|| d|jvrt"j!||3d|jvrt$j!||]d |jvrt&j!||t(j!||y) N)r~ post_autograd pre_autogradCompositeImplicitAutogradr~z is a CompositeImplicitAutograd op, we shouldn't register meta function for it. Instead, we should let the decomposition run and write meta kernels for the base operators.> aten::clone aten::copy_ aten::rot90aten::_to_copyaten::empty_stridedaten::constant_pad_ndaten::as_strided_scatterzmkldnn::zmkl::zonednn::z quantized::)ritemsrnrT_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_tablerxregistryopo op_overloadrDs r6 activate_metar s:9-d3 9C--+3C=#C( 99/4466N R k5::#A#A B +z2226 EHH00556r: 88 9 9     ; 8@@""m$;;         [--//BGG UWXK,,..?DD[RTU{//11BGG UWX+"2"2"44EJJ8<<["Mm6Nr8)FrrB)NNNFrr1rr)Tr)r)rT)FF)TT)rR)FTN)TFF)TF)r)g?N)r>str)r2rrF)r2rFTN)FrFNFr)NF)rF)g?gUUUUUU?FN)NNNNN)rNNr1)NNF)FFN)Nr FFN)r FNN)Nr FNN)r FN)FN)FNNNN)NNNF)NNNNF)NrFNN)NNNN)rTT)NNrN)drr)r)rFF)r )NTTN(rcollections.abcrrenumr functoolsrtypingrtyping_extensionsrrTtorch._prims_commonrrMr r r torch._decompr r rr torch._opsr torch._primsrrrrrrrrrrrrrtorch._prims_common.wrappersrrr r!r"rr#r$torch.fx.experimentalr%rWtorch.nn.functionalr&r' torch.utilsr(rFr)r*opsr+libraryLibraryr rrrrr7r=rIrRr^rglinspacelogspacerrtaker&rrrrcummaxcumminrrrrrrr_fft_c2crrr_fft_r2crrandperm generator_outrrr randintrrlow_outrrandr_fft_c2rrrr*r1 unsqueeze_r5_sparse_semi_structured_linearr r\rB_sparse_semi_structured_mmrH_sparse_semi_structured_addmmrL_cslt_sparse_mmr\ index_reducerc index_reduce_re index_selectrisegment_reducervr unary_outrzrrrrrrrr _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_eigrrrrrrr r linalg_inv_exr linalg_ldl_factor_exrdr&linalg_ldl_solver5 linalg_lur<linalg_lu_factor_exr@linalg_lu_solverI lu_unpackrOrX linalg_qrr_rcr` _linalg_svdrlr2rr|rlinalg_solve_triangularrrr _linalg_detrrrrreflection_pad1drreplication_pad1drrreflection_pad1d_backwardrreplication_pad1d_backwardrr reflection_pad2dr replication_pad2dr_weight_norm_interface_backwardrreflection_pad2d_backwardrreplication_pad2d_backwardr r.reflection_pad3dr0replication_pad3dr3reflection_pad3d_backwardreplication_pad3d_backwardr:_pdist_forwardrXr>_pdist_backwardrDbaddbmmr] bernoullir` bernoulli_rcr;rfpoissonrh_fused_moving_avg_obs_fq_helperrymmrrrrrmiopen_batch_normr convolutionrr _has_mkldnnr r_convolution_pointwiser_linear_pointwiserhas_mklr r _mkl_linearrr rqconv2d_pointwiseqconv_pointwiser-rbinary binary_tensorrqlinear_pointwiserrlinear_dynamic_fp16linear_relu_dynamic_fp16rr r max_pool2drint4mm_packed_weight_cpur r avg_pool2dr3r8avg_pool2d_backwardr@ avg_pool3drZavg_pool3d_backwardre_adaptive_avg_pool2drj_adaptive_avg_pool3drm_adaptive_avg_pool2d_backwardrw_adaptive_avg_pool3d_backwardr|rzadaptive_max_pool2drrradaptive_max_pool3drrrrepeat_interleaverrorrrr _unsafe_indexrconvolution_backwardraddbmmr randint_liker _fused_adam_ _fused_adamw_r _fused_adamr_int_mmr _convert_weight_to_int4packr#_convert_weight_to_int4pack_for_cpur_weight_int4pack_mmr#_weight_int4pack_mm_for_cpur)r2r4r\_dyn_quant_pack_4bit_weightrh_dyn_quant_matmul_4bitrn_weight_int8pack_mmru_cdist_forwardr_cdist_backwardr_embedding_bagr_embedding_bag_forward_onlyrrnansumrmedian nanmedianr dim_valuesrPrr logical_not_rrepeatrzero_rmul_Scalardiv_ logical_and_ logical_or_ logical_xor_radd_sub_rrsubrrounddecimalsrr __rshift__r __lshift__rzerorrrfillrrelu_r _add_relurrrelu_with_noiserrrelu_with_noise_functionalrrrelu_with_noise_r index_put_unsafe_index_putr  masked_fill_r _masked_scalermasked_scatter_rmasked_scatterrmasked_scatter_backwardr index_put_rr+bmmr.r0r3r8r'r(rOr_rar max_pool2d_with_indices_backwardrlmax_pool2d_with_indicesrnfractional_max_pool2drmax_pool3d_with_indicesr max_pool3d_with_indices_backwardrrrrgrid_sampler_2d_backwardrrrrronesrzerosrselect_scatterr slice_scatterrrrrgatherrrrrrr scatter_addr scatter_add_rrr'rr_ value_reducerscatter_r#_scaled_dot_product_flash_attentionr6r?#_scaled_dot_product_cudnn_attentionrI0_scaled_dot_product_fused_attention_overrideablerL,_scaled_dot_product_flash_attention_backwardrV+_scaled_dot_product_flash_attention_for_cpurY4_scaled_dot_product_flash_attention_for_cpu_backwardr]*_scaled_dot_product_attention_math_for_mpsrt'_scaled_dot_product_efficient_attentionrx0_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_mmrr/ _scaled_mm_v2r2scatter_reducetwotwo_outr6scatter_reduce_r8 multinomialr=rArOrW_upsample_nearest_exact1dr`_upsample_nearest_exact2drh"_upsample_nearest_exact2d_backwardrl_upsample_nearest_exact3drrn values_stablerqr_thnn_fused_lstm_cellrrrrrrargminrrtopkr_segment_reduce_backwardrkthvaluerrrrrr pixel_shufflerr bucketize Tensor_outr histcr _upsample_bilinear2d_aa_upsample_bicubic2d_aar _upsample_bilinear2d_aa_backwardr r r# rr* searchsortedr3 r8 embedding_dense_backwardr; _embedding_bag_backwardr@ _embedding_bag_dense_backwardr> *_embedding_bag_per_sample_weights_backwardrH isinrO polygammarR _local_scalar_denserT rV rX rd 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.specialr r2r8r6r s`K .' #++ "U    <78) T]t_ yy~~*/--*?*?PV*W' %a)X" 8BF#3"4hr2v6F"FG 3(* t}}-.     ==5/5p !!499==12  '3 '!!))4+<+<+@+@AB %'%C%$t%%& I'I  [[$++//4;;+>+> P Xy!"!!))4+<+<+@+@AB ICI3lV$s)4 %%t}}'8'89: K;K$s) %%t}}'8'89: 8 ;8 vt}}**+"&3,3t}}$$% **   &  $$dll&6&678  **  9&   $,,"6"678  **  9& !!499==12 %)$tPT3  %%t}}'8'89: $Dv$DDI$Dc$DC$D;$DNtzz!!"#0 (t&&'(t223 "&$(     4-  4Z  {{T! 4Bt../ %)   {{T! 00t112 $(    {{T!3Dt##$$("+<,,+< \\+< 4-+< D= +< {{T! +<  +< +<+<+<%+<\t  (() I  I  I  I LL I  I I  I* It!!))*        LL      +  t  (()'*'t""**+ "!!  W  W  Wd] W d] W d] W  W  W  W, WF  $(("4"456 7txx||  $(("4"456 7txx||tzz!!"6#6tzz~~(( t!!))* + t!!%%& ' t{{""# $ t##++,    )-)t''//0,1,t33;;<= t008896:6&t<<DDEF t,,0012 FC  F # N(,    !% $VSC  F $  "        C!!))4+<+<+H+HIJ ]N+sT,K"$$,,d.A.A.E.EFG BB6BHB ! ]N+ 6 ,"  QQFQt**+ ))F)4)F),)t""#  J JF J4 JF J$ Jt}} )6)$)6))t$$% )6)$)6)&) t&&../&T0" $$,,d.M.M.Q.QR .f.6.f..dt!!))*&+))1143L3L3P3PQR T8V$     666 !" %S&%%--t/D/D/H/HIJ   '' ''  '  'K'T&&(:(:;< S#s/3ffff>T8U=4((00$2J2J2N2NOP T8V$      666 !" %QD$$,,d.B.B.F.FGH   44 44  4  4 4I4nt~~ S#s $$ $$ $  666 !" $$PtTz!2*&&(:(:;< S#fCffn8M=4$$,,d.B.B.G.GHI V[$1 'v '%(F"G '2J '$t''(  " """ $J ")"J. . . 49d3i  .". . . *. 66> .(fVt$$%   7 7 7  7  7 TM 7  7 TM7 4-7 6666 )*7&7t,,44d6R6R6V6VWX          $ Y2t$$% S#4( +( +( +( +( +(  +(  66> +()&+(^t''( ) tzz   W W W W  W  W  WWt>#;Lt$$% =&=t%%& >'>(<t--. \S/St../ \T0T2Ejt$$% =&=t%%& >'>((00  &&.. &&11 ''// ''22  \&&:<G~t$$% =&=t%%& >'> &&.. &&11 ''// ''22  \$($(Nt""#   f    v  $  t##$ PvPVPPfPQWP%P $$dll&6&678 /0'9':&&(:(:;< &*I=I t$$%&t~~ I!I  $$dll&6&678 "9"t33;;<*=*.tww 3U[[4/33,B *.2r,,r LLr IOr#Y_ r 3i#o r  r rIOd*rjQt%%--."$,,"$ LL"$ ,, "$,,% "$ $ "$  "$!&"$"$/"$Jt''(,, LL ,, I  #Y  3i I )> 889>9N9N&&:6599##::BBCD,599##55==>S?S  xx:?--:O:O 66; 7 uyy}}00 1  2  :?9N9N&&:6599##55==>599##33;;<599##33::;0<=?0d599##55<<=599##55CCDE>6599##55==>599##55<<=>?8599##55<<=599##55CCD!E>!F599##77??@599##<<DDE FA =BMM@?@F))1123H4H: ##T[[__56 ./q'7'0!!(()*'+'!!))4+=+=+E+EFG ! H 2  (()* !!+!H ~ BB*00123&889: ; (()* @+ @0012 <3 <>>?@ <A <"3"3"3"k\0012D!'$D3D4++,-;.;((()* <+ <t""**+ &, &Ft##$ G%G*t""**+  `5,`5Ft//778595  ##T[[__56 =$=7= ##T^^%;%;<=)>)    !!     Xy!" t  (()*t{{""#'$'&tzz!!"#     !!    !!        **Z       ""DJJ$7$789:  "&&(>(>?@A&&(>(>?@Atyy  !&"& !!4::#4#4567   $))"2"234"5"tzz!!"#t~~$$% F& %%&' RV"(" 0012RV;3; &&'(KO) &&(>(>(F(FGH"I"t  ''() t!!))*+t##$ % t""# 6$6 t++,!-!t&&'((Vtxx 5!5txx~~JJ 6 ;h#-Y Y Y Y Y Y Y Y Y Y Y YYYY Y !Y" #Y$ %Y&'Y( )Y*+Y,-Yx;4|38 383838 38 38 38 38 38 38 38 38 38 3838 38 !38"#38$ %38&'38lI2Xt44<<=(>(Vt++334    #5#Lt))112Q3Qht++, UI    d-dNt445 \b6bJ% V% 6% Pt v3$t,,445#6#$t##$ 8%8"t,,- \;'!(.!, !!"#.$.t&&'    )()X !!499==12    3( ""DJJNN34    5(t""**+.,.t!!))*.+.  C d /  t{{""# '$ '6   - b4t''(&)& t  !"     !!  &&     ""  889: #< < < < <  <  < 4<<;<~ S#X&889:#( ( ( (} (  (  ((( 4<(;(VEEFG $#' ' ' '} '  '  '' 4<'H'T 99("" " "  "  "  """ " """"" 4<"  ". 88#       }  4<  : AA $!"!" !" !"  !"  !"  !"!"!"}!" 4<!"  !"H??@A $"&3& 3& 3& 3&} 3&  3&  3&4-3& 4<3& 66>3&B3&l<<=>)) )) )) ))} ))  )))) 4<))?))X =="4-4- 4- 4-  4- } 4-  4-4-4-4-4-$Z4-4- 4<4-  4-n 99*!"" " "  "  "  """""" " """ 4*>?@   A  ,* $$d&D&D&L&LM   $$d&D&D&L&LM. ((00 //77"! # ,-u||+,dl  dl   : $$d&D&D&L&LM       &%Pt))112  3 t&&'4/(4/nt$$,,-$%.$%N    ##T[[%8%89:4;4t!!))*+tyy  ! Q" Qt,,- LP .   %%t}}';';<= Xy! K"> K #("9"9 Qt77??@ *A *t##++, 0- 0t!!))*+>t--556F7FD%%t~~'@'@AB 27uC | EE4 !!))4+F+F+N+NO  &55==>?   @8t>>FFGH$'')<)<=> "?"uyy~~(()*8uyy~~  !6"6$t  !    -K"-K`t,,-.t++,' -' Tt1123,t>>?@.tyy 8=e 8 8 t~~ 6c66F66t''(I&I)Ityy "v"&""t|| 6&6V66'^(,$( PK PK PK\\D PK\\D PK 4- PK 4- PK,,%PK{{T!PKPKft  $(    4- 4-  {{T!   !&''() !% $(,$(  << <<\\\\  ,,   ,,  ,,%{{T!*6t}}  v C    t##$ 3%3lt~~ $ 8 8 88 8  8  88,t33;;<  *  * &\ *c * *= * t00889 -":-"`d223d445d445d==>d==>d==>d==>dDDEdDDEtBBCtBBCtBBCtBBCtJJKtJJKtJJKtJJKt@@AtAABtAABtAAB tyy $t|| ./ t|| ./, tyy) !$,, / !$,, /  BNJr8