iT-~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'm(Z(m)Z)dd l*m+Z+m,Z,m-Z-m.Z.m/Z/dd l0m1Z1m2Z2ddl3m4Z5ddl6m7Z7m8Z8ddl9m:Z;e dZ<e dZ=e j|j~Z?e jjdddZBeCd\ZDZEZFdZGdZHdeee=eZxeIe?j d>d?ed@edAedBedzdCezdzdDe jdzf dEZ|eIe?j d?dFedGedHedDe jdzfdIZ~eIe?jdJdJddKd?edFedGedHedDe jdzf dLZeIe?j d@dMe j$dNe j$dBedzdOedzdDe jdzdPe^dQe]dRe]dSe]fdTZeIe?jjdUdVd+ed,e]dWedXe j$dYezdZe^defd[ZeIe?j jdUdVd+ed,e]dWedXe j$dYezdZe^defd\Ze/eIe?jjd]ZeIe?jjddddddd^d_edYezd`edzd'edzdaedzdbe]dce^defddZeIe?jje?jjge/deZeIe?jjd=dfZeIe?j je?j jge/dgZeIe?j jd=dhZeIe?j&jdiZeIe?j&jdjZeIe?j,jdkZeIe?j,j0dlZeIe?j4jdmZeIe?j8jddddddndoZeIe?j<jdAdpZeIe?j>jd>dqZeIe?jBjdAdrZeIe?jDjdsZeIe?jHj0dtZd+eduezfdvZd+edwedxezfdyZ dBdzeduezd{e^fd|ZdCdweduezd}ezfd~Zdwedede^duezfdZ dDdezded?edezfdZdezfdZeIe?jZje?jZj\ge/dddEdwedezde^fdZeIe?j`je?jbjge/d?edefdZeIe?jfge/ddd?efdZdedefdZeIe?jle/d+edwede^defdZeIe?jne/d=d+edwede^defdZeIe?jpe/d=d+ede^defdZeIe?jre/d=d+ede^defdZeIe?jtjdFdwede^de^fdZeIe?jvje?jvjge/d?ededefdZeIe?jxjd=dwede^fdZeIe?j|je?j|jge/ddddddd+ede^de^deeeeffdZeIe?jje?jjge/dddededede^def dZeIe?jje?jjge/ddddUddwede^deeeeffdZeIe?jje?jjge/ddddUdddwede^de^deeeeffdZeIe?jje?jjge/dUdddededede^de^def dZeIe?je/ddd dGdedede^de^deeeeff dZdezdee^e^ffdZeIe?jje?jjge/dddHdwedezdeeeffdZeIe?jje?jjge/dddddwedeeeeeffdZeIe?jj dIdwede^de^dezdzfdZdededee\e]e\e]ffdZdededxezdzdeeeffdZd?edede^fd„ZeIe?jdUddddddÜdwedede^de^dedzdedzdedzdedzdeeeeeffdĄZeIe?jje?jjgdUdddŜdwedede^de^de^dedzdefdȄZeIe?je/dddU˫ dJd+edwede^de^de^deeeff d̈́ZeIe?jjd΄ZeIe?je/ dKd?ededede^de^def dτZdЄZdфZeIe?je/d҄ZeIe?je/dӄZdԄZeIe?je/dիdքZeIe?je/dիdׄZd؄ZeIe?je/dلZeIe?je/dڄZeIe?jjdۄZeIe?jje?jje?jje?jjge/dիd܄Zd݄ZeIe?je/dބZeIe?je/d߄ZeIe?jje?jje?jje?jjge/dիdZeIe?je/dLd+ededefdZeIe?je/ded+edededef dZeIe?jje?jjge/dU˫dJdJddZeIe?jje?jjge/dd2dZeIe?jjdMdZeIe?jj dMdZeIe?jje?jjge/d?dZ eIe?jj dFdZ eIe?je/dU˫d?dDe jdzfdZ dZdNdZ d?de j$d@e j$de\e]e]zde\e]e]zde\e]e]zde^de]de\e]e]zdzfdZdZeIe?j$jde j$d@e j$dBe j$dzde j$dzde j$dzde^dedefdZeIe?j(jde j$d@e j$dBe j$de\e]de\e]de\e]de^de\e]de]fdZe j,j.re jjdddZeIe j|j2j4jdZeIe j|j2j8jdZe j,j<rOe jjdddZeIe j|j@jBdZ"e jjdddZ#eIe j|jHjJjeIe j|jHjLjeIe j|jHjLjNdZ(eIe j|jHjJjReIe j|jHjJjTdZ+eIe j|jHjXjeIe j|jHjXjNdZ-eIe j|jHjXjReIe j|jHjXjTd Z.eIe j|jHj^jeIe j|jHj`jd Z1e jjd ddZ2eIe j|jfjh dOdZ5eIe j|jfjldZ7dZ8eIe?jrj dPdZ:dZ;eIe?jxjdZ=eIe?j|e/ dPdZ?eIe?je/dիdZAeIe?jjdZCeIe?jjdZEeIe?jjdZGeIe?je/dիdZIded}ezfdZJeIe?je/dd'dZLeIe?je/dիdZNeIe?je/dd'dZPeIe?je/dիdZReIe?jj$d?d ZTeIe?jje?jjge/d!ZVeIe?jje?jjge/d"d"d#e]fd$ZWeIe j|j~jje j|j~jjge/d%ZXeIe?jj$e?jj$gd&Z[eIe?jjgd'Z]eIe?jje?jjge/dU˫dJdJdd(Z_eIe?jj$gd)ZaeIe?jje?jjgddd*d+ZdeIe?jjgddd*d,ZfeIe?jge/d-ZheIe?jgd.ZjeIe?jgd/ZleIe?jgd0ZneIe?jgd1ZpeIe?jgd2Zqd3e]d4e]de]fd5Zrd6ZseIe?jgdBedzfd7ZueIe?jgd8ZweIe?jgd9ZyeIe?jjd:Z{eIe?je/d;Z}eIe?jj dQd<ZeIe?jjd=ZdBd>ZeIe?jje?jjge/dRdd?d@ZeIe?j je?j jgdAZeIe?j je?j je?j je?j je?jje?jjge/d&d'dSdBZeIe?jjdCZeIe?jjdDZeIe?j jdEZeIe?j$j&e?j(j&e?j$j$e?j(j$e?j*je?j,je?j.jgdFZeIe?j2j&e?j4j&e?j2j$e?j4j$gd dGZeIe?j8j&e?j:j&gd dHZeIe?j>je?j>j@gdIZdJZeIe?jFj$e?jFj&gdKZeIe?jJj$e?jJj&gdLZeIe?jNjdMZeIe?jRj$e?jRj&gdNZeIe?jVj$e?jVj&gdOZeIe?jZjdPZeIe?j^j$e/d defdQZeIe?jbge/ dTdRZeIe?jfg dTdSZeIe?jjg dTdTZeIe?jnje?jpjgd=dUZeIe?jtj&dVZeIe?jxjdWZeIe?j|dXZeIe?je/dYZeIe?jdZZeIe?jjd=d[ZŐdAd\ZeIe?jjd]ZeIe?jjd^Zɐd_Zʐd`ZːdaZ̐dbZ d=d?edce]dde]dee]dfe]dge]dhe]die]dje]dke]dle]dme]dne]doe]dpe]dqe]dre]dse]dte]due]dezdve^f,dwZΐdxZd?ededce]dde]dee]dfe]dge]dhe]die]dje]dke]dle]dpe]dqe]dre]dse]dte]due]dezf&dyZАdzZeIe?jjd{ZeIe?jj dOd|ZeIe?jjd}ZeIe?je/dd' dOd~ZeIe?je/dիdZd?edefdZGddeZd?edede]fdZeIe?jjdZeIe?je/dZeIe?je/dՐddZeIe?jjgdZeIe?jj dUdZeIe?jje?jjge/dddddddZeIe?jje?jjge/dddddddZeIe?jjdZeIe?jjdVdZdBd,e]de]de^fdZdZdZeIe?jjd=dZd=dZd?dZdZd?dZdWdZeIe?jjdZeIe?jdZeIe?jje?jje?jje?jjge/d?dZeIe?jje?jje?jje?jjgd?dZeIe?jjg dXdedededede^de^dedzfdZeIe?jjfg dXdededededzdedzdedzdede^de^dedzfdZdedee]dffdZeIe?jg dXdededededzde^dede^de^dedzfdZeIe?jg dYdededededzdede^de^dedzfdZ eIe?jg d?dedededededededede]de]dede^dedededzfdZ eIe?jg dZdedededede^dedzdedzfdZeIe?jg dAdededededededede^dedzdedzfdZeIe?j"g d[dededededzdede^dedzdedzdeeeffdZeIe?j&g d\dededededzde^de^dedzfdZeIe?j*g d]dedededededzdededededede\e^de^dedzfdÄZeIe?j.g d?dededededededededededede]de]dede^dedzf dĄZeIe?j2jg dUdededededzdedzde]de]dede^de^dedzde]dzde]dzdedzdedzfdɄZeIe?j2jfg dUdededededzdedzde]de]dede^de^dedzdedzdedzdedzde]dzde]dzdedzdedzf$dʄZeIe?j8g d>dedededededededede]de]dede^dedededzde]dzde]dzf"d˄ZeIe?j<g d^dedededBedzdedzdedzde]dzde]dzdede]de^dedzdedzdedzde]dzfdԄZeIe?j@g dWdededededBedzdedzdedzde jBde jBdedededede]de^dedzde]dzde^f$d؄Z" d_d+e j$dHe j$de j$de j$dBe j$dzde j$dzdDe jdzde^fd݄Z#eIe?jHjg d_d+e j$dHe j$de j$de j$dBe j$dzde j$dzdDe jdzde^fdބZ% d`d+e j$dHe j$de\e j$de\e7de\e j$de\e7dBe j$dzdDe jdzde\e8dzde\e8dzde^fdZ&eIe?jNjg d_d+e j$dHe j$de\e j$de\e7de\e8de\e j$de\e7de\e8dBe j$dzde jdzde\e]dzde^fdZ(eIe?jRjTe?jRjVge/dBdZ,eIe?jZjTdBdZ.eIe?j^je?j^jge/d=dd2dZ0dZ1dZ2eIe?jfje?jhjgd?dZ3eIe?jjje?jljgdAdZ5eIe?jnje?jpjg dAdedee]e jBzdee]e jBzdedzdedzf dZ7eIe?jrje?jtjgd>dZ9eIe?jvje?jvjxe?jvje?jvjzgdadZ>dZ?eIe?jj dAdZAeIe?jjdZBeIe?jjdZCeIe?jjdZDdZEdZFeIe?jje?jjgdRdZIeIe?jjdbdZJeIe?jjdcdZLeIe?je/ dddZNeIe?jje?jjge/d&d'dSdZPe jZRdZSeIe?jjdZTeIe?jjdZUeIe?jjdZWeIe?jjdZXeIe?jj$e?jjge/ddddZ[eIe?jj&e?jjgdddd+e&ded e^d e^fd Z]eIe?jge/ded Z_eIe?jje?jjg dAd ZbeIe?jjg dAdZdeIe?jjdZeeIe?jje?jjge/d>dZfeIe j|j~jdZgeIe j|j~jdZheIe?je/ddddddZjdZkeIe?jdZmeIe?j dfdZoeIe?j dfdZqeIe?j dfdZseIe?je/ddddZueIe?je/de]d+edefdZweIe?jd+efdZyeIe?je/dU˫d+edefdZzeIe?je/d+edefd Z{d!Z| d`d"ed#ede j$dzde j$dzd$edzdBedzde j$dzdDe jdzde^fd%Z}eIe?je/ d>d"ed#ed$edzdBedzdDe jdzdef d&ZeIe?jg d`d"e j$d#e j$de j$de j$d$e j$dzdBe j$dzde j$dzdDe jdzde^fd'ZeIe?jj&dgd(ZeIe?je/d)ed,e]d*e^defd+ZeIe?j e/d d,ZeIe?je/ dhd@ed'ed-e]d.e^d/e^def d0ZeIe?jj did&edae\ed1e\e]d2efd3Zd4ZeIe?jj djd5Zd6Zee?jee?j ee?j"ee?j$ee?j&ee?j(ee?j*ee?j,ee?j.ee?j0ee?j2ee?j4ee?j6ee?j8ee?j:ee?j<ee?j>ee?j@ee?jBee?jDee?jFd7ZeIe?jJe/d8ZeIe?jLe/dJd9d:ZeIe?jNe/dJd9d;Zee?jJZee?jLZee?jNZddl0Z ddlZ ddlZ d<Zey(kN)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_forNumber NumberTypesuggest_memory_formatsym_min 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 j/var/www/vps2.regionflexible.com/Desarrollo/venv/lib/python3.12/site-packages/torch/_meta_registrations.pyceil_divr99s EAI! c||zdz |z|zS)z$Rounds up x to nearest multiple of yr3r4xys r8round_upr?=s UQY1  !!r:returncfd}|S)NcVtfd}tj|S)Nc(tt|yN)r r)opfns r8registerz0register_meta..wrapper..registerFs  B 3r:)r pytree tree_map_)rFrGrEs` r8wrapperzregister_meta..wrapperCs)  $ 4 2& r:r4)rErJs` r8 register_metarKBs Nr:type_promotionctj|d|i\}}|Dcgc]}t||}}t|}t |dt j iScc}w)Ntype_promotion_kindrL)utilsrr r&rrDEFAULT)rLargs_ result_dtyper=s r8elementwise_metarTOss .. *OA|?C C #A| 4 CD C T "D "  BJJ  DsActjtjtjtjtj tj i}|j||SrD)torch complex32halfcfloatfloatcdoubledoubleget)dtype from_complexs r8toRealValueTyper`csE  ekk u||L   E5 ))r:clttg|tjk(fdy)NcddS)Nzoutput with shape z# doesn't match the broadcast shape r4)broadcasted_shape self_shapesr8z)check_inplace_broadcast..ps$ZL0STeSfgr:)tupler%rV_check)rd args_shapercs` @r8check_inplace_broadcastrils0/ HZHI LLZ'gr:Fc  ttjr(tjj dk(dttjr(tjj dk(dt dfDrZt jtj  nFtjt j fdnxstjttjstdttjttfdttstdttjdk\d tjf|d || S) NrcyNz:linspace only supports 0-dimensional start and end tensorsr4r4r:r8rez(meta_linspace_logspace..r:cyrlr4r4r:r8rez(meta_linspace_logspace..rmr:c3<K|]}t|tywrD) isinstancecomplex).0args r8 z)meta_linspace_logspace..s C:c7 # CscddS)Nzlinspace(): inferred dtype z& can't be safely cast to passed dtype r4)default_complex_dtyper^sr8rez(meta_linspace_logspace..s56K5LLrsxryzr:zdtype must be torch.dtype, got cdtjdtjdtjdS)Nz4received an invalid combination of arguments - got (, ))type__name__)endstartstepssr8rez(meta_linspace_logspace..sDu+r$s),,-RU 0D0D/EQHr:zsteps must be IntLike, got cy)Nz$number of steps must be non-negativer4r4r:r8rez(meta_linspace_logspace..rmr:metar^layoutdevice pin_memory requires_grad)rprVr rgdimanyrOrget_default_dtypeis_complex_dtyper^AssertionErrorrz _check_typerempty) r}r|r~baser^rrrrrvs ``` ` @r8meta_linspace_logspacertsn%& IIK1  P #u||$ GGIN P  CsE/B CC % A A  # # %!  =)E LL&&u-z  2002 eU[[ )>tE{mLMM 5'" H eW %:4;-HII LL!KL ;; #  r:ctjjtjk(fdtj|j dk(xrj dk7 d|j jS)Nc"djS)Nz2take(): Expected a long tensor for index, but got r^indexsr8rezmeta_take..sDU[[MRr:rcy)Nz*take(): tried to take from an empty tensorr4r4r:r8rezmeta_take..rmr:)rVrgr^long _check_indexnumel new_emptyshape)selfrs `r8 meta_takersm LL uzz!R   ZZ\Q  55;;=A#56< >>%++ &&r:rcTj}j}tj||k(dtjjdk(xrjdk(fdt j j }j |S)Ncy)Nz=linalg.cross: inputs must have the same number of dimensions.r4r4r:r8rezlinalg_cross..rmr:r0cVddjdjS)Nzlinalg.cross: inputs dimension z must have length 3. Got  and size)rotherrsr8rezlinalg_cross..s6-cU399S>"% 3'8 :r:)ndimrVrgrr%rr)rrrx_dy_d out_shapes``` r8 linalg_crossrs ))C **C LL s O LL #!4 31 4 "$**ekk:I >>) $$r:c|t|dt|dtj|tjS)Nzlinalg.matrix_exp memory_format)squareCheckInputscheckFloatingOrComplexrV empty_likecontiguous_formatrs r8linalg_matrix_exprs3d/04!45   D0G0G HHr:valuesindicescZtj|j|j|j}tj|j|jtj }|j dk7r%|jdk7rt||j||fS)Nrr^r) rVrrrr^int64rrmaybe_wrap_dim)rrrrs r8 cummaxminrsp [[DKKtzz JFkk$**T[[ LG zz|qTYY!^sDII& 7?r:cxt||jtj|tjSNr)rrrVrr)rrs r8 logcumsumexprs+3 "   D0G0G HHr:cD|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|SrDr4)r= self_stridess r8rez_exec_fft..s <?r:keyreverserrr3r)rlenlistrangeappendstridesortpermuterreshaperresize_rVr as_strided_storage_offset)outr out_sizesrforwardr signal_ndim batch_dims dim_permuterRis_transformed_dimdleftright batch_endtmpinput batched_sizes batch_sizebatched_out_sizesi out_strides batch_numelrs @r8 _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|SrDr4)rrs r8rez_sort_dims..0s l1or:)r)rrrintr)rrr sorted_dimsrs @r8 _sort_dimsr,sLs)K;;=L6#k"S%667<< %= r:c tj|jj|s|j St ||}|j |j}t|||j||S)Nr) rVrgr^ is_complexclonerrrr)rr normalizationrrrs r8 meta_fft_c2cr7sb LL&&' zz|T3'K .. %C S$ [' JJr:cft|tkDst|dk\r|ddk(r |ddk(ryy)Nrr3FT)rcufft_max_ndimrs r8use_optimized_cufft_pathrFs3 3x. SX]s1v{s1vQR{r:cztj|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) Nrrr3cudaxpurTrc|SrDr4)rstridess r8rezmeta_fft_r2c..ts '!*r:rr)rVrgr^is_floating_pointrr device_hintrrOrrrrrrminrrr)rrronesided input_sizesrlast_dimlast_dim_halfsizeonesided_sizesoutputworking_tensor target_sizesrmax_dims last_dimsrs @r8 meta_fft_r2crMs> 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  r:) generatorcBt|tj|gSrD)r!rVSize)nrrs r8 meta_randpermr s S%**aS/ 22r:r^rrrc6tj|||||SNr rVr)rr^rrrs r8meta_randperm_defaultrs  ;; vf r:cxdtjkDfdtj|||||S)NrcddSNz:random_ expects 'from' to be less than 'to', but got from=z >= to=r4highlowsr8rezmeta_randint..LSEQXY]X^_r:r rVrgr)rrr^rrrrs` @r8 meta_randintrs> C LL s _ ;; E&J r:cttjkDfdtj|||||S)NcddSrr4rsr8rez"meta_randint_low..rr:r r)rrrr^rrrs`` r8meta_randint_lowrs9 LL s _ ;; E&J r:c6tj|||||Sr r )rr^rrrs r8meta_rand_defaultrs  ;; E&J r:rlastdimctj|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) NrrrrFrr3r)rVrgr^rrrrrr`rrrrrr) rrrrrrtemprc2c_dimsrs r8 meta_fft_c2rr!sk LL&&'4F"% $ #b' 1LM #C ( )@)@ A  3x!|#D#cr(Aw?zz0G0GzHVT9s2wiO O s8a<3BxH xNEbc(C& $ #b'nnYodjj.InJeYUCCr:cJddlm}||s#tj|dk(r t dt |t ra|j||}|j|jk7r.tjj||j|S)Nr)free_unbacked_symbolsr3zQmore than one element of the written-to tensor refers to a single memory location) %torch.fx.experimental.symbolic_shapesr#rV_debug_has_internal_overlap RuntimeErrorrpr torr- expand_copydefault)rsrc non_blockingr# intermediates r8 meta_copy_r-s L "$ 'E,M,Md,SWX,X _  #vvvdL1 99;,++- -    $ $\499; ? Kr:ct|j}t|j}||jk\rdn ||||z}|j |d|j ||||fSr2)rrrrinsert)tensorr result_sizesresult_strides new_strides r8inferUnsqueezeGeometryr4sq &L&--/*NVZZ\)|C/@>RUCV/VJQ#z*  ''r:ct||jdz}t||\}}|j|||Sr2)rrr4r)rrg_sizes g_stridess r8meta_unsqueeze_r8s> dhhj1n -C/c:GYWi( Kr:rweight_metabias_activation_opt out_dtypecFt|j}|R|jd|jdk7r/td|jdd|jd|jd|jddz k7r2td|jdd|jddz |jd|d<t |jdk7r"td t |jd d|jdf}|K|j t jk(r|t jk(std |j d ||j|| |j n| j||}|S)Nrz%output size mismatch: weight.size(0)= != bias.size(0)=r3rrzweight.size(1)=z != input.size(-1)/2=z0we can only handle the squashed input case, got D inputzKout_dtype is only supported for i8i8->i32 linear operator, got input.dtype= , out_dtype=r) rrrrrr^rVint8int32r as_strided) rr9r:r;r<r= output_sizestransposed_stridesrs r8meta_sparse_structured_linearrG%s $L  ;;q>TYYq\ ) 7 A7GGXY]YbYbcdYeXfg {{1~B!++fkk!n--B5::b>TUCUBV W  {{1~L  5;;1>s5;;?O>PPW X  UZZ]+ uzz)i5;;.F ]^c^i^i]jjvxAwBC __&.ekkIj12 Mr:mat1 mat1_metamat2c2t|jdk7r"tdt|jdt|jdk7r"tdt|jdt|jdk7r"tdt|jd|jd|jddz k7r2td|jdd |jddz |jd|jdg}|K|jt j k(r|t jk(std |jd ||j|| |jn| }|S) Nrmat1 must be 2D, got Dmat1_meta must be 2D, got mat2 must be 2D, got r3r mat1.size(1)= != mat2.size(0)/2=Jout_dtype is only supported for i8i8->i32 linear operator, got mat2.dtype=rAr rrrrr^rVrBrCr)rHrIrJr=rErs r8meta_sparse_structured_mmrTQsv 4::!4S_4EQGHH 9??q 9#ioo:N9OqQRR 4::!4S_4EQGHH yy|tyy|a''DIIaL>)) z matmul!r)r^rVfloat32float16bfloat16rB float8_e4m3fnrrr is_contiguousrrCr) rYrZr;rUr=r[r\r]r^is_8bit_input_typerm output_shapes r8meta__cslt_sparse_mmrhs}}      OPWP]P] _  W]]*3L4F4F3GuW]]O \   7==Q;C >+ &&r:)lengthsroffsetsaxisunsafeinitialdatarvrwrxryc| tdfd}|||jS|+|jdd|jddz fz} || Std)Nz?segment_reduce(): indices based reduction is not supported yet.ctj|jdzdzjdtjS)Nr3rr^rr)rVrrr^r) lengths_shaperxr{s r8segment_reduce_lengths_tensorz:meta_segment_reduce..segment_reduce_lengths_tensors>{{ DJJtaxz2 2**11   r:rr3z>" r:ctj|j|f}t|||}|j ||j |t j fSNrrOreduction_dimsr_compute_reduction_shaperrVrrrkeepdimrgs r8 meta_max_dimr+R   tzzC6 2C+D#w?L |$ |5::6 r:c$|jdSrrrs r8meta_minr5rr:ctj|j|f}t|||}|j ||j |t j fSrrrs r8 meta_min_dimr;rr:c|jrt|j}nt|tj \}}t j||SNrNr)rrr^rr INT_TO_FLOATrVr)rrSrRs r8 meta_anglerEsI / ; ,  ? L L <   D 55r:ctj||j|j|j tj |SrD)rV_resize_output_rrcopy_angle)rrs r8meta_angle_outrQs6 #tyy{DKK8 99U[[& ''r:cyrDr4)vals r8 assert_asyncrW r:cyrDr4)r assert_msgs r8assert_async_metar\rr:cyrDr4)ss r8 print_metararr:r^rrrrc0tjddS)Nrrrr rs r8make_dep_tokenrfs ;;q ((r:chddlm}t|ttfr t d||||y)Nr)constrain_range'Constraining SymFloat or Symbool is nyirmax)r$rrpr r ValueError)rrrrs r8sym_constrain_rangerrs/F$7+,BCCDcs+r:c6tj||||SNr)r-rrrr dep_tokens r8functional_sym_constrain_ranger|sTs4 r:c.ddlm}||tj|dk\yt |t t fr tdt|tur5|tj||k\|tj||ky||||y)Nr)_constrain_range_for_sizerr) r$rrVrgrpr r rrzr)rrrrs r8sym_constrain_range_for_sizersP {s{ TQY$7+,BCC DzS ? LL % ? LL %d5r:c6tj||||Sr)r-rrs r8'functional_sym_constrain_range_for_sizers%%d%= r:c|SrDr4)rrrs r8functional_assert_async_metars r:f_namecjdkrtdjtjj dj dk(fdy)Nrz8: The input tensor must have at least 2 dimensions, got rcVdjddjddS)Nz5: A must be batches of square matrices, but they are r by r matricesr)rrsr8rez#squareCheckInputs..s16( " d499R=/Dr:)rrrVrgr)rrs``r8rrs^ xxzA~hNtxxzl [   LL " 2& Dr:Anamectjjjk(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.rrrsr8rez(linearSolveCheckInputs..s%H{{m:ahhZy :r:c>djdjdS)Nz=Expected b and A to have the same dtype, but found b of type z and A of type rrrsr8rez(linearSolveCheckInputs..s%Kzzl/!'') =r:rrcRdjddjddS)Nz3A must be batches of square matrices, but they are rrrrrrsr8rez(linearSolveCheckInputs..s0FF2J.sR,TF3 D $TYYr]O4 "  Hr:)rVrgrr^r)rrrs```r8linearSolveCheckInputsrs LL qxx  LL agg  LL r affRj   LL r diim# r:tallow_low_precision_dtypescJ|jtj|jxs|j fd|sYtjtj tj tjtjfvfdyy)NcdS)Nz<: Expected a floating point or complex tensor as input. Got r4r^rsr8rez(checkFloatingOrComplex..s6(VW\V]^r:cdS)Nz*: Low precision dtypes not supported. Got r4rsr8rez(checkFloatingOrComplex..svhHPr:) r^rVrgrrrZr\rYr[)rrrr^s ` @r8rrsn GGE LL /^ & ekk5<<u}}M M P  &r:arg_namec^tj|jdk\fdy)NrcddS)Nz: The input tensor z! must have at least 2 dimensions.r4)rrsr8rezcheckIsMatrix..s6(-hZ7XYr:)rVrgr)rrrs ``r8 checkIsMatrixrs LL 1 Yr:Brctttjr#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=rrryr)rrrrsr8rez#checkInputsSolver..s[hHxX. AaffRj\qvvbzl!AFF2J.s5i{;-/dm4 nU\\N Lr:)rVrgr)rrrrs````r8checkSameDevicers&  LL % r:UPLOcj}tjtdk(xr |dk(xs|dk(fdy)Nr3ULcdS)Nz1Expected UPLO argument to be 'L' or 'U', but got r4)rsr8rezcheckUplo..sCD6Jr:)upperrVrgr)rUPLO_uppercases` r8 checkUplor s<ZZ\N LL D QKNc1J^s5JJr: 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 r8meta__linalg_eighrsa' dO ME{{5!  ;EU ST{{A3 IIK ;;uOAGG$<; =D :r:ct|dtj|jr |jntj|j}|j |j dd|S)Nzlinalg.eigvalsrr)rrOrr^rrr)r complex_dtypes r8meta__linalg_eigvalsr&sce-.  ! !%++ .   . .u{{ ; ??5;;s+=? AAr:ct|dtj|jr |jntj|j}|j |j dd|}|j |j |}t|dk(}|j|j t|j |||fS)Nz linalg.eigrrrr) rrOrr^rrrrrr)rrrvectorsis_cudas r8meta_linalg_eigr2se\*  ! !%++ .   . .u{{ ; __U[["-]_ CFooekko?G% F*G  0P 7?r:r*cv|jjtjj ddS)Nrrr)mTrrVr transpose)r*s r8cloneBatchedColumnMajorrDs* 66< H HR PPr:rct|SrD)r)rrrs r8_cholesky_solve_helperrHs #4 ((r:ctjjdk\fdtjjdk\fdtd\}}t |||S)Nrc$djdS)Nz-b should have at least 2 dimensions, but has  dimensions insteadrrsr8rez cholesky_solve..Ss? {J]^r:c$djdS)Nz-u should have at least 2 dimensions, but has rrrsr8rez cholesky_solve..Ws?xGZ[r:cholesky_solve)rVrgr!_linalg_broadcast_batch_dims_namer)rrrself_broadcasted A_broadcasteds`` r8r r Nsh LL Q^ LL ! ['H a!'#m ""2M5 IIr:c|jdk(r%tj|tjSt |dt |S)Nrrcholesky)rrVrlegacy_contiguous_formatrrrrs r8rr_s@ zz|qE4R4RSSdJ' "4 ((r:c0t|dt|S)Ncholesky_inverse)rrrs r8rrhsd./ "4 ((r: check_errorsct|dt|d|j}t|}t |d}|j |}|j |||j |d|dz tj}||fS)Nzlinalg.choleskyFrrr) rrrrrrrrVrC)rrrA_shaper L_stridesrinfoss r8linalg_cholesky_exrpsa*+1/0ggG w.rmr:rrcy)Nzbtorch.linalg.householder_product: input.shape[-2] must be greater than or equal to input.shape[-1]r4r4r:r8rez,linalg_householder_product..rmr:cy)Nz`torch.linalg.householder_product: input.shape[-1] must be greater than or equal to tau.shape[-1]r4r4r:r8rez,linalg_householder_product..rmr:r3c<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 rrrsr8rez,linalg_householder_product..')), 2Nuzzl \r:cdS)Nzltorch.linalg.householder_product: Expected batch dimensions of tau to be equal to input.shape[:-2], but got r4actual_batch_tau_shapesr8rez,linalg_householder_product..66L5MOr:c<djdjS)Nz,torch.linalg.householder_product: tau dtype z does not match input dtype rr sr8rez,linalg_householder_product..s#:399+*5;;- 9r:z torch.linalg.householder_productrFrrrr^r) rVrgrrrr^r empty_stridedrr)rrexpected_batch_tau_shaper$s`` @r8linalg_householder_productr*sK  LL aZ LL 2%**R.(t LL 2#((2,&r  LL SXX"  zzA~#(;;s#3 !$3B "&> >   LL U[[  6UEJ    [[*5;;%Hkk||  r:c2t|dt|dd|j|j}|j |jt |jd|j|jddt j}||fS)Nz linalg.inv_exF)rrrrrrrrrrrVrC)rrrrs r8linalg_inv_ex_metar-sqa)1o%P AGGAMM!''6qww%PQ KK EKKK 8E e8Or:LDpivotsinfo) hermitianrr1ct|dt|dtj|jt |jd|j |j}|j|jddtj}|j|jddtj}|||fS)Nztorch.linalg.ldl_factor_exFrr'rrr) rrrVr(rrr^rrr)rr1rr.r/r0s r8linalg_ldl_factor_ex_metar3sd894!=>    ZZ*4::Gjj{{  B ^^DJJsO599^ =F >>$**Sb/> ;D vt r:)r1cdtdtdtdtjj 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)rsr8rez'linalg_ldl_solve_meta..&&!4 6r:rc$djdS)Nzjtorch.linalg.ldl_solve: Expected LD.shape[:-1] and pivots.shape to be the same, but got pivots with shape  insteadrr/sr8rez'linalg_ldl_solve_meta..))/h @r:c"djS)Nz.sNv||n]r:c<djdjS)Nz!torch.linalg.ldl_solve: LD dtype z does not match b dtype r)rr.sr8rez'linalg_ldl_solve_meta..s"3BHH:=UVWV]V]U^_r:Frr')rrrrVrgrrrOis_integer_dtyper^_linalg_broadcast_batch_dimsr(rr)r.r/rr1expected_pivots_shapeB_broadcast_sizerRs``` r8linalg_ldl_solve_metarBsb232781b":; LL !  HHSbM LL-  LL v||,] LL AGG_7q"=a    *+;uMggxx  r:Pr)pivotrDchtjjdk\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: r8r9rsr8rez linalg_lu_meta.. sRSTSZSZR[[cdr:rrr)rVrgrrrrr) rrDsizesrfrkrCrrs ` r8linalg_lu_metarIs LL ! d ME b A b A1 AE"I KK  KK E"I EAE"IE"I EA a7Nr:LU)rDrctjjdk\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: r8r9rsr8rez*linalg_lu_factor_ex_meta..,sXYZY`Y`Xaaijr:rrFrr'r) rVrgrrrr(rr^rrrrr) rrDrrGrfrrJr/r0s ` r8linalg_lu_factor_ex_metarM"s LL ! j ME b A b A    *5EBggxx  B IIK1 E"I [[eii[ 0F IIK ;;uEII; .D vt r:)radjointrNctdtjjjk(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 r8r)rrJsr8rez&linalg_lu_solve_meta..Us($$&HH:_QWWIX Or:cy)NzElinalg.lu_solve: pivots should be a Tensor of scalar type torch.int32r4r4r:r8rez&linalg_lu_solve_meta..\rmr:zlinalg.lu_solvercy)NzYlinalg.lu_solve: Number of pivots per batch should be same as the dimension of the matrixr4r4r:r8rez&linalg_lu_solve_meta..drmr:c$djdS)Nzclinalg.lu_solve: Expected LU.shape[:-1] and pivots.shape to be the same, but got pivots with shape r8r9r:sr8rez&linalg_lu_solve_meta..jr;r:rr'r)rrVrgr^rrrrrr?r(rrrrconj)rJr/rrrNrArRrs``` r8linalg_lu_solve_metarVGs.267 LL AGG  LL  !W b12b!T#45 LL  v{{2&k  LL " % 7q"=a  *+;4xPggxx F||~4    [[]F Mr: 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: r8r9)rJsr8rez lu_unpack_meta..sXY[YaYaXbbjkr:c y)Nztorch.lu_unpack: LU_pivots is expected to be a contiguous tensor of torch.int32 dtype. Note: this function is intended to be used with the output produced by torch.linalg.lu_factorr4r4r:r8rez lu_unpack_meta..s pr:rrr) rVrgrr^rCrrrr) rJr/rWrXrGrfrrHrCrrs ` r8lu_unpack_metar\s LL 1 k LLEKK '   NE b A b A Aq AE"I LL  LL! b LL b b LL  LL!  LL!  a7Nr:modecdk(rd}d}||fSdk(rd}d}||fSdk(rd}d}||fStjdfdfS)NreducedTcompleteFrcddS)Nzqr received unrecognized mode 'z=' but expected one of 'reduced' (default), 'r', or 'complete'r4)r]sr8rez _parse_qr_mode..s1$8NOr:rVrg)r] compute_qr_s` r8_parse_qr_moderes y  g     g    g      g r:QRct|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) rrrerrrrrr) rr]rd reduced_moderfrrHQ_shaperfR_shapergs r8linalg_qr_metarls![!1k*,T2I|  A  A Aq Aqww-'aQ KK  g:7eTU KK 177mG#9!!GBK GAMM'6w%PQ a4Kr:sign logabsdetct|dt|dd|j}|j|dd}|j|ddt |j }t j|t|d|j |j}|j|ddt j}||||fS)Nzlinalg.slogdetFrrr'r) rrrrr`r^rVr(rrrC)rrrmrnrJr/s r8_linalg_slogdetrpsa)*1.6 GGE ;;uSbz "D E#2Joagg.F GI    *5%8ggxx  B [[s5;;[ 7F B &&r: 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^)rrqrrrsrrfrrHU_shaperV_shapeVrSs r8_linalg_svd_metarys#!\"1l+aggcrl#J  A  A Aq A 11== KK  g:7eTU]1== KK  a.F* g:7gVW KK  KK  J!$OAGG,D EA a7Nr:arg1arg2c|jdd}|jdd}t||}t|}||jd|jdgz }t|}||jd|jdgz }||fS)Nrr)rr%rr)rzr{arg1_batch_sizesarg2_batch_sizesexpand_batch_portionarg1_expand_sizearg2_expand_sizes r8r?r?s zz#2zz#2,-=?OP012 " 66012 " 66 - --r:c|r t|||t||\}}||jk(r|n|j|}||jk(r|n|j|}||fSrD)rr?rexpand)rzr{rrrarg1_broadcastedarg2_broadcasteds r8r r &sv  tT40)EdD)Q&&!DJJ.DKK@P4Q!DJJ.DKK@P4Q - --r:rc|jdd}|jdk(xs-|jdz |jk(xr|j|k(}|S)Nrr3)rr)rrexpected_batched_rhs_shape vector_cases r8linalg_solve_is_vector_rhsr:sS!&Sb!1**/ Q%**$R8R)R r:)rrrrJr/r0ctdtjjjk(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 rQr8r)rrsr8rez"_linalg_solve_ex..Qs%Ywwiqwwix 9r:rc 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)r4r4r:r8rez"_linalg_solve_ex..\s  Kr:r'Frrc3$K|]}|du ywrDr4)rrr=s r8rtz#_linalg_solve_ex..ss &Q1D= &s) copy_fromcopy_to exact_dtype)rrVrgr^r unsqueezerr?r(rrrrrCallzipr!rrr#)rrrrrrJr/r0rB_ B_broad_shaperR result_shaperesult_rLU_pivots_info_rresraos`` r8_linalg_solve_exrBs1n- LL 177 -Q2K'RQBaT>23B:M1 LL K *5="%-L!! *.s))$7 9r:c$djdS)NzMtorch.triangular_solve: Expected A to have at least 2 dimensions, but it has rrrsr8rez'triangular_solve_meta..r6r:triangular_solveFrr'rcy)Nz+triangular_solve: Got an unexpected layout.r4r4r:r8rez'triangular_solve_meta..rmr:)rVrgrrrstridedr?r(rr^r sparse_csr sparse_bsrrr) rrrrrself_broadcast_sizeA_broadcast_sizesolutioncloned_coefficients `` r8triangular_solve_metarsL LL Q  LL !  4$67xx5== 0LTST0U--&&$./BeT**;;   #00!./?5Q''88   ' '' U%% %U5E5E)E##D)!^^QC0 ' ''  UQR ' ''r:clt|dt|d|j|jdd}|j|j}|j |jt |jd|j|jddt j}|||fS)Nz linalg.detrFrrrr,)rdetrJr/s r8_linalg_det_metarsa&1l+ ++aggcrl #C QWW BNN17775QR [["U[[[ 9F F?r:c0tjjdk\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.r4r4r:r8rezormqr..rmr:cy)Nz3torch.ormqr: other must have at least 2 dimensions.r4r4r:r8rezormqr..rmr:rrcddS)Ntorch.ormqr: other.shape[z0] must be greater than or equal to tau.shape[-1]r4left_size_conditionsr8rezormqr..s+,?+@@pqr:cddS)Nrz"] must be equal to input.shape[-2]r4rsr8rezormqr..s+,?+@@bcr:cy)NzHtorch.ormqr: tau.shape[-1] must be less than or equal to input.shape[-1]r4r4r:r8rezormqr..rmr:r3c<djdjS)Nz[torch.ormqr: Expected tau to have one dimension less than input, but got tau.ndim equal to rrr sr8rezormqr..r!r:c<djdjS)Nzhtorch.ormqr: Expected other to have the same number of dimensions as input, but got other.ndim equal to rrrrsr8rezormqr..s+++0::,6RSXS]S]R^ `r:cdS)NzWtorch.ormqr: Expected batch dimensions of tau to be equal to input.shape[:-2], but got r4r#sr8rezormqr.. r%r:cdS)NzYtorch.ormqr: Expected batch dimensions of other to be equal to input.shape[:-2], but got r4)actual_batch_other_shapesr8rezormqr..s66N5OQr:c<djdjS)NzPtorch.ormqr: Expected input and tau to have the same dtype, but input has dtype z and tau has dtype rr sr8rezormqr..s'##(;;-/B399+ Or:c<djdjS)NzRtorch.ormqr: Expected input and other to have the same dtype, but input has dtype z and other has dtype rrsr8rezormqr.."s'##(;;-/DU[[M Sr:z torch.ormqrrrFrr') rVrgrrr^rr(rr) rrrrrexpected_batch_shaperr$rs ``` @@@r8ormqrrs 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||  r:crtjtdzk(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)rpaddingsr8rez,_padding_check_valid_input..6s1!c'+c'l^Tr:r3rc:ddzddzdjS)N Expected r3zD or rzcD (batch mode) tensor with possibly 0 batch size and other non-zero dimensions for input, but got: r9)rrsr8rez,_padding_check_valid_input..Ks2aycAgY/AAF  Or:)rVrgrrrr)rrr input_dim is_batch_modevalid_batch_modevalid_non_batch_moders``` r8_padding_check_valid_inputr3s LL G CT  I#'*M$,,q)$ GA/FEJJqMQ4F  Gy! OA#7#NEJJqMQ.g2%%*G2eWOE7*UZU`U`Tacr:cddS)Nz input (W: z%) is too small. Calculated output W: r4)input_woutput_wsr8rez_pad1d_common..os*WI%J8*Ur:r)rrrrVrgr) rr is_reflection dim_planenbatchnplanerrrrrs ` @@@@@r8 _pad1d_commonrRsI E F zzQA  Q ug15LE5 ZZ "FjjG&H GO /   LLA U  zzQ1229::r:ct||dSNTr)rrrs r8meta_reflection_pad1drx t <. =ekk>Q>Q>S=TTUXr:Fr)rVrgr^boolrrs` r8meta_replication_pad1dr~5 LL uzz!X u ==r:cd|s#tjt|dk(djdk(rdz |\j }|zz|r&tj|kxr|kfdtjj k(fdj j S)Nr3rcy)Nz padding size is expected to be 2r4r4r:r8rez(_pad1d_backward_common..rmr:r0c4ddddjSrr9rsr8rez(_pad1d_backward_common..rr:c2ddjSNz(grad_output width unexpected. Expected: , Got: rr grad_outputrsr8rez(_pad1d_backward_common..":8*GKL\L\]bLcKder:)rVrgrrrrr) rrrrrrrrrs `` @@@@r8_pad1d_backward_commonrs E  S\Q&(RS zzQ  LE5jjG&H GO /   LLK$$U++e ??5;; ''r: grad_inputc t|||dSrrrrrs r8meta_reflection_pad1d_backwardrs "+ugT RRr:c t|||dS)NFrrrs r8meta_replication_pad1d_backwardrs "+ugU SSr:c 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) Nrr3rrc4ddddjSrr9rsr8rez_pad2d_common..rr:c4ddddjSNzcArgument #6: Padding size should be less than the corresponding input dimension, but got: padding (rxrrr9dim_hrpad_bpad_tsr8rez_pad2d_common..rr:c ddddS)Nz input (H:  W: z%) is too small. Calculated output H: r4)input_hroutput_hrsr8rez_pad2d_common..s* gY/$$,:T( =r:r0rrrrVrgr)rrr dim_slicesrrrrrrrrrrrrrs` @@@@@@@@@@r8 _pad2d_commonrsU E EJ Fug15 ::D qyA    a !(E5% ZZ #FjjGjjG&H&H GO /    GO /   LLA &Q  zzQ(;<<(CDDr:ct||dSr)rrs r8meta_reflection_pad2drrr:ctjjtjk7fdt |dS)Nc@djjdS)Nz)"replication_pad2d" not implemented for 'rrrsr8rez(meta_replication_pad2d..rr:Fr)rVrgr^rrrs` r8meta_replication_pad2drrr:c^tj|}tj|}||fSrDrVr)grad_wsaved_vsaved_g saved_normsrgrad_vgrad_gs r8meta_weight_norm_backwardr%s.  g &F   g &F 6>r:c d d d}|j}|jdk(r dz dz |dz }|\}}}}| } | } | |z|z | |z|ztjj k( fdtj j k( fd|j |jS)Nrr3rr c2ddjSrrrsr8rez%meta_pad2d_backward..rr:c2ddjSNz)grad_output height unexpected. Expected: rrrrrsr8rez%meta_pad2d_backward..!";H:W[M]M]^cMdLefr:)rrrVrgrr)rrrrrdrrrrrrrrrrs` @@@@r8meta_pad2d_backwardr,s E EIJ xxzQ    Q !(E5%GG&H&H LLK$$U++e LLK$$U++f >>$** %%r:c $ 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) Nr0rr3rrc4ddddjSrr9rsr8rez_pad3d_common..Crr:c4ddddjSr r9r sr8rez_pad3d_common..Jrr:c4ddddjS)NzcArgument #8: Padding size should be less than the corresponding input dimension, but got: padding (rxrrr9)dim_drpad_bkpad_fsr8rez_pad3d_common..Qs2%%*G2fX_UG:V[VaVaUbdr:c ,dddddd S)Nz input (D:  H: rz%) is too small. Calculated output D: r4)input_drroutput_drrsr8rez_pad3d_common..Ys7 gYd7)<$$,:T(4z Kr:r)rrrr batch_moderrr2rrr7rrr8rrrr3r4rrrs` @@@@@@@@@@@@@@@r8 _pad3d_commonr:&s E E EIug15qJA      Q 07-E5%v ZZ "FjjGjjGjjG'H&H&H GO /    GO /    GO 0 0   LLA 7Q7(a-  (HMNN(HEFFr:ct||dSr)r:rs r8meta_reflection_pad3dr<err:ctjjtjk7fdt |dS)Nc@djjdS)Nz)"replication_pad3d" not implemented for 'rrrsr8rez(meta_replication_pad3d..prr:Fr)rVrgr^rr:rs` r8meta_replication_pad3dr?krr:cT tjt|dk(d|jdkrt d|jj|jk7r%t djd|jdd d |jd k(rdz  dz dz |\}}}}}}|j } |j } |j } | |z|z| |z|z| |z|ztjj k(fd tjj k( fd tjj k( fd |j |jS) Ncy)Nz padding size is expected to be 6r4r4r:r8rez%meta_pad3d_backward..rmr:r0zinput.ndim must be > 3, got z,grad_output.ndim must equal input.ndim, got  != rr3r.c2ddjSrrrsr8rez%meta_pad3d_backward..rr:c2ddjSr)rr*sr8rez%meta_pad3d_backward..r+r:c2ddjS)Nz(grad_output depth unexpected. Expected: rr)r2rr8sr8rez%meta_pad3d_backward..rr:)rVrgrrrrrr)rrrrrrrr4r3r7rrr2rrr8rrs` @@@@@@r8meta_pad3d_backwardrGus LLW"$NO zzQ;EJJ<HII5::%:;;K;K:LDQVQ[Q[P\ ]   E E E zzQ      07-E5%vjjGjjGjjG'H&H&H LLK$$U++e LLK$$U++f LLK$$U++e ??5;; ''r:pcJtj|jd|jd}|dkr0|j dgj tj S|j ||dz zdzfj tj S)Ncy)Nz(_pdist_forward requires contiguous inputr4r4r:r8rez%meta__pdist_forward..rmr:rr3rr)rVrgrdrrr'r)rrHrs r8meta__pdist_forwardrKs LL P ! AAv~~qc"%%E4R4R%SS~~qAE{a/1255886  r:gradpdistctj|jdtj|jdtj|tjS)Ncy)Nz._pdist_backward requires self to be contiguousr4r4r:r8rez&meta__pdist_backward..rmr:cy)Nz/_pdist_backward requires pdist to be contiguousr4r4r:r8rez&meta__pdist_backward..rmr:r)rVrgrdrr)rLrrHrMs r8meta__pdist_backwardrQsW LL V LL X   D0N0N OOr:)rVrUc 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_eqr3rr0cyNzbatch1 must be a 3D tensorr4r4r:r8rezmeta_baddbmm..rmr:cyNzbatch2 must be a 3D tensorr4r4r:r8rezmeta_baddbmm..rmr:cVdjdjdjS)Nz+Input dtypes must be the same, got: input: z , batch1: z , batch2: r)batch1batch2rsr8rezmeta_baddbmm..s0A$**ZX^XdXdWeeopvp|p|o}~r:c .ddddddd SNz@Expected size for first two dimensions of batch2 tensor to be: [rxz ] but got: [rr3].r4 batch2_sizesbscontraction_sizesr8rezmeta_baddbmm..s:t2&'|LO3DB|TUFWWY [r:)r$rSrTrrVsym_notrrrgr exp_config&skip_dtype_check_in_meta_registrationsr^r)rrZr[rVrUrSrTdim1dim2dim3 batch1_sizesr`rarbs``` @@@r8 meta_baddbmmrjs7L ;;q>D ;;q>D ;;q>DU]]6$**tT46H#IJK{{D$-. LL"$HI LL"$HI  < < JJ&,, 6&,, 6 ~ <>$))+ &&r:cLtj|tjSrrmrrs r8meta_bernoullirms   D0G0G HHr:c|SrDr4rrHrs r8meta_bernoulli_rp Kr:cLtj|tjSrrmros r8meta_bernoulli_prss   D0G0G HHr:c,tj|SrDrrls r8 meta_poissonru   D !!r:c tj| |jkdtj|tj} tj|| fS)Ncy)NzJError in fused_moving_avg_obs_fake_quant_cpu: ch_axis must be < self.dim()r4r4r:r8rez6meta__fused_moving_avg_obs_fq_helper.. rmr:r)rVrgrrr)r observer_on fake_quant_on running_min running_maxscale zero_pointaveraging_const quant_min quant_maxch_axisper_row_fake_quantsymmetric_quantmasks r8$meta__fused_moving_avg_obs_fq_helperrsO  LL$((*\   D 3D   T "D ))r:c@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 2Dr4r4r:r8rezmeta_mm.. rmr:cy)Nz b must be 2Dr4r4r:r8rezmeta_mm.. rmr:c "ddddd S)Nz/a and b must have same reduction dim, but got [rxz] X [r^r4)M1M2NrCsr8rezmeta_mm.. s(A!Brd%PRtSUVWUXXZ[r:cy)NzFout_dtype must be the same as input dtype or fp32 for fp16/bf16 inputsr4r4r:r8rezmeta_mm..% rmr:r) rVrgrrr^r`rarbr)r6r7r=rSrrrrCs @@@@r8meta_mmr s LLA56 LLA56 GGEAr GGEB LL b[   U]]*?GG u~~>> \  (/177YL ;;1v\; 22r:c|r(tfdtjDStjj S)Nc3HK|]}|vrj|ndyw)r3Nr9)rrrdimsrs r8rtz+_compute_reduction_shape..- s$UqatmTZZ]:Us")rfrrrOcompute_reduction_output_shaper)rrrs`` r8rr+ s7UE$))DTUUU  / / D AAr:ct|tjjr|jj St |drEt |jdr/|jj dk7r|jj Sy)Nrrzrr)rprV _subclasses FakeTensor fake_devicerzhasattrr)r0s r8rr6 sg&%++667!!&&&! FMM6 * MM  & (}}!!!r: input_tensorrrdilation is_transposedgroupsoutput_paddingcdtdtdtdtdtdtf 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 }d dl m }t||r |jn |j}|jdk(xrtj j"du}|s1tj$|d dDcgc]}|d kD c}fdScc}w)NlnrHrrHrr@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 rr3r4)rrHrrHrs r8_formulaz+calc_conv_nd_return_shape.._formulaM s.QU Q!a%[(1,2Q66r:rEc<|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 r3rr4)rrHrrHrrEs r8_formula_transposedz6calc_conv_nd_return_shape.._formula_transposed^ s2"Q!|a!e#a1q5k1B6::r:rr3rzInvalid channel dimensions)r)sym_orrc.dtddddS)NzGiven input size per channel: z&. Calculated output size per channel: rz. Output size is too small)r)r ret_shapesr8rez+calc_conv_nd_return_shape.. s,4T$ZLA33   A %#1!#4"5D "A "0  3t9     #GAJQKN1I'*     a'!*hqk;q>6RS9U %69< lJ /     {{f$B):):d)BH  IabM2qQU2 3 (   3s K& cbtjj|tjk(SrDrV _prims_commonr channels_lasttens r8is_channels_lastr s$    4 4S 9U=P=P PPr: running_mean running_vartrainingexponential_average_factorepsiloncrj}| |jn |j} | |jn |j} fd} j|j| } |r#j| } j| }n"jd} jd}| | |fS)NctrtjSjtjrtjStjSr)rrVrrdr)rsr8pick_memory_formatz2meta_miopen_batch_norm..pick_memory_format sF L )&& &  % %E4K4K % L** *&&&r:rr)rrr')rr9r;rrrrrrsave_mean_shapesave_var_shaperr save_meansave_vars` r8meta_miopen_batch_normr s""I-9,Dl((&,,O*5*A[&&v||N'   + . .=O=Q . RC **?; )).9 **40 ))$/  8 ##r:c t||||||||r|nd} d} d} |j| dk(rd| | <|j| } | SNr3r)rrr) rr9r;rrrrrr shape_outinput_channels_dimoutput_channels_dimrs r8 meta_convr sl*'T I+,1)* %&   +C Jr:mkldnnc t|||||d|g} |j| } tj} |j dk(rtj } | j | } | S)NFr.r)rrrVrrchannels_last_3dr') rr9r;rrrrattrscalars algorithmrrout_memory_formats r8meta_mkldnn_convolution_defaultr sp. &&'8UFB $$Y/!//     " % 6 6 ff#4f5 r:cb|jg|jdd|jdSNrrrr)rr9r;rrrs r8meta_linear_pointwise_defaultr. s5%%&Q (:(:3B(?&Qa&QRRr:mklcb|jg|jdd|jdSrr)r packed_weight orig_weightr;rs r8meta_mkl_linearr9 s:))@,$$Sb)@;+<+ * Ec)nEUUVW &&""%%  i.  ff6f* r:c.|dk7rtd|d|S)Nsumz#binary_op_name must be 'sum', got 'rr)r=rrrrraccumr;rrrrrrr accum_scaleaccum_zero_pointbinary_op_namerU unary_op_name unary_op_argsunary_op_algorithms r8meta_qconv2d_pointwise_binaryr| s-4 U " 5n5EQG  r:c 4t|j} |jd| d<| tjtjtj tj tjfvrtd| |j| | }|S)Nr3rzOoutput_dtype must be one of float32, bfloat16, int8, uint8, float8_e4m3fn, got r) rrrVr`rbrBrrcrr)r=rrrrrr;rrr post_op_name post_op_argspost_op_algorithmrgrs r8meta_qlinear_pointwiser s"AGG} 771: R  MM NN JJ KK      !abnaop kk,lk; r:cB| dk(r|St|j}|jd|d<| tjtjtj tj tjfvrtd| |j|| }|S)Nrr3rrr) rrrVr`rbrrBrcrr)r=rrrrrx_2r;rrrx2_scalex2_zprrUrrrrgrs r8meta_qlinear_pointwise_binaryr s, U "JAGG} 771: R  MM NN KK JJ      !abnaop kk,lk; r:cvt|j}|jd|d<|j|}|S)Nr3r)rrr)r=rr;rgrs r8meta_linear_dynamic_fp16r s6AGG} 771: Rkk,' r: quantizedrr3c.t||||||\}}}|jdk(r|jdnd} tj} |jdk(r|||g} n| |||g} tj | |j |j| SNr r3r0r~)#max_pool2d_checks_and_compute_shaperrrVrrr^r rrrrr ceil_mode nInputPlane outputHeight outputWidthrrrs r8meta_quantized_max_pool2dr  s 0 ;9    $)99;!#3B++ 99;! {;DK{CD{{ ++<<'   r:ctjjdk(fdtjjdk(fdtjjtjtj tj fvfdtjjtjk(fdtjjtjk(fdtjjjk(fdjjdjdj S) Nrc,djdS)Nzx must be a 2D tensor, got rMrr=sr8rez/meta_int4mm_packed_weight_cpu.. -H QR+Sr:c,djdS)Nzw must be a 2D tensor, got rMrrsr8rez/meta_int4mm_packed_weight_cpu.. rr:c"djSNz#expected x to be f32/f16/bf16, got rrsr8rez/meta_int4mm_packed_weight_cpu.. s9!''Cr:c"djSNzexpected w to be uint8, got rrsr8rez/meta_int4mm_packed_weight_cpu.. s.J177),Tr:c"djS)Nz q_group_size must be int64, got r) q_group_sizesr8rez/meta_int4mm_packed_weight_cpu.. s6|7I7I6JKr:c"djS)Nz5q_scale_and_zeros must have the same dtype as x, got r)q_scale_and_zerossr8rez/meta_int4mm_packed_weight_cpu..$ sKL]LcLcKder:rr) rVrgrr^r`rarbrrrrr=rrrs````r8meta_int4mm_packed_weight_cpur s QUUW\#ST QUUW\#ST GG u}}ennE E C   GGu{{ "$T     %++ - K    # #qww . e {{166!9affQiqww{??r:ctjjk(xrjk(fdy)NcjdddddjddjzS)NzExpected a tensor of dimension z and tensor.size[z] == rxzbut got : dimension z] = rr)rdim_sizerr0sr8rez check_dim_size..- sP1#6GzQVW[V\\^_ .?zfll[cNdMe fgr:)rVrgrr)r0rrrs````r8check_dim_sizer * s6 LL > X 6$ > gr:cd}|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)Nr3rcddS)Nz avg_pool2d: 4 must either be a single int, or a tuple of two intsr4rsr8rez1meta_avg_pool2d..unpack..? l4&(\]r:rr3rVrgrrrHWs` r8unpackzmeta_avg_pool2d..unpack< G H  ]  FSQACF!t r:rrr3rcyNzOavg_pool2d: stride must either be omitted, a single int, or a tuple of two intsr4r4r:r8rez!meta_avg_pool2d..H rmr:c@djjdS)Nz""avg_pool2d" not implemented for 'rrrsr8rez!meta_avg_pool2d..L 6u{{7J7J7L6MQQr:rr3rrcyNzdivisor must be not zeror4r4r:r8rez!meta_avg_pool2d..Y rmr:r rrrr0r~)rVrgrr^ruint16uint32uint64rrpooling_output_shaperOrpool2d_shape_checkrr)rrrrrcount_include_paddivisor_overrider,kHkWdHdWpadHpadWrr inputHeight inputWidthr r rrs` r8meta_avg_pool2drE2 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||#  r:ct|||||||dd| | | | | ||j}| }t|||dz |t|||dz | t|||dz | y)Nr3r0r)r:rr )r gradOutputrr=r>r?r@rArBrrCrDr r  mem_formatr nOutputPlanes r8avg_pool2d_backward_shape_checkrJ s"         $ 99;DL:tTAX|<:tTAX|<:tTAX{;r:ctjt|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)Nr3rcy)NzKavg_pool2d: kernel_size must either be a single int, or a tuple of two intsr4r4r:r8rez*meta_avg_pool2d_backward.. rmr:rcyr0r4r4r:r8rez*meta_avg_pool2d_backward.. rmr:cy)NzGavg_pool2d: padding must either be a single int, or a tuple of two intsr4r4r:r8rez*meta_avg_pool2d_backward.. rmr:cyr4r4r4r:r8rez*meta_avg_pool2d_backward.. rmr:r rr5rrr~) rVrgrrrr9rOrrJrr^r) gradOutput_rrrrrr;r<r=r>r?r@rArB input_sizerrrCrDr r rHs r8meta_avg_pool2d_backwardrR 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||  r:ctjt|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)Nr3r0cyNzFavg_pool3d: kernel_size must be a single int, or a tuple of three intsr4r4r:r8rez!meta_avg_pool3d.. rmr:rr3rcyNzJavg_pool3d: stride must be omitted, a single int, or a tuple of three intsr4r4r:r8rez!meta_avg_pool3d.. rmr:c@djjdS)Nz""avg_pool3d" not implemented for 'rrrsr8rez!meta_avg_pool3d.. r2r:cyNzBavg_pool3d: padding must be a single int, or a tuple of three intsr4r4r:r8rez!meta_avg_pool3d.. rmr:r r.cyNz9non-empty 4D or 5D (batch mode) tensor expected for inputr4r4r:r8rez!meta_avg_pool3d.." rmr:cyr4r4r4r:r8rez!meta_avg_pool3d..' rmr:rr5rrz avg_pool3d()T)check_input_sizer ) rVrgrr^rr6r7r8rrr9pool3d_shape_checkr)rrrrrr;r<kTr=r>dTr?r@padTrArBrnslicesitimeiheightiwidthotimeoheightowidths` r8meta_avg_pool3drl 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@AAHIIr:ctjt|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)NrTcyrVr4r4r:r8rez*meta_avg_pool3d_backward..a rmr:rr3rcyrXr4r4r:r8rez*meta_avg_pool3d_backward..i rmr:cyr[r4r4r:r8rez*meta_avg_pool3d_backward..q rmr:r\cyr^r4r4r:r8rez*meta_avg_pool3d_backward..y rmr:cyr4r4r4r:r8rez*meta_avg_pool3d_backward..~ rmr:rr5rrzavg_pool3d_backward()) rVrgrrrr9avg_pool3d_backward_shape_checkrr)rrrrrrr;r<rbr=r>rcr?r@rdrArBrerfrgrhotime_for_shape_checkoheight_for_shape_checkowidth_for_shape_checks r8meta_avg_pool3d_backwardrwS 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;; ''r:c,tjjdk(xsjdk(fdjddt |z}t j }tj|jj|S)Nr0r c"djS)Nz"Expected 3D or 4D tensor, but got r9rsr8rez*meta_adaptive_avg_pool2d.. 4TZZLAr:rr~) rVrgrrrfrOrrr^r)r output_sizergrs` r8meta_adaptive_avg_pool2dr| s| LL Q($))q.A::cr?U;%77L//5M ;;jj{{#  r:ctjjdk(xsjdk(fdjjddt |zS)Nr r.c"djS)Nz"Expected 4D or 5D tensor, but got r9rsr8rez*meta_adaptive_avg_pool3d.. rzr:r5)rVrgrrrrf)rr{s` r8meta_adaptive_avg_pool3dr sO LL Q($))q.A >>$**Sb/E+,>> ??r:cj}td|D].tjj dkDfd0tj|dk(xs|dk(fdtjj j k(fdtj }trtj}jjj|S) Nr3rc*djddS)Nz{adaptive_avg_pool2d_backward(): Expected grad_output to have non-zero size for non-batch dimensions,  with dimension being emptyr9)grad_outrsr8rez4meta__adaptive_avg_pool2d_backward.. s&66>nn5EEUVWUXXdfr:r0r c"djS)NzBadaptive_avg_pool2d_backward(): Expected 3D or 4D tensor, but got r9rsr8rez4meta__adaptive_avg_pool2d_backward.. sTUYU_U_T`ar:c<djdjSNexpected dtype z! for `grad_output` but got dtype r)rrsr8rez4meta__adaptive_avg_pool2d_backward.. s/$**-Nx~~N^_r:r) rrrVrgrr^rrrrrr')rrrrrs`` @r8"meta__adaptive_avg_pool2d_backwardr s ==D 1d^  MM! q  f   LL  TQYa LL hnn$_++M++ >>$** % ( (} ( EEr:cdt|dtj|tjS)Nadaptive_avg_pool3d_backwardr)!_adaptive_pool_empty_output_checkrVrrrrs r8"meta__adaptive_avg_pool3d_backwardr s(&k3QR   D0N0N OOr:rcj}td|D]/tjj dkDfd1y)Nr3rc.djddS)Nzc(): Expected grad_output to have non-zero size for non-batch dimensions, but grad_output has sizes rrr9)rrrsr8rez3_adaptive_pool_empty_output_check.. s/*--8->->,??OPQsR^`r:)rrrVrgr)rrrrs`` @r8rr sG   D 1d^    Q ! #   r:c 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) Nr0r c"djS)Nz:adaptive_max_pool2d(): Expected 3D or 4D tensor, but got: r9rsr8rez*meta_adaptive_max_pool2d.. LU[[MZr:r3rc*djddS)Nzjadaptive_max_pool2d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes rrr9rrsr8rez*meta_adaptive_max_pool2d.. %'',{{m3CA3lTr:rcy)NzCadaptive_max_pool2d(): internal error: output_size.size() must be 2r4r4r:r8rez*meta_adaptive_max_pool2d.. rmr:r r0rr) rrVrgrrrrrrOrr') rr{rdimHsizeBsizeDosizeHosizeWrrrrrs ` @r8meta_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|r:cNj}tj|dvfdtdtjjjk(fdt j }jjj|S)Nrc"djS)NzKadaptive_max_pooling2d_backward(): Expected 3D or 4D grad_output, but got: r9rsr8rez3meta_adaptive_max_pool2d_backward.. s]^i^o^o]pqr:adaptive_max_pool2d_backwardc<djdjSrr)rrsr8rez3meta_adaptive_max_pool2d_backward..% s!/%++.OP[PaPaObcr:r) rrVrgrr^rOrrrr')rrrrrs`` r8!meta_adaptive_max_pool2d_backwardr s   D LL q &k3QR LL {(((c //6M ??5;; ' * * * GGr:c 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) Nr\c"djS)Nz:adaptive_max_pool3d(): Expected 4D or 5D tensor, but got: r9rsr8rez*meta_adaptive_max_pool3d..2 rr:r3rc*djddS)Nzjadaptive_max_pool3d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes rrr9rsr8rez*meta_adaptive_max_pool3d..7 rr:r0cy)NzCadaptive_max_pool3d(): internal error: output_size.size() must be 3r4r4r:r8rez*meta_adaptive_max_pool3d..? rmr:r.r r)rrVrgrrrrr) rr{rdimDrrosizeTrrrrrrs ` @r8meta_adaptive_max_pool3dr, s ::D LL Z1d^  JJqMA     LL KAU D E E qy 1    JJt E(FFF qyFFF3 E666: //) $Cooiu{{o;G <r:cPt|d|j|jS)Nadaptive_max_pool3d_backward)rrr)rrrs r8!meta_adaptive_max_pool3d_backwardrX s"&k3QR ??5;; ''r:c>| td|j|S)Nz:cannot repeat_interleave a meta tensor without output_size)r&r)repeatsr{s r8meta_repeat_interleave_Tensorr_ s%WXX   [ ))r:c~|jjstd|j|jjstd|jt|j t |j|j t |jt j}|S)Nz!real must be floating point, got z!imag must be floating point, got rL)r^rrrTr'rrrP)realimagrs r8 meta_complexrf s :: ' '@ MNN :: ' '@ MNN  +DJJ78 +DJJ786>>F Mr:) fill_valuerct|dvr1|j||jftjStj ||jfd|ftj|j S)N)cpurrr3r^r)rrrrVrr(r)rrrs r8nonzero_staticru si4N*~~tTXXZ0 ~CC"" 488:  I**;;   r:c 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.r4r4r:r8reznonzero.. rmr:r3r) rV_check_not_implementedrdmeta_nonzero_assume_all_nonzeror(rrrrrs r8nonzeror sf   22 S    txxz" DJJLjj{{  r:c 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 providedr4r4r:r8rez#meta_index_Tensor.. rmr:cy)Nz?tensors used as indices must be long, int, byte or bool tensorsr4r4r:r8rez#meta_index_Tensor.. rmr:c"djS)N)too many indices for tensor of dimension rrsr8rez#meta_index_Tensor.. sG {Sr:c NdjddjdzS)NzThe shape of the mask at index z0 does not match the shape of the indexed tensor r9)rrjrHrsr8rez#meta_index_Tensor.. s<"8 ZPQsSJJN**U_`ade`e_f!hr:r3c<djdtdS)Nrz (got ry)rr)rrsr8rez#meta_index_Tensor.. s!;DII;fSQX\NZ[\r:rFrTczz}t|j}dgtz|tt|jtz |j ||S)zI This follows restride_src in TensorAdvancedIndexing.cpp r)rrrrrD)rrr after_shape before_shapereplacement_shapes r8 _restride_srcz(meta_index_Tensor.._restride_src so00;>t{{}%KL#PS Q K L!C Oc+6F$FGug..r:guard_or_false) rVrgr enumerater^rrrBrrrrrrrselect torch._refs_refsrr&rrr$rrrO3compute_elementwise_output_logical_to_physical_perm apply_permr invert_permrDr)rrrrrefsstatehas_contiguous_subspacertransposed_indicesrrrrrestrided_selfpermrR perm_shaper3rrrrrrHrs`` @@@@@@@r8meta_index_Tensorr 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 Jr:c >d} d} d} d}| dr8|||}|j|jj|} | dr8|||}|j|jj|} | dr|j|} | | | fS)Nc&t|}t|}|tjk(s|tjk(rtjS|tjk(s|tjk(rtjStjSrD)rrVrrr)t1t2fmt1fmt2s r8_conv_memory_formatz6meta_convolution_backward.._conv_memory_format9sr%R($R( 5&& &$%2E2E*E&& & 5)) )TU5K5K-K)) )&&&r:rrr3r)rrr') grad_output_input_weight_bias_sizes_optrrr transposedrr output_maskbackend_grad_inputbackend_grad_weightbackend_grad_biasrrs r8meta_convolution_backwardrs  '1~+L'B )33FKKMBEE'F 1~+FLA *44W\\^DGG'H 1~(22>B  35F GGr:cjd}jd}|j||f}tjj dk(dtjj dk(dtjjdjdk(fdtjjdjdk(fdtj|jd|k(xr|jd|k(d |j |jS) Nr3rr0cyrVr4r4r:r8rezmeta_addbmm..Zrmr:cyrXr4r4r:r8rezmeta_addbmm..[rmr:rcPdjddjdS)Nz8batch1 and batch2 must have same number of batches, got rrrrZr[sr8rezmeta_addbmm..^s.J6;;WX>JZZ_`f`k`klm`n_opr:c djddjddjddjdd S)Nz#Incompatible matrix sizes for bmm (r3r=rrryrrsr8rezmeta_addbmm..bsQ1&++a.1A6;;q>BRS;;q>"!FKKN#31 6r:cy)Nz.self tensor does not match matmul output shaper4r4r:r8rezmeta_addbmm..irmr:)rrrVrgrr)rrZr[rVrUrfrgs `` r8 meta_addbmmrTs ;;q>D ;;q>D ;;d| $D LL"$HI LL"$HI LL A&++a.(p LL A&++a.(  LL ! 51!5@ >>$))+ &&r:c @|j|jSrDrr)rrkwargss r8meta_randint_likerns >>$))+ &&r:) grad_scale found_infc n||||||fD])tjttfd+y)Nc dtSNz'exponent must be a tensor list but got rzlsr8rez#meta__fused_adam_..=d1gYGr:rVrgrpr)rgradsexp_avgs exp_avg_sqsmax_exp_avg_sqs state_stepslrbeta1beta2 weight_decayepsamsgradmaximizerrrs @r8meta__fused_adam_rss:&E8[/; O  q$  G  r:c ||||||fD])tjttfd+d}||||||||||fS)Nc dtSrrrsr8rez"meta__fused_adam..rr:cR|Dcgc]}tj|c}Scc}wrDr) tensor_listrs r8empty_like_listz)meta__fused_adam..empty_like_lists!-89  #999s$r)rrrr r r r r rrrrrrrrrs @r8meta__fused_adamrss&E8[/; O  q$  G  : ! $(  r:cjtjjdk(dtjjdk(dtjjtjufdtjjtjufdtjj dj dk(fdj j dj dftj S) Nrcy)Nza must be a 2D tensorr4r4r:r8rezmeta__int_mm..rmr:cy)Nzb must be a 2D tensorr4r4r:r8rezmeta__int_mm..rmr:c"djS)Nzexpected self to be int8, got r)r6sr8rezmeta__int_mm..0 :r:c"djS)Nzexpected mat2 to be int8, got r)r7sr8rezmeta__int_mm..rr:r3rc djddjddjddjdd S)Nz'Incompatible matrix sizes for _int_mm (rr=r3rryrr5sr8rezmeta__int_mm..sM5affQi[!&&)M66!9+Qqvvayk ,r:r)rVrgrr^rBrrrCr5s``r8 meta__int_mmr!s LLA>? LLA>? LL 5::: LL 5::: LL q QVVAY  ;;q 166!9-U[[; AAr:cftjjdk(dtjjtjufdj d}j ddz}j |dz||dzzd|dzftj S) NrcyNzw must be a 2D tensorr4r4r:r8rez2meta__convert_weight_to_int4pack..rmr:c"djSrrrsr8rez2meta__convert_weight_to_int4pack...qwwi8r:rr3 r)rVrgrr^rrrrCr inner_k_tilesrrHs` r8 meta__convert_weight_to_int4packr,s LLA>? LL 5;;8 q A q A A ;; F -"$ %  Q   kk  r:cJtjjdk(dtjjtjufdj d}j d}j ||dzftjS)Nrcyr$r4r4r:r8rez:meta__convert_weight_to_int4pack_for_cpu..rmr:c"djSNzexpected w to be int32, got rrsr8rez:meta__convert_weight_to_int4pack_for_cpu..r&r:rr3r)rVrgrr^rCrrrr*s` r8(meta__convert_weight_to_int4pack_for_cpur1s LLA>? LL 5;;8 q A q A ;; AF kk  r:ctjjdk(djjdk(rdndtjjk(fdtjj tj tjtjfvfdtjj tjufdjjdk(rjdnjdd z}jjd|j S) NrcyNzx must be a 2D tensorr4r4r:r8rez*meta__weight_int4pack_mm..rmr:rr cddS)Nz w must be a zD tensorr4) expected_dimsr8rez*meta__weight_int4pack_mm..sLh2Wr:c"djSrrrsr8rez*meta__weight_int4pack_mm..5aggY?r:c"djSr0rrsr8rez*meta__weight_int4pack_mm..r&r:rr'r) rVrgrrrzr^r`rarbrCrr)r=rrrdim_nr6s`` @r8meta__weight_int4pack_mmr;s LLA>? **e31L LLL(*WX LL EMM5==%..AA? LL 5;;8++u4AFF1I!&&)a-E ;;qvvay%qww; 77r:c(tjjdk(dtjjdk(dtjjtjtj tj fvfdtjjtjufdjjdjdjS)Nrcyr4r4r4r:r8rez2meta__weight_int4pack_mm_for_cpu..rmr:cyr$r4r4r:r8rez2meta__weight_int4pack_mm_for_cpu..rmr:c"djSrrrsr8rez2meta__weight_int4pack_mm_for_cpu..r8r:c"djSrrrsr8rez2meta__weight_int4pack_mm_for_cpu..r&r:rr) rVrgrr^r`rarbrrrrs`` r8 meta__weight_int4pack_mm_for_cpurA LLA>? LLA>? LL EMM5==%..AA? LL 5;;8 ;;qvvay!&&)177; ;;r:c(tjjdk(dtjjdk(dtjjtjtj tj fvfdtjjtjufdjjdjdjS)Nrcyr4r4r4r:r8rez;_weight_int4pack_mm_with_scales_and_zeros..rmr:cyr$r4r4r:r8rez;_weight_int4pack_mm_with_scales_and_zeros..rmr:c"djSrrrsr8rez;_weight_int4pack_mm_with_scales_and_zeros..r8r:c"djSr0rrsr8rez;_weight_int4pack_mm_with_scales_and_zeros..r&r:rr) rVrgrr^r`rarbrCrr)r=rrqScaleqZeross`` r8)_weight_int4pack_mm_with_scales_and_zerosrJ rBr:r6r7c||zdz |z|zSr2r4r5s r8 kai_rounduprLs UQY1  !!r:c |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)Nr r'r(rc8t||zd}t||S)Nr rL)rHkrsrkr_sr_roundedup4s r8kai_k_roundedupz3get_kai_packed_weight_size..kai_k_roundedup+s#$/rBw#: "1&677r:cf|||}|dzdk7rtd|||dzzzzzS)Nrrzk_internal must be even, got r) rHnrrPrQ k_internalrSkai_num_bytes_biaskai_num_bytes_multiplier_rhskai_num_bytes_sum_rhss r89kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0z]get_kai_packed_weight_size..kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4cxp_qsu4cxs1s01s`-QB7 Nq((+H )UVV1_23+,))r:c>t|||z}|||||zSrDrO)rrHrUrPrQnum_rowsrZs r87kai_get_rhs_packed_size_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0z[get_kai_packed_weight_size..kai_get_rhs_packed_size_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0@s6'q"-3O2r2r:r)rc||zdk7rtd|d|d| zdk7rtd|d d|zdk7rtd|ddt|||z}||||||zSNrbl (z) must be divisible by kr (ryznr (z+) must be divisible by kai_nr_multiple_of (+) must be divisible by kai_bl_multiple_of (rrL) rrHrUrPrQblr\kai_bl_multiple_of;kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0kai_nr_multiple_ofs r89kai_get_rhs_packed_size_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0z]get_kai_packed_weight_size..kai_get_rhs_packed_size_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0XsG>(4t3NrdRS)TUU++1(rd"MN`Maabc++1(rd"MN`Maabc'q"-3Q2r2rr:c||zdk7rtd|d|d| zdk7rtd|d d|zdk7rtd|dd } ||} ||}|||zz zzSr_r)rHrUrPrQrcnum_bytes_multiplier_rhsnum_blocks_per_rownum_bytes_per_blockrd#kai_get_bf16_datatype_size_in_bytesrfkai_num_blocks_per_rowrWkai_num_bytes_per_blockrYs r8rez_get_kai_packed_weight_size..kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0osG>(4t3NrdRS)TUU++1(rd"MN`Maabc++1(rd"MN`Maabc ,O+P(%;Ar%B"&=0'#(+==+,()r:cy)Nrr4r4r:r8rlzGget_kai_packed_weight_size..kai_get_bf16_datatype_size_in_bytessr:cV|zdk7rtd|ddt|||zS)Nrr`raryrb)rHrcrds r8rmz:get_kai_packed_weight_size..kai_num_blocks_per_rowsH++1(rd"MN`Maabc#1b)R//r:cH|zdk7rtd|dd|dz|zS)Nrr`raryrr)rcrirds r8rnz;get_kai_packed_weight_size..kai_num_bytes_per_blocksF++1(rd"MN`Maabca#;;;r:r4)n_bitsrK groupsizekai_nrkai_krkai_srr]rgrdrlrerZrSrfrmrWrXrnrYs @@@@@@@@@@@r8get_kai_packed_weight_sizerxs { >FFF$% !+, (!"  8  K1fff ^q Q]a%7FFF$% !!" !" !#  .  :  0 <M1fffi ]&8 _r:ctjjtjufdtjj j r||k(r|jtjk(s2||kre|dzdk(r]||zdk(rU|jtjk(r8td|||}jt|tjSj|jz}|||jz }j|tjS)Nc"djSrr)weightssr8rez2meta__dyn_quant_pack_4bit_weight..s.w}}o>r:r)rr r) rVrgr^rbackendskleidiai is_availablerZrbrxrrr)r{ scales_zerosr; block_size in_features out_featurespacked_weight_sizes` r8 meta__dyn_quant_pack_4bit_weightrs  LL $> ~~++- { "|'9'9U[['H  $R1$j(A-""enn48 |[*   %7!8  LL <+=+=+?? djjl*   /u{{  CCr:cntjjdk(dtjjtjk(xs$jtj k(xrk(fdj d}j||jS)Nrcy)Nzinput must be a 2D tensorr4r4r:r8rez-meta__dyn_quant_matmul_4bit..rmr:c.djddS)NzPexpected input to be f32 or bf16 (bf16 requires block_size == in_features), got z with block_size=z and in_features=r)rrinpsr8rez-meta__dyn_quant_matmul_4bit..s*99+.zl:KK= Zr:rr)rVrgrr^r`rbrr)rpacked_weightsrrrrs` `` r8meta__dyn_quant_matmul_4bitrs LLa!DE LL emm # G II ' EJ+,E   A ==L = ::r:c(tjjdk(dtjjtjtj tj fvfdtjjdk(dtjjtjufdjjdjdjS)Nrcyr4r4r4r:r8rez*meta__weight_int8pack_mm..rmr:c"djSrrrsr8rez*meta__weight_int8pack_mm..r8r:cyr$r4r4r:r8rez*meta__weight_int8pack_mm..rmr:c"djS)Nzexpected w to be int8, got rrsr8rez*meta__weight_int8pack_mm..s-aggY7r:rr) rVrgrr^r`rarbrBrr)r=rq_scaless`` r8meta__weight_int8pack_mmrs LLA>? LL EMM5==%..AA? LLA>? LL 5::7 ;;qvvay!&&)177; ;;r:cntjjdk\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: rMrx1sr8rez$meta_cdist_forward..CBFFH:QOr:c,djdS)Nz1cdist only supports at least 2D tensors, X2 got: rMrx2sr8rez$meta_cdist_forward..rr:rcPdjddjdS)Nz4X1 and X2 must have the same number of columns. X1: rz X2: r)rrsr8rez$meta_cdist_forward..s,Frwwr{mSXY[Y`Y`acYdXefr:c"djS)Nz3cdist only supports floating-point dtypes, X1 got: rrsr8rez$meta_cdist_forward..EbhhZPr:c"djS)Nz3cdist only supports floating-point dtypes, X2 got: rrsr8rez$meta_cdist_forward..rr:rcy)Nz)cdist only supports non-negative p valuesr4r4r:r8rez$meta_cdist_forward..rmr:)Nrr3rcdS)Nz(possible modes: None, 0, 1, 2, but was: r4) compute_modesr8rez$meta_cdist_forward..s:<.Ir:r) rVrgrrrOis_float_dtyper^rrbroadcast_shapesextendr) rrrHrr1r2 batch_tensor1 batch_tensor2rgs `` ` r8meta_cdist_forwardrsJ LL A O LL A O LL  rwwr{"f LL RXX&P LL RXX&P LLaLM LL'I B BHHSbMMHHSbMM..}mLMLR! << %%r:c4|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) rrrVrcopyrmathprod zeros_likerrr) rLrrrHcdistc1rrrrrtensor1_expand_size batch_products r8meta_cdist_backwardrs "B "B "BHHSbMMHHSbMM 6 6}m TU.335Bx(II23M Qw"'R1W (:##d288n, YY* +   Be.E.E FFr:c 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)rsr8rez$meta_embedding_bag..*:7==/Jr:c"djS)Nz(expected offsets to be long or int, got r)rwsr8rez$meta_embedding_bag...rr:c"djS)Nz/expected weight to be floating point type, got r)r9sr8rez$meta_embedding_bag..2sA&,,Pr:rr3cyNz1include_last_offset: numBags should be at least 1r4r4r:r8rez$meta_embedding_bag..9rmr:cy)Nz@embedding_bag: per_sample_weights only supported with mode='sum'r4r4r:r8rez$meta_embedding_bag..Brmr:c$djdS)Nz1expected per_sample_weights to be 1D tensor, got rMr)per_sample_weightssr8rez$meta_embedding_bag..FsGHZH_H_G``abr:cNdjdjdS)Nz%expected per_sample_weights.numel() (z$ to be the same as indices.numel() (ryr)rrsr8rez$meta_embedding_bag..Js478J8P8P8R7ST66=mmo5FaIr:cD|||xr|jddk(SNrr3r)r*r}r padding_idxis_fast_path_index_selects r8is_fast_path_index_select_scalez;meta_embedding_bag..is_fast_path_index_select_scalePs( %c6; ? XELLQROWXDX r:c|jtjk(xs|jtjk(xr1|j ddk(xr|j ddk(xr|dkSr)r^rVrZrXr)r*rrs r8rz5meta_embedding_bag..is_fast_path_index_selectUsb YY%++ % @ejj)@  1 "  a A% a  r:c2| ||||S|||SrDr4)r*r}rrrrs r8 is_fast_pathz(meta_embedding_bag..is_fast_path]s)  23v{S S,S&+F Fr:rcyrr4r4r:r8rez$meta_embedding_bag..wrmr:)rVrgr^rrrOrrrMODE_SUMrrrMODE_MAX MODE_MEANr)r9rrwscale_grad_by_freqr]sparserinclude_last_offsetrnum_bagsrr offset2bagbag_size max_indices fast_path_sumnumBagsrrs``` ` @@r8meta_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 ;;r:rct||d}tj|j|}t |||}|j ||S)NT)rr)rrOrrrr)rrrr^rrgs r8 meta_nansumrsI(u$OL    T 2D+E4AL ??<|? <.rmr:rcddS)Nz"Repeats cannot be negative, found rr4)rrepsr8rezmeta_repeat..s8ZsKr:r) rVrgrrrrfrrr)rrnum_new_dimensions padded_sizer target_sizers ` @r8 meta_repeatrs LL G  "lG$ 3 1H K  W 2++eDJJ.??K8=c'l8KL1;q>GAJ.LKL >>+ &&Ms1Cc|SrDr4rs r8 meta_zero_rrqr:czt|tjr t|j|j|SrD)rprVr rirrrs r8meta_binop_inplacers)%& EKK8 Kr:cd}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 SrD)rprrOr>r^rrss r8 is_integericz.meta_binop_inplace_alpha..is_integerics. c: &))#))4 4c7+ +r:ct|trtj|jSt|t SrD)rprrOrr^rrs r8 is_floaticz,meta_binop_inplace_alpha..is_floatics. c: &'' 2 2c9- -r:ct|trtj|jSt|t SrD)rprrOis_boolean_dtyper^rrs r8 is_booleanicz.meta_binop_inplace_alpha..is_booleanic s. c: &))#))4 4c8, ,r:z]Promotion of int.add/sub_(float) in in-place ops are not possible due to element size change.z_Promotion of book.add/sub_(others) in in-place ops are not possible due to element size change.)r&rprVr rir)rrrUrrrs r8meta_binop_inplace_alpharsz$, . -Dj/ k  D,u"5 m  %& EKK8 Kr:c:t||tjSNrrTrrPrrrUs r8meta_binop_alphar"s  e$C$K$K r:c 8t|tjSrr)rrs r8 meta_roundr.s  <DD r:cltjtjjfdt tj r8tjtjjfdytjt tfdy)Nc&djS)Nz7: Expected input tensor to have an integral dtype. Got r)rrsr8rez#shift_dtype_check..8s7)RSWS]S]R^_r:c&djS)Nz6: Expected shift value to have an integral dtype. Got rrrsr8rez#shift_dtype_check..=swiUVYV_V_U`ar:cdS)Nz): Expected shift value to be an int. Got r4rsr8rez#shift_dtype_check..BswiHNr:)rVrgrOr>r^rpr r)rrrs```r8shift_dtype_checkr5sp LL tzz*_#u||$  " "399 - a  sG $ N r:cTtd||t||tjS)NrshiftrrrTrrPrs r8 meta_rshiftsr F)he,  e$C$K$K r:cTtd||t||tjS)Nlshiftrr rs r8 meta_lshiftsrNr r:c8|j|jSrDrrs r8 meta_zerorVs >>$** %%r:c|SrDr4rrs r8 meta_fill_r[rqr:c,tj|SrDrrs r8 meta_fillr`   D !!r:c|SrDr4rs r8 meta_relu_rerqr:c:t||tjSrrrs r8meta__add_relurjs  e$C$K$K r:c,tj|SrDrrnoiselowerrrrs r8meta_rrelu_with_noiser rs   D !!r:cVtj|tj|fSrDrrs r8 meta_rrelu_with_noise_functionalr"zs%   D !5#3#3E#: ::r:c|SrDr4)rrrrrs r8meta_rrelu_with_noise_r$s  Kr:c,tj|SrDrrrr accumulates r8meta_index_putr(rr:cFt|j|j|SrDrir)rrvalues r8meta_masked_fill_r,sDJJ 3 Kr:c|j|jjtj|}|Sr)rrr'rOr)rrr} masked_scales r8meta__masked_scaler/s<>>$))+.1111$72L r:ctj|jtjtjfvdtjjjk(fdS)Ncy)NzMask must be bool or uint8r4r4r:r8rez&meta_masked_scatter_..rmr:c<djdjS)NzEmasked_scatter: expected self and source to have same dtypes but got rr)rrjsr8rez&meta_masked_scatter_..s"**U6<<.:r:)rVrgr^rr)rrrjs` `r8meta_masked_scatter_r3sU LL uzz5;;//1U LL fll" : Kr:ct||\}}tj|tj}t |||Sr)r&rVrrr3)rrrjrs r8meta_masked_scatterr5s;"$-JD$   d%2I2I JF f 55r:c$|j|SrDr)rrrGs r8meta_masked_scatter_backwardr7s >>%  r:c|SrDr4r&s r8meta_index_put_r9rqr:c  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_andrTr0cyrVr4r4r:r8rez)common_meta_baddbmm_bmm..rmr:cyrXr4r4r:r8rez)common_meta_baddbmm_bmm..rmr:rr3c .ddddddd Sr]r4r_sr8rez)common_meta_baddbmm_bmm..s5RSURV  l<?*;2l1o=NbRr:cy)Nzfout_dtype only supported for torch.float32 output with float16/bfloat16 inputs or same as input dtypesr4r4r:r8rez)common_meta_baddbmm_bmm..rmr:cy)Nzself must be a 3D tensorr4r4r:r8rez)common_meta_baddbmm_bmm..rmr:c0ddjS)Nz*Expected an input tensor shape with shape z but got shape: r)r{ self_baddbmmsr8rez)common_meta_baddbmm_bmm..s!@ M]^j^o^o^q]rsr:) r$r;rTrVrgrrr^rarbr`rr')rZr[is_bmmrBr=r;rTrires_rowsres_colssupported_out_dtyperr`rarbr{s ` @@@@r8common_meta_baddbmm_bmmrGsE 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 Mr:ct||dS)NTrG)rrJs r8meta_bmmrJs "4t 44r:c t||d|S)NT)r=rI)rrJr=s r8meta_bmm_dtyperLs "4ty IIr:ch||z}||z}|dk7r"t|dkt|dkk7r|dz}|Sr)r)r=r>qras r8div_rtnrOsB QA AA Av4A;$q1u+- Q Hr:ct||z|z||dz zz dz |r|dz ndz|dz}|r|dz |z||zk\r|dz}|Sr)rO) inputSize kernelSizerrrrr outputSizes r8pooling_output_shape_pad_lrrTs    *q.) *   'vzA  /     Nf $ E(9 9 !OJ r:c tj|dk7dtjdk\fdtjdz zdzdzkfdt|||S)Nrcy)Nzstride should not be zeror4r4r:r8rez&pooling_output_shape..rmr:cdS)Nz'pad must be non-negative, but got pad: r4padsr8rez&pooling_output_shape..s%LSE#Rr:r3rcdddS)NzApad should be at most half of effective kernel size, but got pad=z, kernel_size=z and dilation=r4)rrRrYsr8rez&pooling_output_shape..s'OPSuU%,nXJ @r:)rVrgrT)rQrRrYrrrs `` ` r8r9r9ss LL1AB LLRS LL a8+a/A55  ':sC9 r:c \ 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: r4)r=r>sr8rez$pool2d_shape_check..9sGt6RTQUVr:cddS)Nz0stride should be greater than zero, but got dH: , dW: r4)r?r@sr8rez$pool2d_shape_check..=sB2$fRDQr:cddS)Nz9dilation should be greater than zero, but got dilationH: , dilationW: r4) dilationH dilationWsr8rez$pool2d_shape_check..AsKI;Vcdmcnor:r3rr r0c*djS)NzExpected 4D (batch mode) tensor expected for input with channels_last layout with optional 0 dim batch size for input, but got: rrsr8rez$pool2d_shape_check..IsCCH::<.Rr:c*djS)NzYExpected 3D or 4D (batch mode) tensor with optional 0 dim batch size for input, but got: rrsr8rez$pool2d_shape_check..Psopupzpzp|o}~r:c ddddS)NzKpad should be smaller than or equal to half of kernel size, but got padW = z , padH = z, kW = z, kH = r4)r=r>rArBsr8rez$pool2d_shape_check..Us&ygbT>r:c .ddddddd SNzGiven input size: (r=z). Calculated output size: (z). Output size is too smallr4)rCrDrrIr r sr8rez$pool2d_shape_check..[s:%k]!K=*N$$0><.+O##r:)rrVrgrr)rr=r>r?r@rArBrbrcrrCrDr r rr valid_dimsrIs`````````````` @r8r: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. # #r:rerbr=r>rcr?r@pTpHpW dilationTrbrcrfrgrhrirjrkr`c 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: r]r4)r=rbr>sr8rez$pool3d_shape_check..}s#$fRDrd ,r:cdddS)Nz0stride should be greater than zero, but got dT: z, dH: r_r4)r?rcr@sr8rez$pool3d_shape_check..s >rd&FSURV Wr:cdddS)Nz9dilation should be greater than zero, but got dilationT: z , dilationH: rar4)rbrmrcsr8rez$pool3d_shape_check..s$#M)M) Vr:r\c&djS)Nz/: Expected 4D or 5D tensor for input, but got: r9)rrsr8rez$pool3d_shape_check..s7)J5;;-Xr:r.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)rrrsr8rez$pool3d_shape_check..s3)--2[[M+EJJqM?:KMr:c .ddddddd S)Nzinput image (T: r6rz ) smaller than kernel size (kT:  kH:  kW: ryr4)rgrfrhr=rbr>sr8rez$pool3d_shape_check..s9"5'gYd6(C$$&4uRDbT<r:rc ,dddddd S)NzHpad should be smaller than or equal to half of kernel size, but got kT: rvruz padT: z padW: z padH: r4)r=rbr>rkrjrlsr8rez$pool3d_shape_check..s6$eB4uRDt72$gbT Kr:r3c:dddddddddSrhr4)rgrfrhrerjrirksr8rez$pool3d_shape_check..sI!'!E7!G9AfXF((/y%'!F8L' (r:)rrVrgrr)rrerbr=r>rcr?r@rjrkrlrmrbrcrfrgrhrirjrkrr`rrs````````````````````` @r8raraas0 ::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!|  r:cj|j}t|||||||| | | | | |||||||||t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |t|||dz |y)Nr r0rr3rrar )rrrrerbr=r>rcr?r@rjrkrlrmrbrcrfrgrhrirjrkrrs r8max_pool3d_backward_shape_checkr{s2 ::D            +0;dQh8;dQh6;dQh8;dQh77D$(G47D$(E27D$(G47D$(F3r:c|j}t||||||||| | | ddd| | |||||dt|||dz |t|||dz |t|||dz |t|||dz |y)Nr3Tr r0rrz)rrrerbr=r>rcr?r@rjrkrlrfrgrhrirjrkrrs r8rsrss* ::D                -2;dQh8;dQh6;dQh8;dQh7r:cd}|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)Nr#cddS)Nz max_pool2d: r%r4r&sr8rezEmax_pool2d_checks_and_compute_shape..unpack..=r'r:rr3r(r)s` r8r,z3max_pool2d_checks_and_compute_shape..unpack:r-r:rr.cy)NzOmax_pool2d: stride must either be omitted, a single int, or a tuple of two intsr4r4r:r8rez5max_pool2d_checks_and_compute_shape..Grmr:rrrrr5rrr cy)NzMnon-empty 4D (batch mode) tensor expected for input with channels_last layoutr4r4r:r8rez5max_pool2d_checks_and_compute_shape..Xrmr:rcy)Nz9non-empty 3D or 4D (batch mode) tensor expected for inputr4r4r:r8rez5max_pool2d_checks_and_compute_shape..]rmr:Fcy)NzAUnsupported memory format. Supports only ChannelsLast, Contiguousr4r4r:r8rez5max_pool2d_checks_and_compute_shape..brmr:) rVrgrrrOrrrrr9r:)rrrrrrr,r=r>r?r@rArBrbrcrrCrDrr r s r8rr1sM; /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 11r:c| t|||||\} tjjjk(fd| j fd} | | |t j } tjjjj| S)Nc<djdjS)NzExpected dtype z for `gradOutput` but got dtype rrsr8rez7meta_max_pool2d_with_indices_backward..s /$**-MkN_N_M`ar:clt|dz t|dz t|dz y)Nr0rr3)r )rrIrr r s r8_check_dim_sizez>meta_max_pool2d_with_indices_backward.._check_dim_sizes9q$q,7q$q,7q$q+6r:r~) rrVrgr^rrOrrrr)rrrrrrrrrrrrIrr r s`` @@@@r8%meta_max_pool2d_with_indices_backwardr}s , k67Hi    LL k'''a L 99D7 K G//5M ;; jj{{#  r:ct||||||\}}}|jdk(r|jdnd} tj|} |jdk(r|||g} n| |||g} t j | |j|j| t j | t j|j| fSr) rrrrOrrVrr^rrrs r8meta_max_pool2d_with_indicesrs , {FGXy   %yy{a/UZZ^QF//6M yy{a\;7 \;? ++<<'   ++<<'    r:ct 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: rrsr8rez,meta_fractional_max_pool2d..sLTYYKXr:r0rc2djddS)Nz_fractional_max_pool2d: Expected input to have non-zero size for non-batch dimensions, but got rz emptyr)rrsr8rez,meta_fractional_max_pool2d..s'77;yy{mCSTUSVV\^r:rcy)NzNfractional_max_pool2d: kernel_size musteither be a single int or tuple of Intsr4r4r:r8rez,meta_fractional_max_pool2d..rmr:cy)NzOfractional_max_pool2d: output_size must either be a single int or tuple of Intsr4r4r:r8rez,meta_fractional_max_pool2d..rmr:r5rrr r3cy)Nz6Expect _random_samples to have the same dtype as inputr4r4r:r8rez,meta_fractional_max_pool2d..rmr:c"djS)Nz1Expect _random samples to have 3 dimensions got, r)random_samplessr8rez,meta_fractional_max_pool2d..sCNDWDWCXYr:cy)Nz=Expect _random_samples.size(0) no less then input batch size.r4r4r:r8rez,meta_fractional_max_pool2d..rmr:cy)Nz.rmr:cddS)Nz/Expect _random_samples.size(2) equals to 2 got .r4)rsr8rez,meta_fractional_max_pool2d..s#RSTRUUV!Wr:cdddS)Nz%fractional_max_pool2d: kernel height rz' is too large relative to input height r4) input_heightrsr8rez,meta_fractional_max_pool2d.. s7 A7GGno{n|}r:cdddS)Nz$fractional_max_pool2d: kernel width r3z& is too large relative to input width r4) input_widthrsr8rez,meta_fractional_max_pool2d..s6{1~6FFlmxlyzr:r) rVrgrrrrr^rrrr) rrr{rrinput_channels input_batchrcrrrrs `` ` @@@r8meta_fractional_max_pool2drss 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?  **;;  ++;;  r:ctjt|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(r||||f}n|||||f}|j|}|j|tj}|r@|jtj}|jtj}||fS)NrTcyNzMmax_pool3d: kernel_size must either be a single int, or a tuple of three intsr4r4r:r8rez.meta_max_pool3d_with_indices..1rmr:rr3rcyNzQmax_pool3d: stride must either be omitted, a single int, or a tuple of three intsr4r4r:r8rez.meta_max_pool3d_with_indices..9rmr:cyNzImax_pool3d: padding must either be a single int, or a tuple of three intsr4r4r:r8rez.meta_max_pool3d_with_indices..Armr:cyNzJmax_pool3d: dilation must be either a single int, or a tuple of three intsr4r4r:r8rez.meta_max_pool3d_with_indices..Irmr:r\cyr^r4r4r:r8rez.meta_max_pool3d_with_indices..Qrmr:r.rr5rrzmax_pool3d_with_indices()r rr) rVrgrrrr9rarOrrrrr')rrrrrrrbr=r>rcr?r@rjrkrlrmrbrcrrerfrgrhrirjrkrrrrs r8meta_max_pool3d_with_indicesr%s 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            #+4  aXE77>%BXBXX zzQeWf5 WeWf= //) $Cooiu{{o;Gff5#9#9f:**5+A+A*B <r:ctjt|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(}|j|j}|r |jtj}|S)NrTcyrr4r4r:r8rez7meta_max_pool3d_with_indices_backward..rmr:rr3rcyrr4r4r:r8rez7meta_max_pool3d_with_indices_backward..rmr:cyrr4r4r:r8rez7meta_max_pool3d_with_indices_backward..rmr:cyrr4r4r:r8rez7meta_max_pool3d_with_indices_backward..rmr:r\cyr^r4r4r:r8rez7meta_max_pool3d_with_indices_backward..rmr:rr5rrz"max_pool3d_with_indices_backward()r.r) rVrgrrrr{rOrrrrr')rrrrrrrrrbr=r>rcr?r@rjrkrlrmrbrcrerfrgrhrirjrkrrs r8%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#            ,/8  aXE77>%BXBXX-J]]1G1G]H r:gridcztjjjk(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 rrrsr8rez+check_grid_sampler_common..s'\\N"24;;- Ar:c<djdjS)NzTgrid_sampler(): expected input and grid to have torch.strided layout, but input has z and grid has )rrsr8rez+check_grid_sampler_common..s&nT[[M Cr:rc<djdjS)NzZgrid_sampler(): expected grid and input to have same batch size, but got input with sizes  and grid with sizes r9rsr8rez+check_grid_sampler_common..s' % },A$** Or:rrcBdjdz djS)Nz+grid_sampler(): expected grid to have size rz, in last dimension, but got grid with sizes )rrrsr8rez+check_grid_sampler_common..s,9%**q.9IJ226** ?r:c*djddS)NzYgrid_sampler(): expected input to have non-empty spatial dimensions, but input has sizes rrr9rsr8rez+check_grid_sampler_common.. rr:)rVrgrrrrrr)rrrs``@r8check_grid_sampler_commonrs LL  #  LL  %F$++*F  LL A$**Q-'  LL 2%**q.( 1ejj !  KKNQ    r:ceZdZdZdZdZy)GridSamplerInterpolationrr3rN)r{ __module__ __qualname__BILINEARNEARESTBICUBICr4r:r8rrsHGGr:rinterpolation_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 rr9rsr8rez'check_grid_sampler_3d..s&449KK=#DJJ< 1r:cy)Nz.&rmr:)rVrgrrrr+)rrrs`` r8check_grid_sampler_3drsp LL a3EJJ$))3  LL JJ!O M"&>&F&F&L&LL  O r:c|d}|r&tj|tj}nd}tj|tj} || fSNrr)rVrrr rrrr padding_mode align_cornersrinput_requires_gradr grad_grids r8grid_sampler_2d_backward_metar*sQ&a.%%e5;R;RS    U5L5LMI  ""r:ct||t||||jd}|jd}|jd}|jd}|jd} |j||||| fS)Nrr3rr0)rrrr) rrrrrrCout_Dout_Hout_Ws r8grid_sampler_3dr=sveT*%'9: AA AA JJqME JJqME JJqME ??Aq%6 77r:rct||t||||d}|r&tj|tj}nd}tj |tj} || fSr)rrrVrrrrs r8grid_sampler_3d_backwardrPsmeT*%'9:%a.%% !?!?    U5S5STI y  r:c|jd}|stj|}||d<tj|g|i|S)Nr^)r]rO get_dtyperVr)rrrQrr^s r8fullrhsC JJw E  +F7O ;;t -d -f --r:cN|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 tensorsr4r4r:r8rezzeros_like..rmr:rr Tr)rV sparse_coorgrr^r is_sparsesparse_resize_and_clear_r sparse_dim dense_dimr _coalesced_r-rr)fill_)rr^rrrrrs r8rrss!!! T ! O kk % $**5"(.4;;f!   >>  ( ( T__.0@   ( (dhhj! D  // ! ! # " CIIaL Jr:rc|tj}|tj}|tj}tj|||||Sr rVrget_default_devicerrrr^rrrrs r8 meta_onesrT }'') ~))+ ~ ;; E&J r:c|tj}|tj}|tj}tj|||||Sr rrs r8 meta_zerosrrr:c,tj|SrDrOclone_preserve_strides)rr*rrs r8meta_select_scatterr  ' ' --r:c,tj|SrDr)rr*rr}r|steps r8meta_slice_scatterrrr: dim_post_expr wrap_scalarc |dkr|std|dd}| }|dz }||ks||kDrtd|d|d|d|dkr||z }|S) Nrzdim_post_expr=z <= 0 but wrap_scalar is Falser3zdim z out of bounds (rxryr)rrrrrs r8rrs  /MN  .C ! C SyC#ItC5(8RuAFGG Qw } Jr:cJ|jdk(rdS|j|Srr)rrs r8ensure_nonempty_sizers!1 1.!''#,.r:c:tjd}tjd}tj||k(dt |D];k7s tjt t kfd=y)Nr3cy)NzDIndex tensor must have the same number of dimensions as input tensorr4r4r:r8rez$gather_shape_check..rmr:cNddjdjdzS)Nz!Size does not match at dimension z expected index  to be no larger than self  apart from dimension r9)rrrrsr8rez$gather_shape_check..s7;A3>Nu{{m\/ |;QRUQVWXr:)rrrVrgrr)rrr self_dims index_dimsrs``` @r8gather_shape_checkrsDHHJ"IUYY[!$J LLZV9  8 LL$UA.2FtQ2OOX r:cpddlm}t||j}|j dk(}|s`t j jt jk(xsjt jk(fdt|||jjS)Nrrc"djS)Nz8gather(): Expected dtype int32/int64 for index, but got rrsr8rezmeta_gather..sNu{{m\r:) r$rrrrrVrgr^rrrrr)rrr sparse_gradr wrapped_dimis_index_emptys ` r8 meta_gatherrsD dhhj1K#EKKMQ$67N  KK5:: % A )A \  4e4 >>%++ &&r:c|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.r4r4r:r8rez#get_operator_enum..rmr:addmultiplycy)Nz/reduce argument must be either add or multiply.r4r4r:r8rez#get_operator_enum.."rmr:rc)reduce_use_new_optionss r8get_operator_enumr s{ e   $    #  #  S   e   "$ UUVr:cPddlm}||jdk7rStj|j tj k(xs|j tjk(fd|1tj|j |j k(fdyy)Nr)rScdS)Nz((): Expected dtype int32/int64 for indexr4 method_namesr8rez,scatter_gather_dtype_check..-s{m#KLr:cdS)Nz0(): Expected self.dtype to be equal to src.dtyper4rsr8rez,scatter_gather_dtype_check..3s{m#STr:)r$rSrrVrgr^rr)rrrsrc_optrSs` r8scatter_gather_dtype_checkr'svCU[[]a'( KK5:: % A )A L  JJ'-- ' T r:ct|dSr2)rrs r8ensure_nonempty_dimr7s sA;r:cddlm}|jdk(rytjt j t j k(dt j }t|D]>}|k(r t|}t|}tj||kfd@tjt j t j k(dt|D]7}t|}t|} tj|| kfd9yy)Nrrcy)NzCIndex tensor must have the same number of dimensions as self tensorr4r4r:r8rez%scatter_shape_check..Crmr:cHdjdjdzS)NExpected index rrr9)rrrsr8rez%scatter_shape_check..Qs,oekk]2Mdjj\Z&se,-r:cy)NzBIndex tensor must have the same number of dimensions as src tensorr4r4r:r8rez%scatter_shape_check..Yrmr:c<djdjS)Nr#z to be no larger than src r9)rrsr8rez%scatter_shape_check..`s!/%++6PQXQ^Q^P_`r:) r$rrrVrgrrrr) rrrrrrr index_d_size self_d_size src_d_sizes ```` r8scatter_shape_checkr)<s-Dekkmq() LLDHHJ'+>uyy{+KKU $DHHJ/I9    8 +E15 *43  K ' -    +/B7;;=/Q Q X y! A/q9L-gq9J LL *`   r:ct||j}td|||t||||| t ||yy)Nscatter)rrrr)r)rrrr*rrrs r8scatter_meta_implr,esE dhhj1Ky$s;k5#6'?3r:cVt||||d|j|jSNrr,rrrrrr*s r8meta_scatter_addr1ns%dCU3 >>$** %%r:c$t||||d|Sr.r,r0s r8meta_scatter_add_r4tsdCU3 Kr:ct|tjr|nd}t||||||j |j SrD)rprVr r,rrrrr src_or_valuerkr*s r8 meta_scatterr8zs;%\5<<@,dCdCV4 >>$** %%r:c`t|tjr|nd}t||||||SrD)rprVr r,r6s r8 meta_scatter_r:s-%\5<<@,dCdCV4 Kr:queryrr+ dropout_p is_causalreturn_debug_maskr}c .|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) Nrr3rr0r@rr4r)rrVrrrZrrceilr^rrrr~rrr8)r;rr+r<r=r>r}r num_headsmax_seqlen_batch_qhead_dimmax_seqlen_batch_k attention logsumexp blocksize_c max_seqlen_k debug_maskseedoffsets r8(meta__scaled_dot_product_flash_attentionrOsAJ 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      r: q_descale k_descale v_descalec |jtjk(r|jtj}t ||||||| SrD)r^rVrcr'rbrO) r;rr+rPrQrRr<r=r>r}s r82meta__scaled_dot_product_flash_attention_quantizedrTsJ {{e)))( 3     r: res_shape.ctj|k(rtj}|St gdfdd}|Dcgc]}|| }}t t |Dcgc]}|j|}}tj|jjj|}|Scc}wcc}w)N)rr3rr0c*j|SrDr)idxr;s r8rez,alloc_with_matching_layout..s%,,.*=r:Trr) rfrrVrsortedrrrrr^rr)r;rUr dim_orderrXpermuted_shaper final_permutes` r8alloc_with_matching_layoutr]s 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 SNrr3rrrr4rrr]rVrrZrr)r;rr+r^r_r<r=r>r}rr*S_QS_KVD_VrUr logsum_exprMrNs r8(meta__scaled_dot_product_cudnn_attentionrgs  1 A 1 A **Q-C 88A;D **R.CAsC I $UI 6C AsAkk||J ;;rF ;D [[5::f =F         r:c |jd}|jd} |jd} |jd} |jd} || | | f} t|| }tj|| | ftj|j }tjdtj d}tjdtj d}||dd| | ||df Srarb)r;rr+r^r<r=r>r}rH_QrcrdrerUrrfrMrNs r85meta__scaled_dot_product_fused_attention_overrideablerj1s  1 A **Q-C **Q-C 88A;D **R.CCc"I $UI 6C C kk||J ;;rF ;D [[5::f =F         r:rrI cum_seq_q cum_seq_kmax_qmax_k philox_seed philox_offsetctj|}tj|}tj|}|||fSrDr)rr;rr+rrIrkrlrmrnr<r=rorpr}grad_qgrad_kr#s r8'meta__scaled_dot_product_flash_backwardrt\sA,  e $F   c "F   e $F 66 !!r: attn_maskc |jd}|jd}|jd} tj|} tj|| |ftj|j j dd} | | fS)Nrr3rr)rrVrrrZrr) r;rr+r<r=rur}rrDrErHrIs r80meta__scaled_dot_product_flash_attention_for_cpurwxsAJ 1 IA  'I    kk||i1o  r:c ntj|jd|j|j} tj|jd|j|j} tj|jd|j|j} | | | fS)Nrrr3r0r)rVempty_permutedrr^r) rr;rr+rrIr<r=rur}rrrsr#s r89meta__scaled_dot_product_flash_attention_for_cpu_backwardr{s& ! ! kk|| F  ! !  iizz F  ! ! kk|| F 66 !!r: 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) Nr0rTr r3r5rrF)rrrrviewr)r=rrs r8 ensure_4dzBmeta__scaled_dot_product_attention_math_for_mps..ensure_4ds 557a<;;q>4' ' UUWq[J1557Q;' )aggaj(  )66*affRj!&&*affRjI4O Oe8Or:cXjj} r|j}jf} rajdk(r|j d}||fSt jdd|jddz}|j |}||fS)Nr0rr5r3r )rrview_asrsqueezerr) rattnrrmax_seq_lengthnum_headq_q_sizer; unsqueezeds r8sdpa_vector_fast_mpszMmeta__scaled_dot_product_attention_math_for_mps..sdpa_vector_fast_mpssll288$ ++e$C||Z6>JK yy{a||ADyU[["-.Aa@yy'Dyr:crd}jj}j|f}||fS)Nr)r)blocksrr,r head_sizerrrs r8sdpa_vector_2pass_mpszNmeta__scaled_dot_product_attention_math_for_mps..sdpa_vector_2pass_mpss>ll288$||Z669$UV L  r:iir9)r;rr+rur<r=r|r}rk_rRv_k_sizerrrrrrrrrs` @@@@@@@r8/meta__scaled_dot_product_attention_math_for_mpsrs u%NB cNEB e EB.0hh+J&)#%88 Av~q  !! $FVO$8N$&&#%%r:c@|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) Nr3rrrrrr)r4r)rrrVrr^rrrrr~rrCrZr)r;rr+r^r_r<r=r}rrrDKvr logsumexp_dimrfrMrNs r8,meta__scaled_dot_product_efficient_attentionrsD 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& ((r:grad_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) Nrr3rr0ryrrr(.)rrVrzr^rrr)rr;rr+r^rrIrorpr<rr=r}rrDrmrF head_dim_vrnrrrsr# grad_biaslastDimlastDimAligned new_sizess r8+meta__scaled_dot_product_efficient_backwardr's{(AJ 1 I JJqMEzz!}HAJ HHQKE  ! ! Yx0kk|| F  ! ! Yx0iizz F  ! ! Yz2kk|| F I!3..$$+bLA$57R<'TV,;V)* & " KK Y__Y5E5E c8G8m, 669 ,,r:ctj|}tj|}tj|}|||fSrDr)rr;rr+rrIrorpr^rkrlrmrnr<r=r}rrrsr#s r8'meta__scaled_dot_product_cudnn_backwardrcsA.  e $F   c "F   e $F 66 !!r:window_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) Nrr3rrrr@rArBNNr4rr)rrrVrrrZrrrCr^rrrr~rr8)r;rr+rkrlrmrnr<r=r>r}rrrrrrErGrDrFrHrItotal_qrJrKrLrMrNs r8meta__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   r:c|jtjk(r|jtj}t |||||||||| | ||||SrD)r^rVrcr'rbr)r;rr+rkrlrmrnr<r=r>rPrQrRr}rrrrs r8'meta__flash_attention_forward_quantizedrsb* {{e)))( (       r:ctj|}tj|}tj|}|||fSrDr)rr;rr+rrIrkrlrmrnr<r=rorpr}rr grad_querygrad_key grad_values r8meta__flash_attention_backwardrsA0!!%(J$H!!%(J x ++r: cu_seqlens_q cu_seqlens_k max_seqlen_qrKcustom_mask_typecausal_diagonalseqlen_k window_sizec|jd}|jd}|jd}|jd}|jd}tj|||||j|j}||jddz n|}|}|| t d|}||n|}| rt j|dz dznd}tj|||ftj|j}tjdtjd }tjdtjd }||||||fS) Nrr3rrrz;max_seqlen_q must not be None when cu_seqlens_q is providedr)r4r) rrVrr^rrrrCrZr)r;rr+r;rrrrKr<rr_r}rrrrrrrDrrlogsumexp_batch_dimactual_max_seqlen_qactual_max_seqlen_krrfrMrNs r8!meta__efficient_attention_forwardrsE,  1 A 1 A  A 2I BB ++aIrU\\ RC7C7O,++A.2VW   M +*6*B,4F %*+b0A i7kk||J ;;rF ;D [[5::f =F  D&*=?R RRr:bias_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)Nr3cy)Nz,seqlen must match for `shared_storage_dqdkdvr4r4r:r8rez4meta__efficient_attention_backward..mrmr:r0cy)Nz3embedding dim must match for `shared_storage_dqdkdvr4r4r:r8rez4meta__efficient_attention_backward..qrmr:rrrrr5rr(.r4r) rVrgrrr^rrrrr)rr;rr+r;rrrrKrIr<rorprrr}rrchunkrrrrrrrs r8"meta__efficient_attention_backwardrQs2 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 66r:scale_ascale_b scale_resultuse_fast_accumc d}tjjdk(xrjdk(fdtj|jxr|jfdt dk(st dk(rad} 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(dnb| rjtjk(rd } dznd} d}t| }t|ddz}|t|z|z|t|z|zjk(r_jk(rLtjjdtjjdntjdfdnhtjjtjk(xrjtjk(dtjjdk(xrjdk(fdj dk(ruj d d k(raj dd k(rMj d k(r9tjjxrjdnNj dk(rUj d j dcxk(rtdk(r"nnj d tdk(rnj dk(rKj d j dcxk(rtdk(rnnj d k(rnj dtdk(rKj d j dcxk(rtdk(rnnj d k(rntjdfd||n j}tjj dj d |jS)Nc|tjtjtjtjtj fvSrDrVrc float8_e5m2float8_e4m3fnuzfloat8_e5m2fnuzfloat4_e2m1fn_x2rs r8is_fp8_or_fp4_typez2_check_scaled_mm_sizes..is_fp8_or_fp4_typeA        ! !  ! !  " "    r:rcLdjdjSNz%Inputs must be 2D but got self.dim()=z and mat2.dim()=rrJrsr8rez(_check_scaled_mm_sizes..'7 |CSTXT\T\T^S_`r:c<djdjSNz?Expected both inputs to be fp8 or fp4 types but got self.dtype=z and mat2.dtype=rrsr8rez(_check_scaled_mm_sizes..&QRVR\R\Q]]mnrnxnxmyzr:rrc,|d|dkDxr|ddk(Srr4rs r8 is_row_majorz,_check_scaled_mm_sizes..is_row_major"!9vay(;VAY!^ ;r:c&|ddk(xr|ddkDSrr4rs r8 is_col_majorz,_check_scaled_mm_sizes..is_col_major!9>3fQi!m 3r:cV|jddk(xs|jddk(Srr tensor_2ds r8 has_zero_dimz,_check_scaled_mm_sizes..has_zero_dim)>>!$)CY^^A->!-C Cr:c*djSNz#self must be row_major, got stride rrsr8rez(_check_scaled_mm_sizes..9$++-Ir:c*djSNz#mat2 must be col_major, got stride rrJsr8rez(_check_scaled_mm_sizes..rr:r3r(rc,djdSNzBExpected self.size(1) to be divisible by 16, but got self.size(1)=r3rrsr8rez(_check_scaled_mm_sizes..XY]YbYbcdYeXfgr:c"djSNz?Expected both dimensions of mat2 to be divisible by 16 but got r9rsr8rez(_check_scaled_mm_sizes..UVZV`V`Uabr:cy)NzNFor tensorwise scaling, both scale_a and scale_b must be float (fp32) tensors.r4r4r:r8rez(_check_scaled_mm_sizes..rmr:r)rAr cy)Nzscale_a must be contiguousr4r4r:r8rez(_check_scaled_mm_sizes..rmr:cy)Nzscale_b must be contiguousr4r4r:r8rez(_check_scaled_mm_sizes..rmr:Fc Zddjddjd S)NzTInvalid blockwise scaling configuration. For blockwise scaling, scale_a should have  elements, got z, scale_b should have rr)expected_a_sizeexpected_b_sizerrsr8rez(_check_scaled_mm_sizes..sHFFUEVVefmfsfsfuevw//>.?w}}N__`br:cy)NzKFor rowwise scaling, both scale_a and scale_b must be float (fp32) tensors.r4r4r:r8rez(_check_scaled_mm_sizes.. rmr:cLdjdjS)NzLFor non-tensorwise scaling, scale tensors must be 2D, but got scale_a.dim()=z and scale_b.dim()=rrrsr8rez(_check_scaled_mm_sizes..s,gY`YdYdYfXhh|nunynyn{m}~r:cy)Nz@Both scale_a and scale_b must be contiguous for rowwise scaling.r4r4r:r8rez(_check_scaled_mm_sizes..rmr:cddddtdd dtddtdddtdd zdtddd tddtdd zdtddd jd djd d jd djd d zS)N}Invalid scaling configuration. For tensorwise scaling, both scales should be scalar. For rowwise scaling, scale_a should be (, 1), scale_b should be (1, >). For (BlockWise1x128, BlockWise128x128), scale_a should be (rxrA), scale_b should be (<). For (BlockWise1x128, BlockWise1x128), scale_a should be (z>). For (BlockWise128x128, BlockWise1x128), scale_a should be (). Got scale_a.size()=(rr3) and scale_b.size()=(ryr9r)_krfrrrsr8rez(_check_scaled_mm_sizes..7sLCCD#EabcadeVVWUXXZ[cdfhk[lZmmpr0S0A/B"XaQTEUDVWTTUSVVXYabdfiYjXkknp p 0S0A/B"QCHVV^_`beVfUggijrsuwzj{i||A A0S0A/B"QCH//6||A.?r',,q/ARS//6||A.?r',,q/ARRSU Ur:r)rVrgrr^rrrrfloat8_e8m0fnurcrr`r9rdrr)rrJrrr;rr=rrrrris_blockwise_scaling block_size_k block_size_mn num_k_blockspadded_num_k_blocks _out_dtyperrrrfrs```` @@@@@r8_check_scaled_mm_sizesr s   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( Q8As#33LLOw||AK(2s:KKLLOq(  "(3J ;;tyy|TYYq\DKK XXr:c &t||||||||SrD)r )rrJrrr;rr=rs r8meta_scaled_mmrJs" " dGWdL)^ r:scale_recipe_ascale_recipe_b swizzle_a swizzle_bc   !"#$%&d} d} tjjdk(xrjdk(fdtj| jxr| jfdjd jdjd!| 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}dtt dtt fd}|||rtjdjdk(xrZdjdk(xrBdjtjk(xr djtjk(d n|||rtjdjd k(xrrdj k(xrZdjtjk(xr8dj!k(xr djtjk(fdn|||rqdjtjk(xr djtjk(}d%%jd k(xrk%jddzk(xrT%jddk(xr>%jd k(xs(%jddk(xr%jddk(}d&&jd!k(xrk&jddzk(xrT&jddk(xr>&jd!k(xs(&jddk(xr&jddk(}tj|xr|xr| !%&fdne|||rdjtjk(xr djtjk(}tdz d d%%jdk(xrk%jd dzk(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 d d%%jd k(xrk%jddzk(xrT%jddk(xr>%jd k(xs(%jddk(xr%jddk(}d&&jdk(xrk&jd!dzk(xrT&jddk(xr>&jdk(xs(&jddk(xr&jddk(}tj|xr|xr| !%&fd"nO|||rtjj rct#jdd#jdz"t#jdd#jdz#t j$ntjddtt#jdd#d z"tjddtt#jdd#d z#t j$$tjdj"k(xrndjtj&k(xrLdj#k(xr4djtj&k(xrd$k(xr d$k("#$ fd$n|||rt dtt#dd z"t!dtt#dd z#t j$$tjdj"k(xrndjtj(k(xrLdj#k(xr4djtj(k(xrd$k(xr d$k("#$ fd%n|||rnt dtt#dd z"t!dtt#dd z#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 fvSrDrrs r8rz5_check_scaled_mm_sizes_v2..is_fp8_or_fp4_typegrr:c(|tjk(SrD)rVrrs r8 is_fp4_typez._check_scaled_mm_sizes_v2..is_fp4_typeps....r:rcLdjdjSrrrsr8rez+_check_scaled_mm_sizes_v2..urr:c<djdjSrrrsr8rez+_check_scaled_mm_sizes_v2..yrr:rr3rrc,|d|dkDxr|ddk(Srr4rs r8rz/_check_scaled_mm_sizes_v2..is_row_majorrr:c&|ddk(xr|ddkDSrr4rs r8rz/_check_scaled_mm_sizes_v2..is_col_majorrr:cV|jddk(xs|jddk(Srrrs r8rz/_check_scaled_mm_sizes_v2..has_zero_dimrr:c*djSrrrsr8rez+_check_scaled_mm_sizes_v2..rr:c*djSrrrsr8rez+_check_scaled_mm_sizes_v2..rr:r(c,djdSrrrsr8rez+_check_scaled_mm_sizes_v2..rr:c"djSrr9rsr8rez+_check_scaled_mm_sizes_v2..rr:recipe_arecipe_bct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr( TensorWiser r!s r8 is_tensorwisez0_check_scaled_mm_sizes_v2..is_tensorwisesTH ":MQ&:QK;#9#99:QK;#9#99  r:ct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr(RowWiser$s r8 is_rowwisez-_check_scaled_mm_sizes_v2..is_rowwisesTH "7MQ&7QK;#6#667QK;#6#66  r:ct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr( BlockWise1x32r$s r8is_mxz(_check_scaled_mm_sizes_v2..is_mxsTH "=MQ&=QK;#<#<<=QK;#<#<<  r:ct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr( BlockWise1x16r$s r8is_nv_single_levelz5_check_scaled_mm_sizes_v2..is_nv_single_levelsVH "=MQ&=QK;#<#<<=QK;#<#<<  r:ct|dk(xrnt|dk(xr^|dtjk(xrF|dtjk(xr.|dtjk(xr|dtjk(S)Nrrr3)rr(r-r#r$s r8is_nvz(_check_scaled_mm_sizes_v2..is_nvsH ":MQ&:QK;#<#<<:QK;#9#99:QK;#<#<< : QK;#9#99  r:ct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr(BlockWise1x128r$s r8is_1x128_1x128z1_check_scaled_mm_sizes_v2..is_1x128_1x128sTH ">MQ&>QK;#=#==>QK;#=#==  r:ct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr(r2BlockWise128x128r$s r8is_1x128_128x128z3_check_scaled_mm_sizes_v2..is_1x128_128x128sWH "@MQ&@QK;#=#==@QK;#?#??  r:ct|dk(xr>t|dk(xr.|dtjk(xr|dtjk(Sr)rr(r5r2r$s r8is_128x128_1x128z3_check_scaled_mm_sizes_v2..is_128x128_1x128sTH ">MQ&>QK;#?#??>QK;#=#==  r:cy)Nz\For Tensorwise scaling, both scale_a and scale_b must be single element float (fp32) tensorsr4r4r:r8rez+_check_scaled_mm_sizes_v2..rmr:c djdddjdjdddjd S)Nz'For Rowwise scaling, scale_a must have rz elements (got: z), and scale_b must have r3ry)rr)rJrrrsr8rez+_check_scaled_mm_sizes_v2.. sa=djjm_L\]def]g]m]m]o\pq//3zz!}o=MgVWjN^N^N`Maabdr:rAcdddzdjddjdddzdjddjdS)Nz>For 1x128 x 1x128 blockwise scaling, scale a must have shape [rxrA] (got: ) and stride [1, )scale b must have shape [ryrr)rsrrsasbsr8rez+_check_scaled_mm_sizes_v2..'s~001s"Q#XJhrxxjPabcaddlmomvmvlwx001s"Q#XJhrxxjPabcaddlmomvmvlwwxzr:r cdddzdjddjdddzdjddjdS)Nz]For 128x128 x 1x128 blockwise scaling, L4 = {round_up(K / 128, 4)}, scale a must have shape [rxrAr<r=r>ryr?rsL4rrr@rAsr8rez+_check_scaled_mm_sizes_v2..Fs~002t2a3hZxzQbcebffnoqoxoxnyz001s"Q#XJhrxxjPabcaddlmomvmvlwwxzr:cdddzdjddjdddzdjddjdS)Nz]For 1x128 x 128x128 blockwise scaling, L4 = {round_up(K / 128, 4)}, scale a must have shape [rxrAr<r=r>ryr?rCsr8rez+_check_scaled_mm_sizes_v2..es~001s"Q#XJhrxxjPabcaddlmomvmvlwx002t2a3hZxzQbcebffnoqoxoxnyyz|r:r)cdddjdddjdtjddjddjdd dddd S) Nz!for MX scaling scale_a must have  (got: r) and scale_b must have z). Scales must have types z (for self: , mat_b: z) Must have swizzle type (got self: ry)rrVrr^expected_scale_a_elemsexpected_scale_b_elemsexpected_swizzlerrrrsr8rez+_check_scaled_mm_sizes_v2..s78N7OwW^_`WaWgWgWiVjk--C,DGGTUJL\L\L^K_`""'"6"6!7|GAJDTDTCUU^_fgh_i_o_o^pq..>-?|IVWL>Ybclmncobppqsr:cdddjdddjddddddS) Nz.for single-level NV scaling scale_a must have rGrrH). Must have swizzle type rJrIryrrKsr8rez+_check_scaled_mm_sizes_v2..swDE[D\\cdklmdndtdtdvcwx--C,DGGTUJL\L\L^K_`$$4#5\)A,yYbcdYeXffgir:cdddjdddjddddddS) Nz!for NV scaling scale_a must have rGrrHrPrJrIryrrKsr8rez+_check_scaled_mm_sizes_v2..sv78N7OwW^_`WaWgWgWiVjk--C,DGGTUJL\L\L^K_`$$4#5\)A,yYbcdYeXffgir:Fcbddddtdd dtddtdddtdd zdtddd d jd dd jd d d jd dd jd d zS)NrrrrxrArrrrrr3rryr)rsrrrrsr8rez+_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 Wr:r)rVrgrr^rr(r) NO_SWIZZLErrrrrr`r?rrr9SWIZZLE_32_4_4rrcrr)'rrJrrrrr;r=rrrrrK_packed_multipliersirrrr%r(r+r.r0r3r6r8types_ok scale_a_ok scale_b_okr rsrDrrrLrMrNr@rAs'``` ` `` @@@@@@@@@r8_check_scaled_mm_sizes_v2rZZs$  / 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  ;' 37 3D  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   ?%-a%5!RRS8T%T "%-a%5!RRS8T%T "*99  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 BBm @?: :s0ww"w wrcontraction_dimsc .t|||||||| |||  S)N)r;r=rrr)rZ) rrJrrrrrrr;rr[rs r8meta_scaled_mm_v2r]s2 %   %  r:cZt|||||d|j|jSNT)rr/rrrr*rkris r8meta_scatter_reduce_tworas)dCVTJ >>$** %%r:c(t|||||d|Sr_r3r`s r8meta_scatter_reduce__tworcsdCVTJ Kr:cttjdjcxkxrdkncfdjdk(r0tj|tjj Stjj d|tjj S)Nrrc*djS)NzAThe probability distributions dimensions must be 1 or 2, but got rrsr8rez"meta_multinomial..sSTYT]T]T_S`ar:r3r)rVrgrrrrr)r num_samples replacementrs` r8meta_multinomialrhs| LL EIIK1a yy{a{{;ejjNN ;; 1 {%**U\\ r:c"d}|D]}||z} |Sr2r4)vsravs r8multiply_integersrls$ A  Q Hr:cLtjtk(fddztjtk(fdtjtdddDxrtdDfddd\}}||gS)Nc&ddtS)Nz%It is expected output_size equals to , but got size r)num_spatial_dimsr{sr8rez'upsample_common_check..s78H7IY\]hYiXjkr:rc&ddtS)Nz$It is expected input_size equals to ror)expected_input_dimsrQsr8rez'upsample_common_check.. s67J6K?[^_i[jZklr:c3&K|] }|dkD ywrNr4rrrs r8rtz(upsample_common_check.. s *aAE *c3&K|] }|dkD ywrtr4rus r8rtz(upsample_common_check.. s2NQ1q52NrvcddS)NzDInput and output sizes should be greater than 0, but got input size z and output size r4)rQr{sr8rez'upsample_common_check..s \!2;-Ar:)rVrgrr)rQr{rprchannelsrrs``` @r8upsample_common_checkrzs LL K,,k+Q. LL J..l  LL *:ab> **Ns2N+2N/N A ""1~FH H +{ ++r:c4tjjdk7xstj ddfdt j |d}j |jtjS)Nrr3c*djS)Nz>Non-empty 3D data tensor expected but got a tensor with sizes rrsr8rez$upsample_nearest1d..PQVQ[Q[Q]P^_r:rpr rVrgrrlrrzrr'rOr)rr{scalesfull_output_sizes` r8upsample_nearest1dr LL A/ QR0@A_- kA ??+ , / /11%8 0 r:ctjjdk7xstj ddfdt j |d}j |}tj}j\}}}}jjdk(r|dkrtj}|j|}|S) Nrr3c*djSNz>Non-empty 4D data tensor expected but got a tensor with sizes rrsr8rez$upsample_nearest2d..,r}r:rr~rr r)rVrgrrlrrzrrOrrrrzr contiguous) rr{scales_hscales_wrrrrR n_channelss ` r8upsample_nearest2dr&s LL A/ QR0@A_- kA__- .F//6M ++Az1a ||F"zA~//   ]  ;F Mr:r{rQrrcXt||dtjjdk(fdt dD]2tjj k(fd4j |jtjS)Nrr~r c"djSNzFExpected grad_output to be a tensor of dimension 4 but got: dimension rrsr8rez-upsample_nearest2d_backward..RXYdYiYiXjkr:c DddddjS)NzCExpected grad_output to have the same shape as output; output.size() = z but got grad_output.size(rrrrsr8rez-upsample_nearest2d_backward..Ws> !s$'7':&;,QCtK4D4DQ4G3HJr:r) rzrVrgrrrrr'rOr)rr{rQrrrrs` @@r8upsample_nearest2d_backwardr@s-K! LLAk1X    Q #3A#6 6      , / /11+> 0 r:c4tjjdk7xstj ddfdt j |d}j |jtjS)Nrr3c*djS)Nz>Non-empty 5D data tensor expected but got a tensor with sizes rrsr8rez$upsample_nearest3d..ir}r:r0r~rr)rr{scales_drrrs` r8upsample_nearest3drcrr:ctj|tj|tj}}||t|tst dt |t|tst dt ||j}|j} t||}t||}|j|| |j|| t||t||||fS||fS)Nrzvalues must be TensorLike, got z indices must be TensorLike, got )rr) rVrrrprrrzrrr!rr#) rstabler descendingrrrkrr out_strides r8 meta_sortrss   D !5#3#3D #LqA g1&*- #B4<.!QR R':. #CDM?!ST TGG XXZ "695#GY79j1Iz2F3G4w a4Kr:c tjjdk(fdtjjjk(fdj dtjjdk(fdtjj k(fdtjjjk(fdtjjdk(fdj d zztjj k(fd tjt fd fDd y) Nrc"jdSNz != 2r) input_gatessr8rez%rnn_cell_checkSizes..;3C3C2DE0Jr:c:jdjSNrCr9) hidden_gatesrsr8rez%rnn_cell_checkSizes..s ;$$%T,*<*<)=>r:r3c"jdS)Nz != 1r) input_biassr8rez%rnn_cell_checkSizes..sjoo5Fe3Lr:c.jdSrr) gates_sizersr8rez%rnn_cell_checkSizes..sz'')*$zl;r:c:jdjSrr9) hidden_biasrsr8rez%rnn_cell_checkSizes..s z''([->->,?@r:c"jdSrr) prev_hiddensr8rez%rnn_cell_checkSizes..rr:rc `jdjddddd S)NrCrz * z // z (aka ry)rr)expected_prev_hidden_numelfactorrrrsr8rez%rnn_cell_checkSizes..sB;$$&'tK,<,.s(  HH ** * s#&cy)Nz%expected all inputs to be same devicer4r4r:r8rez%rnn_cell_checkSizes..rmr:)rVrgrrrrr)rrrrrrrrs``````@@r8rnn_cell_checkSizesrs@ LL!!Q&(JK LL\///>!!!$J Z__)+LM    * , ;      1 1 1 @  LL!!Q&(JK!,!1!1!!4z!AV!K LL99 LL  "J [I   8 r:c t||||d|tj|tj}tj|tj}tj|tj}|||fS)Nr r)rrVrr)rrcxrr workspacehycys r8_thnn_fused_lstm_cell_metarsk \:{ArR  E ?j(^*CD  __Y 'F~-z;GJ z [[5<< 0 \\* % zN2JI JBAAMoom5;;o?G 2r7J ..r:cd}|D])}|jdkDs||jz }+|j|f}t||||||||d||| | | | | Srs)rrr)rr9rrrr]rrrrrrrrtotal_weight_elemsrrs r8 miopen_rnnrs& , 779q= !'') + ,"4!67J       ! r:c(|r|jdn|jd}|r|jdn|jd}| }|r|||gn|||g}|j|}|"tjd|j}n|j|j}|"tjd|j}n|j|j}tjd|jtj }||||fS)Nr3rrr)rrrVrrr)rw0w1w2w3hx_cx_rrr]rr has_biasesrrrrroutput_chanelsrrrrrs r8mkldnn_rnn_layerr+s&$/QEKKNJ#.QEKKNJ N  Z0*n 5 __Y 'F { [[5<< 0 ]]399 % { [[5<< 0 ]]399 % Aell%++FI 2r9 $$r:c|jdk(r%tjdk(xsdk(fdytj|jdk7fdy)NrrcdS)Nz4: Expected reduction dim -1 or 0 for scalar but got r4rrsr8rez'zero_numel_check_dims..WswiSTWSXYr:cddS)Nz: Expected reduction dim z to have non-zero size.r4rsr8rez'zero_numel_check_dims..\swi8=TUr:)rrVrr)rrrs ``r8zero_numel_check_dimsrSsR yyA~  1H !r Y  IIcNa  U r:c|(t||j}t||ytj|j dk7fdy)NrcdS)Nz@: Expected reduction dim to be specified for input.numel() == 0.r4r&sr8rez%check_argmax_argmin..hstf\]r:)rrrrVrgr)rrrs` r8check_argmax_argminrasC S$((*-dC. JJLA  ] r:ctd||tj|j||fnd}t |||}|j |t jS)Nargmaxr)rrOrrrrrVr)rrrrrs r8argmax_argmin_metarlsQ$,    coSF4 PD $T4 9E >>%u{{> 33r:c||tjk(rtj}tjd||||S)Nr4r )rVjaggedrr)rr^rrrs r8 scalar_tensorrts5  ;; %v* r:ct||jd}|jdk(rdn|j|}tj|dk\tj||kdt |j }t|dkDr|||<|j||j|tjfS)NTrrr3cy)Nzk not in range for dimensionr4r4r:r8reztopk_meta..rmr:r) rrrrVrgrrrrr)rrHrlargestrY sliceSizetopKSizes r8 topk_metars dhhjd ;CXXZ1_$))C.I LLa LLi!GHDJJH 8}q >>( #T^^HEKK^%P PPr:c| | td|j}|j} tj|| j| j | j S)Nz;segment_reduce(): Either lengths or offsets must be defined)r^rr)rrrVrr^rr) rLrr{rkrvrwrxrz data_contig grad_contigs r8meta__segment_reduce_backwardrsh 7? I  //#K//#K   !!!!  r:cddlm}t|jd|jdkDr|j nd}t j ||dk\||kfdt|jd|jdzdz}|r%|jdkDr|jd|j||j|t jfS)Nr)r;Trr3cdS)Nz9kthvalue(): selected number k out of range for dimension r4rsr8rezkthvalue_meta..sKC5Qr:r) r$r;rrrrVrgrrr/rr)rrHrrr;dimSizers ` r8 kthvalue_metars> dhhjd ;C $ QdiinAG LLQW %Q DS!DJJsQwy$99 :E488:> S! >>% $..ekk."J JJr:c||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) NrcyNr4r4r:r8rez(checkLSTMBackwardSizes..rmr:cyrr4r4r:r8rez(checkLSTMBackwardSizes..rmr:cyrr4r4r:r8rez(checkLSTMBackwardSizes..rmr:cyrr4r4r:r8rez(checkLSTMBackwardSizes..rmr:cyrr4r4r:r8rez(checkLSTMBackwardSizes..rmr:cyrr4r4r:r8rez(checkLSTMBackwardSizes..rmr:rr3r cyrr4r4r:r8rez(checkLSTMBackwardSizes..rmr:)rVrgrrr)grad_hygrad_cyrrr defined_gradexp_sizes r8checkLSTMBackwardSizesr s%17wL LL!!#q(*5  "H W\\^x/< W\\^x/< LLh& 3 LLh& 3 LLA%z2 LL"hqkHQK&?!&CCZPr:c||yt|||||tj|t}tj|t}|r|j ddnd}|||fS)NNNNrrF)r)r rVrlegacy_contiguous_memory_formatr) r r rrrhas_bias grad_gatesgrad_cxrs r8#_thnn_fused_lstm_cell_backward_implr sl7?7GRY?!!!@Jr1PQG4< q%0$I w ))r:cd}d}d}|dr|j|j}|ds|drQ|j|jd|jdf}|j|jd}|||fS)Nrr3rrr)rrrrr grad_weightrs r8linear_backwardr sJKI1~!++FKKM: 1~Q",,l.?.?.CV[[QS_-UV  **<+<+.is_channels_lasts$""88=ATATTTr:c2r.tdk(rtjStjSj tjrtjSj tj rtj Sy)Nrr)rrVrrrdpreserve_format)rrsr8rz.meta_pixel_shuffle..pick_memory_formats| D !4 F*...***   e.E.E  F** *   e.C.C  D(( (Er:rrr)rrrrr') rupscale_factorrrHrWrrrrs ` @r8meta_pixel_shuffler s DJJ! 2.>2Q RVW W5djj\AXYgXh i  U ) 2>N:;A B. (B B. (B-$**Sb/-1-b-"-I .. #C &&13& 4C Jr:cX|j|j}|j|j}|j|j}|j|j}|j|j}|j|j}|||||||fSrDr)rweight0weight1weight2weight3rcx_tmprhy_cy_grad_output_r_opt grad_hy_r_opt grad_cy_r_optrr]rrrrrrrrdiff_xdiff_hxdiff_cxdiff_w1diff_w2diff_bs r8mkldnn_rnn_layer_backwardr, s4__U[[ )FmmCII&Gv||,G .G .G   w}} -F 7GVVWg EEr:) out_int32rctj||rtjntjtjS)Nr^r)rVrrCrrr boundariesr- rs r8meta_bucketizer2 %s2    &ekkEKK-- r:r1 r- rcr|jd|rtjStjS)Nr4r)rrVrCrr0 s r8meta_bucketize_scalarr4 /s>    &ekk  ,1KK  r:cdtdk(r'tjjfdtdk(r%jrt j dtjt tfdtjdkDfdtjt tfd tjt tfd tjk\fd tjjj S) Nzhistc()rc$djdS)Nz%"histogram_cpu" not implemented for 'rrrsr8rezmeta_histc..Ds=ekk]!Lr:rz%_histc_cuda with floating point inputc$dtS)Nz#: argument 'bins' must be int, not rbinsrsr8rezmeta_histc..Js7)>tDzlKr:rcdS)Nz: bins must be > 0, but got r4r8 sr8rezmeta_histc..LsgY.J4&#Qr:c$dtS)Nz%: argument 'min' must be Number, not r)rrsr8rezmeta_histc..O7)@c Lr:c$dtS)Nz%: argument 'max' must be Number, not r)rrsr8rezmeta_histc..Sr< r:cdS)Nz: max must be larger than minr4)rsr8rezmeta_histc..Usy0M%Nr:r) rrVrgrrOrrprrrrr^)rr9 rrrs````@r8 meta_histcr? =sG5U"  # # % L 5V#(?(?(A %%&MN LL4!K LLQR LL3L LL3L LLNO ;;tELL DDr:cBtj|d}tjj dk7xs#t djddDfdj |jtjS)Nrr~rc3&K|] }|dkD ywrtr4)rrrs r8rtz,meta_upsample_bimode2d_aa..gs!Ht$(!Hrvr3c*djSrrrsr8rez+meta_upsample_bimode2d_aa..hr}r:r) rzrrVrgrrrr'rOr)rr{rrrrs` r8meta_upsample_bimode2d_aarC Ys- kA LL Hc!Huzz|AB7G!HH_ ??+ , / /11%8 0 r:cTt||dtjjdk(fdt dD]0tjj k(fd2j |jtjS)Nrr~r c"djSrrrsr8rez4meta_upsample_bimode2d_aa_backward..}rr:c DddddjS)NzD Expected grad_output to have the same shape as output; output.size(rz but got grad_output_size(rrsr8rez4meta_upsample_bimode2d_aa_backward..s@DDE3dK[\]K^J_`D!1!1!!4 59r:r) rzrVrgrrrrr'rOr)rr{rQrrrrrs` @@r8"meta_upsample_bimode2d_aa_backwardrG os-K! LLAk1X    a $4Q$7 7 9     , / /11+> 0 r:cPtj|jdk(dtj|jdk(dtj|jjdtj|jjdy)Nr3cy)Nz%found_inf must be a 1-element tensor.r4r4r:r8rez<_amp_foreach_non_finite_check_and_unscale_..rmr:cy)Nz%inv_scale must be a 1-element tensor.r4r4r:r8rez<_amp_foreach_non_finite_check_and_unscale_..rmr:cy)Nz!found_inf must be a float tensor.r4r4r:r8rez<_amp_foreach_non_finite_check_and_unscale_..rmr:cy)Nz!inv_scale must be a float tensor.r4r4r:r8rez<_amp_foreach_non_finite_check_and_unscale_..rmr:)rVrgrr^r)rr inv_scales r8*_amp_foreach_non_finite_check_and_unscale_rN s| LLQ O LLQ O LL))3 LL))3r:c,tj|SrDr)rnanposinfneginfs r8 nan_to_numrS rvr:c|jtjtjtjtj hvrt d|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) rrVr sparse_cscr sparse_bscrrrrrrr)rdim0rfndimsrrs r8rrs {{      LT[[MY` a   IIE $ &D $ &D t|  D $++- F!'vd|F4L&,!$ZdDJT T6" Kr:c|j}|jr<|j}|j}|dkr|dk(s4t d|d|d|j dkDrt d|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 rMr3)rrrrrrr)rrX rrs r8t_rZ s IIE ~~__& NN$ aIN %,l9+=NP  88:> HqQ  dAEAIq 551 55r:)r- rsidesorterc^tjtjdkxsjddjddk(fdtjduxsjjk(fdtj|dk7xs| d|rtjntj }t tjr&tj|tjStjd|j S) Nr3rc`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_sequencesr8rez#meta_searchsorted..s8 3378M8M3N2OP""&tzz"2!3 5r:cldtjdtjSgS)Nz[torch.searchsorted(): boundary and sorter must have the same size, but got boundary tensor z and got sorter tensor r_ )r` r\ sr8rez#meta_searchsorted..sO ##'(=(=#>"??V%+%7tFLL!@ B>@@ Br:rcy)Nzetorch.searchsorted(): side and right can't be set to opposites, got side of left while right was Truer4r4r:r8rez#meta_searchsorted..rmr:r/ r4r) rVrgrrrCrrpr rrrr)r` rr- rr[ r\ r^s`` ` r8meta_searchsortedrc s LL O ! !"a' 9  " %CR 8   LL$?///6<<?  LL #e) $ %EKK%++E$ % U-D-D  {{2U?3I3IJJr:ctjtjtjtjfvfdy)NcdS)Nz/Unsupported input type encountered for isin(): r4rsr8rez3_check_for_unsupported_isin_dtype.. sA%Ir:)rVrgr complex128 complex64rs`r8!_check_for_unsupported_isin_dtyperh s/ LL ejj%"2"2EOODDIr:cJ|j||jdf}|S)Nrr)rr num_weightsrrr s r8meta_embedding_dense_backwardrk s*''k6F6Fr6J(KLK r:c j| rtj||||||||| | St||||||||| | SrD)r-_embedding_bag_sparse_backward!meta_embedding_bag_dense_backward) rLrrwrrmaximum_indicesrj rr]rrrs r8meta_embedding_bag_backwardrp se22            1            r:c Ntjjtjtjtj tj fvfd|tk(rtj|duj|jdf} | S)Nc"djS)Nz$Unsupported input type encountered: r)rLsr8rez3meta_embedding_bag_dense_backward..Ws6tzzlCr:r3) rVrgr^rarbr`float64rrr) rLrrrro rj rr]rrindex_grad_weights ` r8rn rn Hsv LL u}}ennemmU]]SSC x _D01 TYYq\'BC r:c|jd}tj|tk(dtj|j dk(tj|j dk(|jd}tj|j dk(tj|jd|k(|j |f} | S)Nr3cy)NzHembedding_bag_backward: per_sample_weights only supported for mode='sum'r4r4r:r8rez@meta_embedding_bag_per_sample_weights_backward..lrmr:rr)rrVrgrrr) rLr9rrwrr]rembedding_featuresrfrs r8.meta_embedding_bag_per_sample_weights_backwardrx _s1 LL Z LLq! LL!#$,,q/K LL"# LLQ#556 ^^[N +F Mr:) assume_uniqueinvertctjt|txst|tdt|ts!tj||j }t|ts!tj||j }t |jt |jtj|tjS)Ncy)Nz.|rmr:rr) rVrgrpr r0rrh r^rr)elements test_elementsry rz s r8 meta_isinr ws LL8V$I =&(IN h '<<1E1EF mV , ]8??K %hnn5%m&9&9:   HEJJ 77r:rctj|dk\dt|tj\}}tj ||S)Nrcy)Nz,polygamma(n, x) does not support negative n.r4r4r:r8rez meta_polygamma..rmr:rr)rVrgrrrr)rrrRrSs r8meta_polygammar sF LLaOP( ;HHOA|   D 55r:ctd)Nz.Tensor.item() cannot be called on meta tensors)r&rs r8meta_local_scalar_denser s G HHr:c,tj|SrDrrs r8silur rvr:clt|tj\}}tj||Sr)rrrrVr)rrRrSs r8sigmoidr s3) ;HHOA|   D 55r:cR|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) Nrrr3cyNz matrix batch sizes have to matchr4r4r:r8rez2_create_grouped_mm_output_tensor..rmr:rcyr r4r4r:r8rez2_create_grouped_mm_output_tensor..rmr:cy)Nzbatched dimension has to matchr4r4r:r8rez2_create_grouped_mm_output_tensor..rmr:r(r) rrrVrgr^rritemsizer(rr) rHrJoffsr= mat1_is_2d mat2_is_2dr alignment size_paddedrrs r8 _create_grouped_mm_output_tensorr sqJqJ  ! diilDIIaLAH LL !  ! ,.X  ! diim4H  LL !  ! ,.X  ! diil3H LL !  ! ,.V  ! diilDIIbMBH'TZZI }}),,, |i/!3 AIM  #"1+ 3[!DJ%q)J!! j $++  Jkk()DKKH Jr:mat_amat_br c tjduduk(dduxrdu} | rtj} tjjrYtj j r;dtj jdjvrtj} tjj| k(xrj| k(fdnTtjjtjk(xrjtjk(fdtjjdvxrjdvfdjdk(} jdk(} | r| s8tjjd jd k(d | rDd } d }tj| fdtj|fdd}|d|dtjjtjk(xrjtjk(xs<jtjk(xrjtjk(fdjtjk(xrjtjk(dfd }| r| rj dnd}|dd||dd|tj|dud| s| r}tjdufdxtjjdk(fdtjjtj"k(fdntjdudtj|dudtj|duxs|tjk(dt%|S) Ncy)Nz,Either both scale factors are given, or noner4r4r:r8rez)_meta_grouped_mm_common..rmr:gfx94rc>djdjdS)Nz5Expected inputs of E4M3 FP8 type but got mat_a.dtype= and mat_b.dtype=rrr r sr8rez)_meta_grouped_mm_common..s%KEKK=Xijojujuivvwxr:c>djdjdS)Nz1Expected inputs of BF16 type but got mat_a.dtype=r rrr sr8rez)_meta_grouped_mm_common..s%G }Tefkfqfqerrstr:)rr0cLdjdjS)Nz3Multiplicands must be 2D or 3D but got mat_a.dim()=z and mat_b.dim()=rr sr8rez)_meta_grouped_mm_common..s'Eeiik]Rcdidmdmdocpqr:rrrcy)Nz3contraction dimension of mat_a and mat_b must matchr4r4r:r8rez)_meta_grouped_mm_common.. rmr:cF|j}|ddkDxr|ddk(SNrr3rrmat mat_strides r8rz-_meta_grouped_mm_common..is_row_major s*Jb>A%=*R.A*= =r:cF|j}|ddk(xr|ddkDSr rr s r8rz-_meta_grouped_mm_common..is_col_major s*Jb>Q&=:b>A+= =r:c0djddS)NzNExpected mat_a tensor to be row major in the last two dimensions, got strides rr)r sr8rez)_meta_grouped_mm_common.. s#dejeqeqestvtwexdyzr:c0djddS)NzQExpected mat_b tensor to be column major in the last two dimensions, got strides rr)r sr8rez)_meta_grouped_mm_common.. s#ghmhththvwywzh{g|}r:cjdz djz}jdz dk(rGtdjdz k\r%t j |zdk(fdydk(rJdz tdjk\r(t j dz |zdk(fdyt j dfdy)Nr3r(rc"ddddS)Nr stride along % dim to be multiple of 16 bytes, got rr4end_dimmat_namer sr8rezF_meta_grouped_mm_common..check_valid_strides..$ s))H:^G9Dijtu|j}i~~Ar:c.dddz ddz dS)Nrr r3r rr4r sr8rezF_meta_grouped_mm_common..check_valid_strides..+ sK)H:^GaK=HmnxzADEzEoFnGGHIr:Fc*ddjdS)NzInvalid strides/sizes, got z for strides and z for sizes.r9r sr8rezF_meta_grouped_mm_common..check_valid_strides..0 s!5j\ARSVS\S\R]]hir:)r element_sizerrrrVrg)r r r r r s`` @@r8check_valid_stridesz4_meta_grouped_mm_common..check_valid_strides s'')a-#**,, ZZ\ gk "a 'Jw,?3 syy1%D -  LL7#i/14A  A %*Wq[*AS syy!F +  LL7Q;')3q8I  LLi r:r r c>djdjdS)NzhFor FP8 scales must both be float32, or for MXFP8 both scales must be float8_e8m0fnu. Got scale_a.dtype=z and scale_b.dtype=rrrsr8rez)_meta_grouped_mm_common..= sT~@G@M@MNNabiboboappqrr:r3c 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.r4 scale_namesr8rez>_meta_grouped_mm_common..check_scale..H si |3EFr:cBddjdjS)NzKFor MXFP8, scale must have same number of dimensions as target tensor, but  has mat.ndim= and scale.ndim=rr r}r sr8rez>_meta_grouped_mm_common..check_scale..Q sZ"mnxmyzHILIQIQHRRbchcmcmbn!or:r3c2ddjdS)Nrz to be 1D tensor, but got D tensor.rr}r sr8rez>_meta_grouped_mm_common..check_scale..V #)J<7QRWR[R[R]Q^^g hr:rcVddjzdjddS)Nrz to have rrz elements.r9)r r}scale_multiplierr scaled_dimsr8rez>_meta_grouped_mm_common..check_scale..Z sW)J_meta_grouped_mm_common..check_scale.._ si |3[\r:cPddjddjddS)Nrz batch dimension to be r, got rr9r sr8rez>_meta_grouped_mm_common..check_scale..c s6i |3J399UV<.X^_d_j_jkl_m^nnopr:cBddjdjS)Nz0For MXFP8, 3d tensor should have 2d scales, but r r rr sr8rez>_meta_grouped_mm_common..check_scale..j sF"RS]R^^lmpmumulvwGHMHRHRGS!Tr:r)r rAcNdjddzdjS)NzFor MXFP8, expected mat.shape=z to have scale shape of (,z ), but got r9)G blocked_K blocked_Nr r}sr8rez>_meta_grouped_mm_common..check_scale..t sQ"@ Kdefdgghiru~i~h@KLQLWLWKX!Yr:c2ddjdS)Nrz to be 2D tensor, but got r rr sr8rez>_meta_grouped_mm_common..check_scale..y r r:cVddjdzdjddS)Nrz non-batch dimension to be r3r rr9)r r}r r sr8rez>_meta_grouped_mm_common..check_scale..} sT)J<7RSVS\S\]^ak]kSlRmmstyttABuCtDDE!Fr:)rrVrgrdrrrr?) r r}r r r rsrrr r r is_mxfp8s ````` @@@r8 check_scalez,_meta_grouped_mm_common..check_scaleD 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*CCFr:rrcy)Nz:Scale result tensor provided, but it is not supported yet.r4r4r:r8rez)_meta_grouped_mm_common.. rmr:cNdjdjdS)Nz/Offsets tensor not provided, but is needed for zD/zD multiplicand layouts.rr sr8rez)_meta_grouped_mm_common.. s*Eeiik]RTUZU^U^U`Taaxyr:c,djdS)Nz.Offsets tensor must be 1D, but got offs.dim()=rrr sr8rez)_meta_grouped_mm_common.. sH TUVr:c$djdS)Nz7Offsets tensor must be integer (int32) tensor, but got rrr sr8rez)_meta_grouped_mm_common.. sQRVR\R\Q]]^_r:cy)NzJOffsets tensor provided, but is not needed for 3D/3D multiplicand layouts.r4r4r:r8rez)_meta_grouped_mm_common.. rmr:cy)Nz2Bias tensor provided, but it is not supported yet.r4r4r:r8rez)_meta_grouped_mm_common.. rmr:cy)Nz4If output dtype provided, it must be torch.bfloat16.r4r4r:r8rez)_meta_grouped_mm_common.. rmr:r)rVrgrcrrrr~get_device_properties gcnArchNamerr^rbrrr`rrrCr )r r rrr r;rr=rscaled fp8_dtype mat_a_is_2d mat_b_is_2drrr r r r s````` @r8_meta_grouped_mm_commonr s LL Dgo.>D 8WD%8F '' MM   '')5::;;A>JJJ--I 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$ JJr:c (t||dd||d|S)N)rrr r;rr=)r )r r r r;r=s r8meta_grouped_mmr s) #      r:c R|xstj}t||||||||| S)N)rrr r;rr=r)rVrbr ) r r rrr r;rr=rs r8meta_scaled_grouped_mmr s:+U^^I "    !%  r:c6t|tdk(r/|D]*}tj|jdkDd,g}|D]K}||n |j}|j r t |}|j|jd|M|S)Ninfrcy)Nz:_foreach_norm cannot compute infinity norm on empty tensorr4r4r:r8rez#meta_foreach_norm.. rmr:r4r) rZrVrgrr^rrrr)tensorsordr^rresultsr=s r8meta_foreach_normr s SzU5\! A LL A T   G 9".EAGG   0;Iq{{2Y{78 9 Nr:r= half_to_floatcP|rE|jtjtjfvrt d|jdt j |t jj\}}|s|n|}tj||tj}|S)Nz%half_to_float is True but x.dtype is z, expected half or bfloat16rr/ ) r^rVrXrbrrOrrrPrr)r=rr computation_dtyperSrs r8softmaxr s 775::u~~6 6 7y@[\ ',&>&> uDDLL'#|(5<:KL   1L@W@W XC Jr:c  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 rrXsr8rez'_constant_pad_nd_meta..!sCCH:Nr:c(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_inprYsr8rez'_constant_pad_nd_meta..!s"225c(;P 'r:c3^K|]%}t|tjxr|dk'ywrt)rprOIntWithoutSymInt)rrrHs r8rtz(_constant_pad_nd_meta..!s) I:a// 0 ;Q!V ; Is+-r3c Fdzdddzdzd S)NzThe input size z, plus negative padding rr3zG resulted in a negative output size, which is invalid. Check dimension z of your input.r4)rrl_diffrYpad_idxsr8rez'_constant_pad_nd_meta..)!sGok&1*&=%>>V7|nE#gk"2!34117! OMr:)r^rrr)rVrgrrrrnarrowrrrrr^rrr) rrYr+l_padc_input new_shapenew_dimrrr r r s ` @@@@@r8_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  r:rrrcD|jdk7rtd|jd|j}|j}|jdk(r|df}n$|jdk(r |d|df}n g||d}|j}|j ||S)Nrz'weight' must be 2-D, got z-Drr3r)rrrrr^r) r9rrrr weight_shape indices_shaperr=s r8 embeddingr 8!szz|q9&**,rJKK<._fi!s =JJ  r:rKr$funcr s r8_create_unary_float_meta_funcr h!*4]  Ir:c |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 tensorsrr3) is_nestedrrrrr)r;rr+ embed_dimr qkv_weightqkv_bias proj_weight proj_biasr need_weightsaverage_attn_weights mask_typerT output_dimr attn_weightss r8 native_multi_head_attention_faker t!s  #--5??! ^   {{} ,eooa.@AA 1 A 1 A !!!$J __Q: .F !??1a3L L !! !??1h1=L L !!q) L !!r:cBt|td}|S)Nc:t||tjSrr r<s r8r z*_create_binary_float_meta_func.._f!s q!@!M!M  r:r r s r8_create_binary_float_meta_funcr !r r:ctfd}jd}||_ttt||}|S)Nc`|g|i|}t|j|j|SrDr*)rrQrrrFs r8_fnz#_register_inplace_meta.._fn!s.''' CII6 r:rR)rr{rKgetattrr-)rFr! inplace_names` r8_register_inplace_metar$ !sO 2Y kk]!$LCL 4-l3 4S 9C Jr:cftjjjk(fdg}ttrQj dk7r1tjjjk(fd|j t|dtjiS)Nc<djdjS)Nrz for `end`, but got dtype r)r|r}sr8rezlerp..!s/%++.H Tr:rc<djdjS)Nrz for `weight`, but got dtype r)r}r9sr8rezlerp..!s!/%++6STZT`T`Sabr:rL) rVrgr^rprrrrTrrP)r}r|r9rQs``` r8lerpr( !s LL syy T 3.!rmr:r)rVrgrOr>r^rTrrPr* s r8addcdivr0 !sb LL  " "7== 1 6&&w}}5    w0O0W0W r:ci}dD]}t|}|D]}||vs||||<!|jD]\}}t|tjj r,t|t stdt||jtjjj|tjj|jdr|tdvst|d|j r|jdvrd|jvrt"j%||d|jvrt&j%||Hd |jvrt(j%||rd |jvrt*j%||t,j%||y) N)r post_autograd pre_autogradz$op_overload must be OpOverload, got CompositeImplicitAutogradrz 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::)ritemsrprV_opsHigherOrderOperatorrrrzpy_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_tabletypregistryopo op_overloadrFs r8 activate_metarO "s99-c2 9C--+3C=#C( 99/4469N R k5::#A#A B +z2 6tK7H6IJ  7 EHH00556r: 88 9 9     ; 8@@""m$;;         [--//BGG UWXK,,..?DD[RTU{//11BGG UWX+"2"2"44EJJ8<<["Ms9Nr:)Fr rD)NNNFrr3rr)Tr)r)rT)FF)TT)r_)FTN)TFF)TF)r)g?N)r@str)r4rrF)r4rFTN)FrFNFr)NF)rF)g?gUUUUUU?FN)NNNNN)rNNr3)NNF)FFN)NrQ FFN)rQ FNN)NrQ FNN)rQ FN)FN)FNNNN)NNNF)NNNNF)NrFNN)NNNN)rTT)NNrN)drr)r)rN)rFF)rQ )NTTN(rcollections.abcrrenumr functoolsrtypingrtyping_extensionsrrVtorch._prims_commonrrOr r r torch._decompr r rr torch._opsr torch._primsrrrrrrrrrrrrrrrtorch._prims_common.wrappersr r!r"r#r$rr%r&torch.fx.experimentalr'rdtorch.nn.functionalr(r) torch.utilsr*rHr+r,opsr-libraryLibraryrI rrrrr9r?rKrTr`rilinspacelogspacerrtaker)rrrrcummaxcumminrrrrrrr_fft_c2crrr_fft_r2crrandperm generator_outr rrrandintrrlow_outrrandr_fft_c2rr!rr-r4 unsqueeze_r8_sparse_semi_structured_linearrP r^rG_sparse_semi_structured_mmrT_sparse_semi_structured_addmmrX_cslt_sparse_mmrh index_reducero index_reduce_rq index_selectrusegment_reducerr unary_outrrrrrrrrr _assert_asyncrmsgr_printr_make_dep_tokenrr_functional_sym_constrain_rangerr(_functional_sym_constrain_range_for_sizer_functional_assert_asyncrrrrrrrr _linalg_eighrr_linalg_eigvalslinalg_eigvalsr linalg_eigrrrr rrrr* linalg_inv_exr-linalg_ldl_factor_exrfr3linalg_ldl_solverB linalg_lurIlinalg_lu_factor_exrMlinalg_lu_solverV lu_unpackr\re linalg_qrrlrprm _linalg_svdryr?r rrlinalg_solve_triangularrrr _linalg_detrrrrreflection_pad1drreplication_pad1drrreflection_pad1d_backwardrreplication_pad1d_backwardrrreflection_pad2drreplication_pad2dr_weight_norm_interface_backwardr%reflection_pad2d_backwardrreplication_pad2d_backwardr,r:reflection_pad3dr<replication_pad3dr?reflection_pad3d_backwardreplication_pad3d_backwardrG_pdist_forwardrZrK_pdist_backwardrQbaddbmmrj bernoullirm bernoulli_rprHrspoissonru_fused_moving_avg_obs_fq_helperrmmrrrrrmiopen_batch_normr convolutionrr@ _has_mkldnnrD r_convolution_pointwiser_linear_pointwiserhas_mklrF r _mkl_linearrrG rqconv2d_pointwiseqconv_pointwiser0rbinary binary_tensorrqlinear_pointwiserrlinear_dynamic_fp16linear_relu_dynamic_fp16rrH r max_pool2dr int4mm_packed_weight_cpurr  avg_pool2drErJavg_pool2d_backwardrR avg_pool3drlavg_pool3d_backwardrw_adaptive_avg_pool2dr|_adaptive_avg_pool3dr_adaptive_avg_pool2d_backwardr_adaptive_avg_pool3d_backwardrradaptive_max_pool2drrradaptive_max_pool3drrrrepeat_interleaverrqrrrr _unsafe_indexrconvolution_backwardraddbmmr randint_liker _fused_adam_ _fused_adamw_r _fused_adamr_int_mmr!_convert_weight_to_int4packr,#_convert_weight_to_int4pack_for_cpur1_weight_int4pack_mmr;_weight_int4pack_mm_for_cpurArJrLrx_dyn_quant_pack_4bit_weightr_dyn_quant_matmul_4bitr_weight_int8pack_mmr_cdist_forwardr_cdist_backwardr_embedding_bagr_embedding_bag_forward_onlyrrnansumrmedian nanmedianr dim_valuesr]rr logical_not_rrepeatrzero_rmul_Scalardiv_ logical_and_ logical_or_ logical_xor_radd_sub_rrsubrrounddecimalsrr __rshift__r  __lshift__rzerorrrfillrrelu_r _add_relurrrelu_with_noiser rrelu_with_noise_functionalr"rrelu_with_noise_r$ index_put_unsafe_index_putr( masked_fill_r, _masked_scaler/masked_scatter_r3masked_scatterr5masked_scatter_backwardr7 index_put_r9rGbmmrJrLrOrTr9r:rar{rsr max_pool2d_with_indices_backwardrmax_pool2d_with_indicesrfractional_max_pool2drmax_pool3d_with_indicesr max_pool3d_with_indices_backwardrrrrgrid_sampler_2d_backwardrrrrronesrzerosrselect_scatterr slice_scatterrrrrgatherrrrrr)r, scatter_addr1 scatter_add_r4r+r*r+rk value_reducer8scatter_r:#_scaled_dot_product_flash_attentionrOrTr]#_scaled_dot_product_cudnn_attentionrg0_scaled_dot_product_fused_attention_overrideablerj,_scaled_dot_product_flash_attention_backwardrt+_scaled_dot_product_flash_attention_for_cpurw4_scaled_dot_product_flash_attention_for_cpu_backwardr{*_scaled_dot_product_attention_math_for_mpsr'_scaled_dot_product_efficient_attentionr0_scaled_dot_product_efficient_attention_backwardr,_scaled_dot_product_cudnn_attention_backwardr_flash_attention_forwardrr_flash_attention_backwardr_efficient_attention_forwardr_efficient_attention_backwardSymIntrr  _scaled_mmrrZ _scaled_mm_v2r]scatter_reducetwotwo_outrascatter_reduce_rc multinomialrhrlrzr_upsample_nearest_exact1dr_upsample_nearest_exact2dr"_upsample_nearest_exact2d_backwardr_upsample_nearest_exact3drr values_stablerr_thnn_fused_lstm_cellrrrrrrrargminrrtopkr_segment_reduce_backwardrkthvaluerrr r r r pixel_shuffler r, bucketize Tensor_outr2 Scalar_outr4 histcr? _upsample_bilinear2d_aa_upsample_bicubic2d_aarC _upsample_bilinear2d_aa_backwardrG rN rS rrZ searchsortedrc rh embedding_dense_backwardrk _embedding_bag_backwardrp _embedding_bag_dense_backwardrn *_embedding_bag_per_sample_weights_backwardrx isinr polygammar _local_scalar_denser r r r r _grouped_mmr _scaled_grouped_mmr _foreach_normr _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- r0 lerp_addcmul_addcdiv_torch._refs.nn.functionaltorch._refs.specialrO r4r:r8rW sM .' #++ "U<78) T]t_ yy~~*/--*?*?PV*W' %a)X" 8BF#3"4hr2v6F"FG 3(* t}}-.     ==7/7t !!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! (4(Vt../ %)   {{T! 0>t112 $() ) ))  ){{T!)3)Xt##$$(";<,,;< \\;< 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 F C   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=6((00$2J2J2N2NOP T8V$        666 !" %Q F$$,,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          $ Y4t$$% 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  \((((Vt""#   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 *.2B,,B LLB IOB#Y_ B 3i#o B  B BIOd*BJQt%%--."$,,"$ 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::;4<=?4l599##55<<=599##55CCDE><599##55==>599##55<<=>?>599##55<<=599##55CCD$E>$L599##77??@599##<<DDE FA =BMMA?AH))11234H44Hn ##T[[__56 ./q'7'0!!(()*'+'!!))4+=+=+E+EFG ! H 2  (()* !!+!H ~ BB*00123&889: ; (()* 8+ 8 0012 <3 <>>?@ <A <"3"3"3"@F0012D!'$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    _-_Dt445 \\6\~% V% 6% Pt v3$t,,445#6#$t##$ 8%8"t,,- \;'!(.!, !!"#.$.t&&'    )()X !!499==12    3( ""DJJNN34    5(t""**+.,.t!!))*.+.  C d  /  t{{""# '$ '6   %R4t''(&)& t  !"     !!  &&     ""  88@@AB #< < < < <  <  < 4<<C<~88BBCD#   }  }  }  4<E4 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       *$Nt))112  3 t&&'4/(4/nt&&')()Xt$$,,-$%.$%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%%t~~'@'@AB        C  | EE4 !!))4+F+F+N+NO  &55==>?   @8t>>FFGH$'')<)<=> "?"uyy~~(()*:uyy~~  !6"6(t  !    -K"-K`t,,-.t++,' -' Tt1123,t>>?@.tyy 8=e 8 8 t~~ 6c66F66t''(I&I)Ityy "v"&""t|| 6&6V66'^(,$( VK VK VK\\D VK\\D VK 4- VK 4- VK,,%VK{{T!VKVKrt  $(    4- 4-  {{T!   !&''() !% $(,$(  << <<\\\\  ,,   ,,  ,,%{{T!*6t!!(() *  t}}  v C     t##$ 3%3lt~~ $ 8 8 88 8  8  88.t33;;<  * * &\*c* *=*,t00889 -":-"`d223d445d445d==>d==>d==>d==>dDDEdDDEtBBCtBBCtBBCtBBCtJJKtJJKtJJKtJJKt@@AtAABtAABtAAB tyy $t|| ./ t|| ./, tyy) !$,, / !$,, /  ENPr: