^j=K!}UddlZddlZddlZddlZddlZddlmZmZddlm Z ddl m Z ddlm Z m Z ddlmZmZddlmZmZddlZddlZddlmZddlmZddlmcmZddlmZm Z m!Z!dd l"m#Z#dd l$m%Z%dd lm&Z&m'Z'm(Z(m)Z)m*Z*dd l+m,Z,m-Z-m.Z.m/Z/dd l0m1Z2ddl3m4Z4ejjjlZ6gZ7e8e9e:d<ejvjxjzZ=Gdde Z> ddedej~de@de@fdZAe eAej~jdZCe eAej~jZDe eAej~jdZEe eAej~jZGde!deHde!fdZIe#e=je/d eDd!e!d"e!fd#ZJe#e=je/d eDd!e!d"e!fd$ZKe#e=je/d eDd!e!de!d%eMd&eMfd'ZLe#e=je/d eDd(e!d)eMd*eMd+eMd,e@d-e!f d.ZNe#e=jjgd/ZQe#e=jjBgd0e!fd1ZRe#e=je/eDd2e!de!fd3ZSe#e=je/d eDd(e!d2e!fd4ZTe#e=je/d eDd(e!d2e!d5eMd6eMfd7ZUe#e=je/eDd2e!de!fd8ZVe#e=je/eDd(e!d2e!de!fd9ZWe#e=je/d d(e!d2e!d&eMfd:ZXe#e=je/d eDd(e!d2e!d;eMde!d2e!d?e9fd@ZZe#e=jeDd(e!dAe!fdBZ[e#e=je/eDd2e!de!fdCZ\e#e=je/d eDd(e!d2e!de!fdDZ]e#e=jd2e!dEe!de!fdFZ^e#e=jd(e!d2e!dEe!de`e!e!ffdGZ_e#e=je/eDd(e!d2e!dHe!dIeMdJeMdKe@djjfd2e!dVe!dReHde!fdWZge#e=je/d eDd(e!dAe!dVe!dReHfdXZje#e=jddYZle#e=je/eDe>jjdZfd2e!dVe!dReHd%eMfd[Zme#e=jjeDd(e!d2e!dVe!dReHd%eMf d\Zne#e=jjeDd(e!d2e!dVe!dReHd%eMd e!f d]Zqe#e=jjeDd(e!d2e!dVe!dReHd^eMf d_Zre#e=jjeDd(e!d2e!dVe!dReHd^eMd e!f d`Ztd(e!d2e!dVe!dEe!dzdReHdaeHdbe!de!fdcZue#e=je/d eDd(e!d2e!deHde!fddZve#e=je/d d(e!d2e!dVe!dEe!dzdReHdaeHdbe!de!fdeZwe#e=je/d d(e!d2e!dVe!dEe!dzdReHdaeHdbe!de!fdfZxe#e=je/eDde>jjfd2e!dVe!dEe!dzdReHde!f dgZye#e=je/d eDde>jjfd(e!d2e!dVe!dEe!dzdReHde!f dhZze#e=je/eDe>jjfdAe!dVe!dReHde!fdiZ{e#e=je/d eDe>jjfd(e!d2e!dVe!dReHde!f djZ|e#e=je/ddAe!dOe!dkeMfdlZ}e#e=je/dme!dne!de!fdoZ~e#e=je/d(e!dpe8eHdeHdqeHdreHdseHf dtZe#e=jjB dd2e!deHdqeHdzdreHdzdseHf dvZde!deHdqeHdzdreHdzde`eHeHff dwZe#e=je/ ddAe!dxe!deHdqeHdzdreHdzdseHf dyZe#e=je/d(e!dpe8eHdeHdzeHfd{Ze#e=j e/d(e!dpe8eHd|eHd}eHd~eHf dZd(e!d e!dejfdZe#e=je/d eCd(e!de!deHdejfdZe#e=je/eCd(e!de!deHdejfdZdZe#e=je/dAe!de8eHde8eHde8eHde8eHde!f dZe#e=je/eDdAe!de8eHde8eHde8eHde8eHde8eHde!fdZe#e=je/d(e!de!d*eMfdZe#e=je/d>e!de8eHdeHdeHdseHde!f dZe#e=jjeD dd(e!d2e!deMdzde!fdZe#e=je=jjj!e6j"e=jjj!e6j$dAe!dkeMde@dzfdZe#e=j&e/dddAe!dkeMde@dzfdZe#e=j(e/de!deHde@fdZde!deHde!fdZe#e=j,e/dde!deHde@fdZe#e=j.e/ ddEe!de!deHde@de@de!f dZe#e=j0e/d(e!de!deHdeHde@f dZde8eHfdZde8e!deHdeHde8e!fdZde8e!fdZde8e!deHfdZde8e!deHdeHfdZe#e=j<je=j<jg dde8e!deHdeHde!dzde!f dZe#e=j>je=j>jg dd2e!de8eHdeHde8e!dzde8e!dzf dZe#e=j@jBddAe!deHdeHde`e!dffdZe#e=jBj ddAe!de8eHdeHde`e!dffdZe#e=jDjBdd2e!deHdeHde`e!dffdZe=jFjHj!e6j" dd2e!de!deHde`e!dffdZe#e=jLje=jLjge/deDdd2e!de!de!d%eHd)eHf dZe#e=jLje=jLjNge/d dd2e!de!de!dejd%eHd)eHf dZe#e=jRe/eD dd2e!de!de!d%eHd)eHde@f dZeDdd2e!de!de!d%eHd)eHf dZe#e=jVe/ddd2e!de!de!d%eHd)eHf dZe#e=jXjeDd(e!dAe!de!de!de!dzdeHdeHdeHdeHde8e@de`e!dze!dze!dzffdʄZe#e=jXjd(e!dAe!de!de!de!dzdeHdeHdeHdeHde8e@dejBdejBdejBde`e!dze!dze!dzffd̄Zde!dzde!dzfd̈́Ze#e=j^jde!dAe!de8eHde!de!dEe!dzde!dzde8e@de`e!dze!dze!dzffdфZe#e=j^jde!dAe!de8eHde!de!dEe!dzde!dzde8e@dejBdejBdejBde`e!dze!dze!dzffd҄Ze#e=jbjdAe!de8eHdEe!dzdeMdzde`e!e!ff dӄZe#e=jdjde!dAe!de8eHde!dEe!dzde8e@de`e!dze!dzffdԄZdAe!dEe!dzde!dzde!dzde!dzdKe@deMdeMde@de`e!e!e!e!dze!dzffdلZe#e=jhe/ddd۫dAe!dEe!dzde!dzde!dzde!dzdKe@deMdeMde`e!e!e!ffd܄Ze=jhjj!e6j$e=jhjj!e6j"dAe!dEe!dzde!dzde!dzde!dzdKe@deMdeMde`e!e!e!ffd݄Ze=jljj!e6j"dde8e!fdބZe#e=jpjdAe!dEe!dzde!dzde!de!deMdeMde`e!e!e!ffd߄Ze#e=jrjdAe!dEe!dzde!dzde!de!dKe@deMdeMde`e!e!e!ffdZe#e=jrjtdAe!dEe!dzde!dzdKe@deMdeMde`e!e!e!ffdZe#e=jxjdAe!dEe!dzde!dzde!de!dKe@deMdeMde`e!e!e!e!e!ffdZdAe!dEe!dzde!dzde!de!deMdKe@de!fdZe#e=j|jdAe!dEe!dzde!dzde!de!deMdeMde`e!e!e!e!ffdZe#e=j~jdAe!dEe!dzde!dzde!de!deMdeMde`e!e!e!e!e!e!ffdZe#e=jjdAe!dEe!dzde!dzde!de!deMdeMde`e!e!e!e!ffdZe#e=je/ddddZe#e=je/dddddddde!e'zdTejdzdejdzde@de@dejdzf dZe#e=je=je=jge/dZe=jjj!e6j$e#e=je/dddddAe!dEe!de!dzde!dzde!dzdKe@deMdeMfdZdZe#e=jjde!dAe!dEe!dzde!dzde!dzde!dzde!dzde@deMde8e@de!de`e!e!dze!dzffdZe#e=jjde!dAe!dEe!dzde!dzde!dzde!dzde!dzde@deMde8e@de`e!e!dze!dzffdZe#e=jjde!dAe!dEe!dzde!dzde!dzde!dzde!dzde@deMde8e@dejBdejBdejBde`e!e!dze!dzffdZe#e=je/ddd˫dAe!d(e!dEe!de!dzde!dzde!dzde!dzdeMfdZe#e=je/ddd˫dAe!d(e!dEe!de!dzde!dzde!dzde!dzdeMde!fdZe#e=je/eDdAe!de`eHeHffdZd2e!de8eHde8eHde8eHde8eHde@deHde`e!e!ffdZe#e=je/ddgdgdugdfd2e!de8eHde8eHde8eHde8eHde@de`e!e!ffdZe#e=je/ddgdgdugdfd2e!de8eHde8eHde8eHde8eHde@de`e!e!ffdZe#e=je/dddAe!de`eHeHfde`e!e!ffdZe#e=je/dddAe!de`eHeHeHfde`e!e!ffdZd2e)de)de8eHdeHfdZe#e=je/d2e)de)de8eHfdZe#e=je/dAe)de)de8eHde8eHde8eHf dZe#e=jdudde)deHdze)de)d)e'f d Ze#e=je/dudde)deHdze)de)d)e'f d Zdudde)deHdze)de)d e@d)e'f d Ze#e=jje=jjj!e6j"dd Ze#e=jde)deHdze)de)fdZe#e=je/de)deHdze)de)fdZde)deHdze)de)d e@f dZe#e=je/ddMeDd2e!de`e!e!ffdZe#e=je/ dde!de@eHzeMzde@eHzeMzdejdzfdZe#e=jddZdZdZe#e=jje#e=jje#e=jje=jjj!e6j"e=jjj!e6j$e=jjj!e6j"e=jjj!e6j$e=jjj!e6j"e=jjj!e6j$dAe!de8eHdzde8eMdzde!fdZe#e=jje#e=jje#e=jje=jjj!e6j"e=jjj!e6j$e=jjj!e6j"e=jjj!e6j$e=jjj!e6j"e=jjj!e6j$dAe!de8eHdzde8eMdzde!fdZdÐdZe#e=jje=jjge=jjj!e6j"e=jjj!e6j$e/dd ddAe!de8eHdeMdzde!fdZe#e=jje=jjge=jjj!e6j"e=jjj!e6j$e/dd ddAe!de8eHdeMdzde!fd Ze#e=jje=jjge=jjj!e6j"e=jjj!e6j$e/dd ddAe!de8eHd!eMdzd"eMdzde!f d#Ze#e=jje=jjge=jjj!e6j"e=jjj!e6j$e/dd ddAe!de8eHd!eMdzd"eMdzde!f d$Ze#e=jje=jjge=jjj!e6j"e=jjj!e6j$e/dd ddAe!de8eHd%eMdzd!eMdzd"eMdzde!f d&Ze#e=jje=jjge=jjj!e6j"e=jjj!e6j$e/dd ddAe!de8eHd%eMdzd!eMdzd"eMdzde!f d'ZeD ddAe!de8eHde8eMdzd(e@de!f d)Zd*Zd+Zd,Zd-Z dÐd.Zd/Zd0ZdÐd1ZdÐd2Zd3Ze#e=jje=jjj!e6j"e=jjj!e6j$d4Ze#e=jje=jjj!e6j"e=jjj!e6j$d5Ze#e=jje=jjj!e6j"e=jjj!e6j$d6Ze#e=jje=jjj!e6j"e=jjj!e6j$d7Zd8ZdÐd9ZdÐd:Zd;Z e#e=jje=jjj!e6j"e=jjj!e6j$d<Z e#e=jje=jjj!e6j"e=jjj!e6j$d=Z d>Z d?Ze#e=jje=jjj!e6j"e=jjj!e6j$d@Ze#e=jje=jjj!e6j"e=jjj!e6j$dAZe#e=j$je=j$jj!e6j"e=j$jj!e6j$dBZe#e=j(je=j(jj!e6j"e=j(jj!e6j$dCZe#e=j,je=j,jj!e6j"e=j,jj!e6j$dDZe#e=j0je#e=j2je=j4jj!e6j"e=j4jj!e6j$e=j0jj!e6j"e=j0jj!e6j$e=j2jj!e6j"e=j2jj!e6j$dEZe#e=j4je=j4jge/ ddAe!de8eHdFe@d"eMdzde!f dGZe#e=j0je=j0jge=j0jj!e6j$e/ ddAe!de8eHdFe@d!eMdzd"eMdzde!f dHZe#e=j2je=j2jge/ ddAe!de8eHdFe@d%eMdzd!eMdzd"eMdzde!fdIZddJZdKZdxee!dLee!dMe!de!fdNZdLe*de!fdOZeDdAe!de8eHdFe@de8eMdzde!f dPZ e#e=jBjdQe!dRe!de@fdSZ!e#e=jDe=jFge/dTZ"e#e=jHgdUZ$e#e=jJgdÐdVZ%e#e=jLgdWZ&e#e=jNgdXZ'd2e!dVe!dEe!dzdReHdaeHde`e!e!ff dYZ(e#e=jRe/ddbd2e!dVe!dEe!dzdReHdaeHde`e!e!ff dZZ)e#e=jTe/ddbd2e!dVe!dEe!dzdReHdaeHde`e!e!ff d[Z*de!d\eMde!fd]Z+de!d\eMde!fd^Z,d_e!de*fd`Z-dae*dbe!de!fdcZ.dbee!de!fddZ/deeHdFe@dTejdejfdfZ0dge!dheHdieHdFe@fdjZ1dge!dkeHdheHdieHdFe@f dlZ2dge!de8eHdFe@fdmZ3dge!de8eHdFe@fdnZ4e#e=jje/eDdge!de8eHdFe@fdoZ5 dƐdQe!dpe!dqeHdreHdFe@dse@de!fdtZ6e#e=jne/eD dǐdQe!dpe!dqeHdreHdFe@de!f duZ7e#e=jpe/deDdvZ8e#e=jre/dde>jjfdwZ9dxejBdyejBdze@de@fd{Z:e=jvjj!e6j"e=jvjj!e6j"e/d|dd}d~Z;e#e=jxje=jxjge=jxjj!e6j$e/eD ddAe!de`eHeHfdFe@deMdzdeMdzde!f dZ=e#e=jxje=jxjj!e6j"e=jxjj!e6j$e/eD ddQe!de`eHeHfdzdFe@de`eMeMfdzde!f dZ>e#e=j~e#e=je#e=jeDe/dQe!de`eHdfde!fdZBe#e=je#e=je#e=jeDe/dQe!de`eHdfde!fdZFdQe!de`eHdfdeeHeHeHge!fde!fdZGe#e=je#e=je#e=je/d dZKe#e=je/ddddddZLe#e=je/ddddZMe#e=jje=jjge/dejddddre'dTejdzdejdejdzde@f dZQe#e=jjgdejddddqe'dre'dTejdzdejdejdzde@f dZSe#e%dZTe#e=je=jjj!e6j$e/dudude>jjfdAe!dVe!dke'de'dEe!dzdReHde!fdZUe#e=je=jjj!e6j$e/dddAe!dVe!dReHde`e!e!ffdZVe#e=jj ddddde!de!d0e!deMde@de!dzd*eMdzde`e!e!ffdZXdZYe#e=jge/deDddZZe#e=je/dZ[e#e=jdZ\e#e=jje=jjgdddd2e!dTejdzde!dzde!fdZ^e#e=jje=jjgdd2e!deHdzfdZae#ejxjzjddZbe#e=je/ddddZce#e=jjddd2ejBdejdzdejBfdZddddZeddddZfe#e=je/dZge#e=jddZhe#e=jd(e!d2e!de@de!fdZie#e=jje=jjge/ddddddeHdTejdzdejdzdejdzde@dzde!f dZje#e=jje=jjge/ dddddddeHde@dTejdzdejdzdejdzde@dzde!fdZmeYe=je=jeYe=je=jLeYe=je=jVeYe=je=jeYe=je=jeYe=je=jeYe=je=jeYe=je=jeYe=je=jeYe=je=jeYe=je=jeYe=je=jeYe=je=jeYe=je=jeYe=je=j eYe=j e=jeYe=je=jeYe=je=jeYe=je=jeYe=je=jeYe=je=j eYe=j"e=j$eYe=j&e=j(eYe=j*e=j,eYe=j.e=j0eYe=j2e=je=j4jj!e6j"dd2e!deHde!fdZy(N)CallableIterable) nullcontext)Enum)partialreduce)chainproduct)Anycast) sym_floatsym_intTensorregister_decomposition out_dtype)IntLike NumberTypesuggest_memory_format TensorLikeTensorSequenceType)_maybe_convert_to_dtype_maybe_resize_out_safe_copy_out out_wrapper)_pytree)tree_map__all__ceZdZdZdZdZy) ReductionrN)__name__ __module__ __qualname__NONEMEANSUMg/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_decomp/decompositions.pyr!r!1s D D Cr+r!Fftype_promotioncompute_dtype_onlyinclude_non_tensor_argscNtjfd}|S)Nc\ r ttjjfntf}t j |i|Dcgc]}t ||r|}}tj|d i\ fd} fd} t||it||} r|St||Scc}w)Ntype_promotion_kindcJt|tr|jS|SN isinstancerto)xcomputation_dtypes r, increase_precz0type_casts..inner..increase_precOs"!V$tt-..r+cJt|tr|jS|Sr5r6)r9 result_dtypes r, decrease_precz0type_casts..inner..decrease_precUs!!V$ttL))r+) rtorchtypes_Numberpytreearg_tree_leavesr7utilselementwise_dtypesr)argskwargs allowed_typesr9 flat_argsr;r>rr:r=r/r-r0r.s @@r,innerztype_casts..inner@s.EVU[[(( )6)  ++TJ;V1W W ' !  r+c.tj||Sr5)r? full_likeselfvalues r, fill_scalarr}s ??4 ''r+r|ctjjdk(fdtj |S)Nrc,djdS)Nz@fill only supports 0-dimension value tensor but got tensor with z dimensionsrOr|sr,zfill_tensor..sRSXS\S\S^R__jkr+)r?_checkrOatencopyrzs `r, fill_tensorrs3 LL qk 99T5 !!r+r{cftjtj|dzdddz SNrminmaxr?clampr{s r, hardsigmoidrs) ;;u{{4!83 ;a ??r+cHtj|dkD|dkz|dzdS)Ngg@gUUUUUU?r?rirlr{s r,hardsigmoid_backwardrs0 ;; $y!  r+min_valmax_valcBtj||k||k\zd|S)Nrr)rlr{rrs r,hardtanh_backwardrs$ ;;DGO L r+c6tj||kd|Srsr)rlr{rds r,threshold_backwardrs ;;ty(![ 99r+negative_slopeself_is_resultc<tj|dkD|||zSrsr)rlr{rrs r,leaky_relu_backwardrs ;;taxkN.J KKr+grad approximatecd}d}d}|dk(ri||zdz}d}||z}||z} |||| zzz} tj| } d|z} d| z} d| z}d| | zz }|dd|z|zzz}| |z|z}|||zzS|}||zdz}ddtj||zzz}|tj||zd zz}||||zzzS) Ng;f?g;f?gmBP ?tanhrgHm?r"r)r?rerfrh)rr{rM_SQRT2 M_SQRT1_2 M_2_SQRTPIkBetakKappax_sqx_cuberK tanh_innerleftrightleft_derivativetanh_derivativeinner_derivativeright_derivativekAlphacdfpdfs r, gelu_backwardrs'%G&I'Jf*$s*d{/0ZZ& TzJ+j:55 AF T(9$9:/14DD)99::Y&,Q4&=112eiit d 233sTCZ'((r+inputctjtj|}tj|}||zd||zz z}|||zzSr\)r?rFsoftplussigmoid)rlrinput_tanh_softplus input_sigmoidouts r, mish_backwardrsV **QZZ%67MM%(M - 1':=P'P#P QC -3 44r+c2|tj|zSr5r?rrs r,silur's %--% %%r+cNtj|}||zd|d|z zzzSr\r)rlr{rs r, silu_backwardr.s0mmD!G  AG (<$< ==r+weightc<tj|dkD|||zSrsr)r{rs r, _prelu_kernelr6s ;;taxv} 55r+c~tj|dkD|||z}tj|dkDd||z}||fS)Nrrr)rlr{r input_grad weight_grads r,_prelu_kernel_backwardr;sE TAX{F[4HIJ++dQhTK-?@K  $$r+noiseloweruppertrainingch|r|j|S||zdz }tj||||SNr#)mulrr)rlr{rrrrrrs r,rrelu_with_noise_backwardrFs@u%%%-1,'' ~~  r+bufferc|dk}tj|dd}tj|dd}|jdkDr|n(tjtj| }||||d|zz zz zS)Nrr"rS)r?rinumelrhabs)rlr{r in_negative max_derivsignrjs r,log_sigmoid_backwardr[sx(K KA.I ;;{Ar *D,,.1$%))UYYt_4D*EA )da1q5k&:: ;;r+otherctj|jstj|jrtj n |j}|j dd|}|t j||zS)Nr*@dtype)rDis_integer_dtyperis_boolean_dtyper?float32new_fullpow)r{r two_dtype two_tensors r,ldexprhsh  ! !$** -1G1G 1S  ZZ r3i8J %))J. ..r+loss reductionc|tjjk(rtj|S|tj jk(rtj |S|Sr5)r!r(r|r?meanr)sum)rrs r,apply_loss_reductionrtsHINN(((zz$ imm)) )yy r+rc|tjk(rtjS|tjk(rtjS|tj k(rtj Syr5)r? complex32float16 complex64r complex128float64rs r, to_real_dtyper}sK }} %// !}} %"" "}} #r+targetc*||z dz}t||Sr)r)r{rrrs r,mse_lossrs 6Ma D i 00r+c||tjjk(rd|jz nd}|||z z|zS)Nr)r!r(r|r)rlrrrnorms r,mse_loss_backwardrs; #,y~~/C/C"C3 D 56> "[ 00r+ctj|||}|jtd}tj||d}tj |}tj |||S)N)rOr-infTrOkeepdim)r?softmaxeqfloatall zeros_likeri)r{rOrrmasked masked_rowszeross r, safe_softmaxrs[ --#U 3C WWU6] #F))FT:K   S !E ;;{E3 //r+rgc||z j}tj||kd|dzz|z |d|zz }t||S)Nrr#)rr?rir)r{rrrcrs r,smooth_l1_lossrsO 6M   D ;;td{C$'MD$8$t:K LD i 00r+c|tjjk(rd|jz nd}||z }t j |}||z}t j ||k||z|z |t j|zSrf)r!r(r|rr?rrir) rlr{rrrcrr9abs_x norm_grads r,smooth_l1_loss_backwardrs{ "+inn.B.B!B3 D v A IIaLE{"I ;;  A EJJqM! r+cht|||||}t||jt||dSNT copy_fromcopy_to exact_dtype)rrshaper)rlr{rrrcrXresults r,smooth_l1_loss_backward_outr s3%[$ 4 PFj&,,/ FJD QQr+deltac |tjjk(rd|jz nd}||z }t j || k| |z|zt j ||kD||z|z||z|zSrf)r!r(r|rr?ri)rlr{rrr!rr9s r,huber_loss_backwardr#s "+inn.B.B!B3 D v A ;; UF  e# AItk1E94!8k;QR r+cht|||||}t||jt||dSr)r#rrr)rlr{rrr!rXrs r,huber_loss_backward_outr%s3!dFIu MFj&,,/ FJD QQr+ ignore_index total_weightc|jdkrdnd}|tjjk(r||z }|jdk(r|jdkDr|d}|j |}t j ||k7|d}t j|} t j| ||d} | j|jcxkDrdkDrnn|j |}|Nt|jD cgc]} d} } |jd| |<|j| }||z}t j ||k7|d}| |zScc} w)Nr#rr"g) rOr!r(r|rUr?rir scatterrTrreshape) rlr{rrrr&r' channel_dim safe_targetrXrV new_shapes r,_nll_loss_backwardr.sCxxzA~!1KINN(((!L0  xxzQ6::H MMr+c ^d|jcxkrdksntd|jd|jdkDrtd|jd|jdk(xr|jdk(}|sE|jd|jdk(s&td|jd|jd |jdk7r*td |jd |jd |+|j|jd k7r td|tj j k(r|jdk(rn|jdk(r|jd|jdk(sztd|jdd|jd|jd|jdkr|jdk(std|jt|||||||S)Nrr#%input tensor should be 1D or 2D, got Dr"A0D or 1D target tensor expected, multi-target not supported, got size mismatch (got input: , target: ):expected total_weight to be a single element tensor, got: z (z elements)rSz J  INN(((TXXZ1_!Q&;+<+!+C I+J[J[I\]  T669lL r+c ^|jdk7rtd|j|jdk7rtd|j|jd|jdk(r>|jd|jdk(r|jd|jdk(s%td|jd |j|jdk7r*td |jd |jd t |||||||S) NzSonly batches of spatial inputs supported (4D tensors), but got input of dimension: rzUonly batches of spatial targets supported (3D tensors) but got targets of dimension: rr#r"rArBrDz ( z , elements))rOr0rrr.)rlr{rrrr&r's r,nll_loss2d_backwardrIes; xxzQabfbjbjblam n  zz|qcdjdndndpcq r  1 a( JJqMV\\!_ , JJqMV\\!_ ,( Jv||n M  q  &&'s<+=+=+?*@  M  T669lL r+c "|dz tjtj| |jddz|tjtj||jddzz }|||z}t ||S)Nr"r*i)r?maximumlog1prlogr)r{rrrrs r,binary_cross_entropyrNs QJ%-- TEDMM"d3 uyy b$0GHH IDf} i 00r+cd}|||z ztj|d|z z|z }|||z}|tjjk(r||j z }|S)Ng-q=r"r)r?rr!r(r|r)rlr{rrrEPSILONrs r,binary_cross_entropy_backwardrQsgG D6M *U[[T9JPW-X XF &INN((($**,& Mr+crtjtj| |z}t||Sr5)r?rLrhr)rrrrs r,soft_margin_lossrSs. ;;uyy%&1 2D i 00r+c||ztj||zdz z}|tjjk(r||j z }|Sr\)r?rr!r(r|r)rlr{rrrXs r,soft_margin_loss_backwardrUsM+%v})E)IJJINN((($**,. r+pc6tj||z |S)N)rV)rr)rrrVs r,distrXs 99UU]a9 ((r+x1x2c|jdjdd}tj|tj}|jdjdd}tj|tj}tj |j d||gd}tj |||gd}|j|j}|jdjS)Nr#rST memory_formatr) rrr? ones_likecontiguous_formatr4rmatmulmT clamp_minsqrt) rYrZx1_normx1_padx2_normx2_padx1_x2_rs r,_euclidean_distrksffQimmB%G __WE4K4K LFffQimmB%G __WE4K4K LF ))RVVBZ&12 6C ))R)2 .C ZZ F   A  # # %%r+ input_sizesstartendstepcX|j|}tj||||||Sr5) new_zerosr? slice_scatter)rlrlrOrmrnrorXs r,slice_backwardrss/&&{3J   z;UC NNr+r"cddlm}|j}|dk(r tdt j |j|}t |j}t |j}|dkr td||nd} ||ntj} | dkr| ||z } | dkr| ||z } | dkrd} n | ||kDr||} || tjk(r||} n| | kr| } n | ||kDr||} |j| ||zz} | | z } | |zdz |z||<||xx|zcc<|jr td|j||| S)Nrstatically_known_truez,slice() cannot be applied to a 0-dim tensor.zslice step must be positiver"z>r+cn|j|dtffd }||dd}|||}||fS)zn Normalize start and end such that both are in the range [0, x.get_size()[dim]] and start <= end. rPcL||S|dkr|z}tt|||Srsrr)valrrdefaultdim_sizes r, clamp_wrapz(_normalize_start_end..clamp_wrap3s0 ;N 7.C3sE?E**r+r)rint)r9rOrmrnrrs @r,_normalize_start_endr*sJwws|H+#+ ua1 -E S%8 4C #:r+srcc |tj|j|}|j|}t ||||\}}t |j}||z |dz z|z||<|j |}|dk(r||k(r|dk(r|jSdg|jz}tj||j} | |z |z||<tj||jtj} |dk7rtj| | |k\} ||k7rtj| | |k} |dk7rtj| | |z |zdk(} dg|jz} d| |<| j| } t j#| t j%|| |d|S)Nr"rdevicerrrS)rDr1rrrryexpandclonerOr?arangeronesbool logical_andviewrri_unsafe_masked_index) rrrOrmrnrorsrc_sizeindicesidxmask mask_shapes r,rrrrAs  S 1C{{3H%eS%=JE3EKK H5[D1H-$6HSM **X C zcXo$!)yy{$(6EIIK#7G ,,x 5C%KD(GCL ::hu||5:: FD z  se|4 h  sSy1 qy  e t';q'@Auyy{"JJsO 99Z D ::dD55c4!Le TTr+indexcT|j|}tj||||Sr5)rqr?select_scatter)rlrlrOrrXs r,select_backwardrjs+&&{3J    Ke DDr+offsetdim1dim2cV|j|}tj|||||Sr5)rqr?diagonal_scatter)rlrlrrrrXs r,diagonal_backwardrqs- &&{3J  ! !*k64 NNr+ input_dtypecF|j|k7r|j|}|Sr5)rr8)rlrXrs r,_cast_grad_to_input_dtyperzs&K']];/ r+outputc~||z}||tj||dzz }t|||jSNTr)r?rr contiguous)rlrrOrnew_grad_outputrXs r,_softmax_backward_datarsK "F*O 6EIIS$-$J %[*k J U U WWr+c~|tj|tj||dzz }t|||Sr)r?rhrr)rlrrOrrXs r,_log_softmax_backward_datarsA uyy0599d4 J %[*k JJr+c||dzz||dz zz }ttjtj|}|d||j d}|d||z|j d} || zS)z/Utility function to implement im2col and col2imr#r"rrrrS)rr?rint64rU) input_dkernel_d dilation_d padding_dstride_drblocks_d arange_kwblocks_d_indices kernel_grids r, _im2col_col2im_indices_along_dimrsQ&x!|)DDH EKKGI!Hh7AA!DAx*4jAKKBOK k ))r+ kernel_sizedilationpaddingrzc tjtdk(dtjtdk(dtjtdk(dtjtdk(ddd}|d|d|d d |d |jt}tj|d vxrt dddDfdt dt ddDtjt dDfd|dk(}|s|jd}|j\}} } } \} } \}}\}}\}}t| |||| |j}t| |||| |j}tj|||||f}|jdjd}|dddd||f}|jdddddd}|jd}|jd}|j|| |z|z||z}|s|jd}|S)Nr#cy)Nz"im2col(): only 2D kernel supportedr*r*r+r,rzim2col..r+cy)Nz$im2col(): only 2D dilation supportedr*r*r+r,rzim2col..rr+cy)Nz#im2col(): only 2D padding supportedr*r*r+r,rzim2col..rr+cy)Nz"im2col(): only 2D stride supportedr*r*r+r,rzim2col..rr+c|rtdDntdD}tj|fdy)Nc3&K|] }|dkD ywrNr*.0rVs r, z1im2col..check_positive..(Q1q5(c3&K|] }|dk\ ywrr*rs r,rz1im2col..check_positive..;RqAF;RrcdSNz& should be greater than zero, but got r*param param_namesr,rz0im2col..check_positive..ZL(NugVr+r r?rrrstrictconds`` r,check_positivezim2col..check_positive4,2s(%((;RE;R8R V r+rrrFrrzrrHc3&K|] }|dk7 ywrr*rds r,rzim2col..:!qAv:rrc dtS)NzmExpected 3D or 4D (batch mode) tensor for input with possible 0 batch size and non-zero dimensions, but got: tuplersr,rzim2col..--25\N<r+c3\K|]$\}}}}}d|d|zz||dz zz dz |zz&yw)r"r#Nr*)rrpaddilkersts r,rzim2col..sF "Cc3 S1s7]SC!G_ ,q 0R 77s*,r^c3&K|] }|dkD ywrr*)rcs r,rzim2col..s 'aAE 'rc Fdtdddddddd S) Nz!Given an input with spatial size r^, kernel_size= , dilation= , padding= , stride=z9, the calculated shape of the array of sliding blocks is z*, but its components must be at least one.r)rr output_sizerrrzsr,rzim2col..sL3E%*4E3FG"m;xj9)9VH-]D Fr+rHrrSr"rT)r?rrrr rziprUrrrrpermuter2r*squeeze)rrrrrzrr batched_input batch_dimr+input_hinput_wstride_hstride_w padding_h padding_w dilation_h dilation_wkernel_hkernel_wblocks_row_indicesblocks_col_indices padded_inputrnum_blocks_rownum_blocks_colrrs ```` @@r,im2colrs~ LL[!Q&(TU LLX!#%ST LLW"$QR LLV!#OP ; .8Z(8Yu568$ KKE u:D LL :3:uRSz:: < &) "#J;' K  LL '; '' F FAIM "/4{{,I{GWHh"Iy%J $Hh9:y(ELL::y(ELL 55Iy) LML+55b9CCBG !Q 24FF GF ^^Aq!Q1 -F',,Q/N',,Q/N ^^;)H4n~6UF " Mr+rc !"tjtdk(dtjtdk(dtjtdk(dtjtdk(dtjtdk(ddd}|d |d |d d |d|d|j"t"}tj|dvxrt d"ddD"fdddz}tj"d|zdk("fdt D cgc]"\} } } } } d| d| zz| | dz zz dz | zz$}} } } } } |d|dz!tj"d!k(!"fdtj!dkD!"fd|dk(}|s|j d}|j"\}}\}}\}}\}}\}}|j"d"d|zgz|z}|jdddddd}t||||||j}t|d}t||||||j}t Dcgc] \}}|d|zz}}}|j"d"dtzg|z}dd||f} tj|| |d}t!j"|| | | | f}|s|j%d}|Scc} } } } } wcc}}w)Nr#cy)Nzonly 2D output_size supportedr*r*r+r,rzcol2im..rr+cy)Nzonly 2D kernel supportedr*r*r+r,rzcol2im..rr+cy)Nzonly 2D dilation supportedr*r*r+r,rzcol2im..rr+cy)Nzonly 2D padding supportedr*r*r+r,rzcol2im..rr+cy)Nzonly 2D stride supportedr*r*r+r,rzcol2im..rr+Tc|rtdDntdD}tj|fdy)Nc3&K|] }|dkD ywrr*rs r,rz1col2im..check_positive..rrc3&K|] }|dk\ ywrr*rs r,rz1col2im..check_positive..rrcdSrr*rsr,rz0col2im..check_positive..rr+rrs`` r,rzcol2im..check_positiverr+rrrFrrzr)r#rc3&K|] }|dk7 ywrr*rs r,rzcol2im..%rrr^c dtS)NzmExpected 2D or 3D (batch mode) tensor for input with possible 0 batch size and non-zero dimensions, but got: rrsr,rzcol2im..&rr+rr"cdddS)Nz|Expected size of input's first non-batch dimension to be divisible by the product of kernel_size, but got input.shape[-2] = r^z and kernel_size=r*)rrsr,rzcol2im..,s#==B2YKH"m%r+rSc:ddddddddd S NzGiven output_size=rrrrz , expected input.size(-1) to be but got rS.r*Lrrrrrrzsr,rzcol2im..9F$[M }M:Zy &B))*9U2YKqBr+c:ddddddddd Sr!r*r$sr,rzcol2im..?r&r+rrHr accumulater)r?rrrr rrUr*rrrrWrqprodr_unsafe_index_putrrr)#rrrrrrzrrprod_kernel_sizerrrrrcolrout_hout_wrrrrr r r r  indices_row indices_colorVoutput_padded_sizerrr%rs# ````` @@r,col2imr4s LL[!Q&(OP LL[!Q&(JK LLX!#%IJ LLW"$GH LLV!#EF ; .8Z(7Ie468$; . KKE u:D LL :3:uRSz:: < #1~ A6 LL b $$) %'* (K'    "Cc3 S1s7]SC!G_ ,q 0R 77 C  AQA LL b Q B B  LL A B B AIM " KKELE5Hh"Iy%J $Hh MM58U1X1A%AB[PSVV WE MM!Q1a +E2 xY%,,K$K3K2 xY%,,K14K0IJ1!a!e)JJ __ q58tK0014FFF {K 0C  # #FC4 # HF UU6YJ YJ K LF " Mk VKs 6'M 9Mrcz||j||zzjtj|}|SNr\)type_asrrDr)rlrrnrJs r,native_dropout_backwardr8hsB  [1E9 :AA11+> B A Hr+ input_size dimensionr2ctjdkDfdtdk(rtjdSt j t|j dtffd }k\r|Sddlm dtffd }|r|Sfd}fd}d td ttdtffd } } j} ttD]} || | z zd z zd} || | }tj |j"tj$ }|z| z}j'd| j)d|}t}||<| ||}t*j-||j/S)NrcddS)Nzstep is z but must be > 0r*)rosr,rz!unfold_backward..|s8D69I)Jr+rPc\tjjtj}|j dj }j ddzj dz}dz|fz}tj||djS)NrrrSr"r5Tr() r?rrint32unfoldflattenmovedimrr+r) rsourcerrOrrXr9r2ros r,_vectorized_index_putz.unfold_backward.._vectorized_index_putsll:c?4;;ekkRjjD$'//1b#'*223a@# &%% v$& *, r+ruc jtdzk7ryjdk(syz zdz}j|k(sytfdt DS)Nr"FrSc3fK|](\}}|k(xsj||k(*ywr5r2)r source_dimrrOrrvs r,rzDunfold_backward.._is_canonical_grad_shape..sF % H #  H$TYYz%:h%FG H s.1)rOrr2r  enumerate)n_foldsrOrr9r2rvros r,_is_canonical_grad_shapez1unfold_backward.._is_canonical_grad_shapes 88:Z1, ,$TYYr]d%:;c?T)d2Q6$TYYs^w%>? )2*(=   r+cd||kr|S||kr|Stj||Sr5)r?sym_minabrvs r,_sym_min_if_neededz+unfold_backward.._sym_min_if_needed6 a (H a (H}}Q""r+cd||k\r|S||k\r|Stj||Sr5)r?sym_maxrMs r,_sym_max_if_neededz+unfold_backward.._sym_max_if_neededrQr+rB target_shapect|}|j}||kDr|j|S||z }dg|z}t|D]s}|k(rn||z}|dks||k\r|j|cS ||dk(r|j|||<Mt j |j|dk(du|j|j|S)Nr"rcy)Nz=unfold_backward grad shape is not broadcastable to input_sizer*r*r+r,rz:unfold_backward.._reshape_source..rr+)rrOrrTrr?rr*) rBrUr source_ndim dim_offset view_shaperG target_dimrOrvs r,_reshape_sourcez(unfold_backward.._reshape_sources< jjl  ==. .K' S4Z  , J *c 1zJ7NJA~t!3}}\22$Z %;q%@A)/j)A :&  LL,1$  ~~j)00>>r+r"rrS)r? _check_valuer squeeze_copyrDr1rqrrwrvrryrr2reversedrTrrr>selectr3r index_addr)rr9r:r2rorCrJrPrTr\input_dim_size grad_dim_sizer max_valid_len valid_lenrrB source_shaperOrXrvs`` `` @@@r,unfold_backwardrgws  tax!JK :!!!$**  Z) &> uDDLL'#|  A '3 /E99QY'L EIIlCF FF <( Mr+c ddlm}|jdk(rtj||dSt j |j|}||j|dkDrtj||dSt|j}d||<tj||j|tdf|}tj||dS)Nrguard_or_falseTr}r"rr) rwrrr?amaxrDr1rryr4rr )r9rOr canonical_dim pad_shapes r,rr&sDvv{zz!S$//**16637Maggm,q01zz!S$// QWW I Im 1ajjE&M:;EA ::ad ++r+)rc|j}|r6|jtjk7rt d|jdt j |t jj\}}|j|}t||}||z }tjtjtj||d}||z }|s|j|}|Sr{)rrr?r~r0rDrErrr8rrMrrh) r9rOryr:r=rshiftedshifted_logsumexprs r, _log_softmaxr:s A 77ejj  7y@UV ',&>&> uDDLL'#|  A '3 /E%iG %))EIIg,>T"RS ( (F <( Mr+r padding_idxscale_grad_by_freqsparsec|jdk7rtd|jd|jdkr4|jd|}|jdk(r|j d}|S||S)Nr#z'weight' must be 2-D, got z-Dr"r)rOr0r index_selectr)rrrrrrs r, embeddingrRsszz|q9&**,rJKK||q!!!W- <<1 ++a.C gr+ num_weightsctj|tjj\}}|j |}t |t j}|dk\||kz}|rm|j|f}t j|} tj|||g| }tj|||gd} || jdz }t|||k7z|j} |j|f|j |jdz} tj| | |g|j |S)Nr|rr"rS)rDrErrr8rr?longrqr_r#_unsafe_masked_index_put_accumulaterrUrWrr) rlrrrrr:r= valid_indicescountsrgrad_weights_scaler grad_weights r,embedding_dense_backwardrhsD',&>&>)N)N)V)V'#|..!23K%guzz:G\g &;2w'99 MG9d "66 MG9a "$6$@$@$DD ]g.DE{GWGW XD'' **7<<>::K  3 3TG9k br+c"d}|D]}||z} |Sr\r*)r9rJis r,r*r*s$ A  Q Hr+tensors num_chunkscZg}|D]}|j}|||zdz |z|z}|||k7r;dgdz|j|z dz zd|||z gz}tj||d}|d|t j |dgz}|j |j||S)Nr"rr#rS)r2rrconstant_pad_ndr?Sizeappendr*) rrOrpadded_tensorstensor tensor_size pad_along_dimr view_sizes r, _pad_chunkrs N 9kkm $S)J6:zIJV K, ,#'V[[3.23 C 007C))&#q9F% J3C(DD fnnY78 9 r+cR|dj}|D]}|j|k7syy)NrFTr)rrrs r,have_same_ndimsrs2 1:??D ;;$  r+c|djd|}|D]-}tj|jd||k(d/y)Nrcy)NzG_chunk_cat expects same sizes of 0,...,dim-1 dimensions for all tensorsr*r*r+r,rz+leading_dimension_matches..rr+)r2r?r)rrOleading_dim_sizesrs r,leading_dimension_matchesrsN )$3/  KKM$3 #4 4 ]  r+ctj|dk\dtjt|dkDd|dj}|dj}|D]r}tj|j dkDdtj|j|k(dtj|j|k(dtt |r(tj|dj|}nEtj|dk\d|D]&}tj||jkd (t|||S) Nr"cy)Nz&_chunk_cat expects positive num_chunksr*r*r+r,rz._preprocess_chunk_cat_inputs..rr+rcy)Nz0_chunk_cat expects a non-empty input tensor listr*r*r+r,rz._preprocess_chunk_cat_inputs..rr+cy)Nz#_chunk_cat expects non-empty tensorr*r*r+r,rz._preprocess_chunk_cat_inputs..rr+cy)Nz8_chunk_cat expects all input tensors with the same dtyper*r*r+r,rz._preprocess_chunk_cat_inputs..rr+cy)Nz8_chunk_cat expects all inputs tensors on the same devicer*r*r+r,rz._preprocess_chunk_cat_inputs..rr+cy)NzK_chunk_cat expects non-negative dim when input tensors have different ndimsr*r*r+r,rz._preprocess_chunk_cat_inputs..rr+cy)Nz3_chunk_cat expects dim < ndim for all input tensorsr*r*r+r,rz._preprocess_chunk_cat_inputs..rr+) r?rrrrrrrDr1rOrr)rrOrexpected_dtypeexpected_devicers r,_preprocess_chunk_cat_inputsrs-  LLq"RS LL G qTQZ%%Naj''O   V\\^a')VW LLN * N   MM_ , N  w$$WQZ^^%5s; 1H a  F LLfkk!M   gs+ Jr+rct|||}t|||}|tj||dzStj||dz||S)Nr")r)rrr?r4)rrOrrrs r, _chunk_catrsS 'wZ @Cj9N {yyq11 .#'s3 r+ split_sizesctj|||}|.|Dcgc]"}|jtj$c}St ||D])\}}t ||jt||d+ycc}w)Nrr\Tr) rsplit_with_sizesrr?r`rrrr)r{rrOrsplitssrsplits r,split_with_sizes_copyrs " "4# " >F {HNO1e&=&=>OO f- NMFE fekk 2 UF M N Ps'B split_size.cDtjj|||Sr5)rrr)rrrOs r, unsafe_splitrs ::  UJ 44r+cDtjj|||Sr5)rrr)rrrOs r,unsafe_split_with_sizesrs  ( ( S AAr+cz|j}||tjdk\fddk(r|jfStjdkDfdzdz z}ddlm}||}t |Dcgc]}}}|zz z |d<tj|||Scc}w)NrcdS)Nz=split expects split_size be non-negative, but got split_size=r*)rsr,rzsplit..sOPZ|\r+cdS)NzKsplit_size can only be 0 if dimension size is 0, but got dimension size of r*)rsr,rzsplit..s%%-J0r+r") guard_intrS)rr?rdetachrwrrTr) r{rrOrlchunksrrrrs ` @r,rr s**K3H LLa\1}  LLQ 0 #a'J 6F@ v F',V}5!:5K5 J$7($BCKO ;;t[# ..6s B8tensor_indices_or_sectionsc |jjdk7rtd|j|jtj k7rtd|j|j t j dk(xs dk( fd dk(rS|j}t|ts!tdt|j|j||St}tjjx}r|j x}r |j"}|5|Dcgc]}|j}}ddd|j|Scc}w#1swY xYw)Ncpuz/tensor_indices_or_sections must be on CPU, got z.tensor_indices_or_sections must be int64, got r"rcddS)Nz{tensor_split expected tensor_indices_or_sections to be a zero-dimensional or one-dimensional tensor, but got a tensor with z dimsr*) split_dimsr,rzAtensor_split_tensor_indices_or_sections_py_impl..:s<X>_>_=` a  "''5;;6<=W=]=]<^ _  +..0I LLQ()q. M A~-224(G, 7X8O8O7PQ   3//779 9I 9",, ,I ,99C U E)CDAqvvxDGD E  #..E E Es:E7?E2E72E77Fmat1mat2c|js&|jst|}t|}|tj||z}|dk(r|S|||zzSrs)ru is_complexrr?mm)r{rrrcrmrs r,addmmr]s]  ! ! #DOO,=4yE  %((4& &C qy   r+rct|tj|||z}|dk(r|S|||j|zzS)Nrr)r?rr8)r{rrrrcrmrs r, addmm_dtyperqs@ %((4; ;C qy  ** **r+use_geluct|||||}|r8|jrtj|dStj|Stj |S)Nr)r)ris_cudargelurelu)r{rrrcrmrrs r,_addmm_activationrsO dD$ .C <<99Sf95 599S> ! 99S>r+vecc|js&|jst|}t|}|tj||z}|dk(r|S|j dk(r||zS|||zzSrs)rurrr?mvr)r{rrrcrmrs r, _addmv_implrsq  ! ! #DOO,=4yE  %((4% %C qy  yy{ad{  r+ctjjjk(xrjjk(fdt||S)NcVdjdjdjS)Nz6addmv input tensors must have the same dtype, but got , z, and r)rr{rsr,rzaddmv..s1H TVW[WaWaVbbhilirirhstr+)r?rrr)r{rrrcrms``` r,addmvrsK LL djj .s'A }I>r+c.dddjS)NzExpect mean to have shape (rz , but got r)rrrsr,rz,native_group_norm_backward..s-aS5'DJJ<Pr+c<ddjSdS)NzExpect gamma to have z elements but got rS)r)rrsr,rz,native_group_norm_backward..s-'s*.s,QC/[\a[bcr+r#rrrgrSr"rrH)rDcheck_same_devicecheck_same_shaper?rrrrrrrUr*rrrWr8r)rlrrrrrrrrrcpgdsdbd_inputd_gammad_biasrds_valdb_valc1c2c3s ` ````` r,native_group_norm_backwardrs UD$ 5+N 4F LL Q$> LL q%j P LL  +!+i u*C LL S5[c ; & + +Aq# 6 : :s : CB   !Q $ ( (aS ( 1B!G!G F1~ 39   YYr5??1#56>>q%MQQRSTFYYr5??1#56>>q%MQQRSTFr" a,B ZZ5#.2215FZZ5#.2215Fr" Auc?4;;?Btmf$ ,t 3d :Q >S4Z&4-!+ + \\"  r1 % r1 % IIk))!UC=r Bii aS92> ?   //%++.11%++>1~E3'"''!UC*@4>>RTCU*UU..$%SaSS\ WQZ 1~QC Wf %%r+out2c  t|||||||||| } | | | f}t| D]2\}}| t|||jt |||d4|Sr)rrHrrr)rlrrrrrrrrrrrrsrrrXrrJs r,native_group_norm_backward_outrsz"(UD$q!S%Fd#J&!Q1 = jmQWW 5 Q 1 4 PQ r+c,||j|S|Sr5r8)r9rs r, _maybe_castr"s}ttE{ Hr+grad_outnormalized_shapebiasc~"|j}|j} tj|j""fd||||fD\} } } } | t d| t |z }||d}|d|}g}g}t| D]*}||k\r|j||j|,t|}t|}ddl m }||dk(s ||dk(rN|dr|j|nd|dr|j||dnd|dr|j||dfSdfSt|| j}t|| j}| t d| |z |z}| | | z}n| }||z}tj||d}tj ||}tj||d}tj ||}||z |z }d}d} d}!|dr||z |z}|dr0| .t |dkDrtj| |z|d } n| |z} |dr8| 6t |dkDrtj| |d }!n| j#}!t%||jt%| | |jndt%|!||jfSdfS) Nc3hK|])}|!|jtjn|+ywr6)r8r?r`rr9r:s r,rz-native_layer_norm_backward..7s>9  =  e.E.EF  9s/2 grad_out_cast should not be Nonerrur"r#zinput_cast should not be NoneTF)rrOrDget_computation_dtyperr0rrTrr*rwrvrqrWr?rrrr)#rrrrrrrr input_shape input_ndim grad_out_cast input_cast weight_cast bias_castaxis inner_dims outer_dimsinner_dim_indicesouter_dim_indicesrrMrvx_hat grad_x_hatrNrOrrrrKr d_weightr r:s# @r,native_layer_norm_backwardr0)s++KJ33EKK@9E640 95M:{I ?@@ ,- -DTU#JUd#J#%#% : ( 9  $ $Q '  $ $Q ' ( ZA ZAKQ!V$(=a1f(E,7NEOOK (3>q>EOOK. /t3>q>EOOK. /  HL  T:>>#3 4D T:>>#3 4D<== $ $ &E"[0 " QA */6A :u %B 2($ /B 5" B EBJE!G"H F1~!8u$1~+1  !A %yy!68I5QH$u,H1~)/  !A %YY}.?GF"((*F GU[[)Hf.@flldKF$*:DJJE AEE r+c t||||||||} || | f} t| D]2\} }| t| | |jt || | d4| Sr)r0rHrrr)rrrrrrrrrrrsrrrXrrJs r,native_layer_norm_backward_outr2sw(%)4vt[Fd#J&!Q1 = jmQWW 5 Q 1 4 PQ r+c&g}tt|D]'}|j|j|z dz )t j |j }|j|}|~|tjtjfvr.tjtjj}n0tjtjj}n|}tjtjj j"j%tj&|dj)|d|} |j+| } || j+|} |j,xs|duxr |j,} t j.|} | tj0tj2fv} | s"| s | j5} | j5} | j7|}|| fS)Nr"r#Tr)rTrrrOrDr rr8r?rrfinforhrrsqrtopsraddScalarrrr is_nestedr channels_lastchannels_last_3drr7)rrrrhdims_to_reducerr:upcasted_inputeps_val rqrst_inputupcasted_resultr9r]is_channels_lastrs r,_fused_rms_normrBs!#N 3'( )3eiikAo12333EKK@XX/0N {  @ @kk%--044Gkk%--044G++  !! IIna ( - -.$ - OQX K%((5O )--f5LF$$6$K6;K;KI//6M$  ) -)446!,,.  $ $U +F ; r+cL|j}|j}tj|j}|j |t j} |j |t j} |!|j |t jnd} | td|t|z } || d} |d| }g}g}t|D]*}|| k\r|j||j|,t| }t|}ddl m}||dk(s ||dk(r4|dr|j|nd|dr|j|| dfSdfSt!|| j}| | | z}n| }d}d}| |z}|dr)t j"||z|d}|||z |zz |z}|dr0| .| |z}t|dkDrt j"||d}n|}t%||jt%||jfS) Nr\rrrr"TrF)rrOrDr rr8r?r`r0rrTrr*rwrrqrWrr)rrrrrrr!r"r:r#r$r%r'r(r)r*r+rrr,rr.r r/r-sum_vald_weight_full_shapes r,_fused_rms_norm_backwardrFsa++KJ33EKK@KK)@)@ M+5;R;RSJ    #53J3J K  ?@@ ,- -DTU#JUd#J#%#% : ( 9  $ $Q '  $ $Q ' ( ZA ZADa1fQ!7,7NEOOK (3>q>EOOK. /  GK  T:>>#3 4D"[0 " !G"H  E1~))EJ.4EtTg 55=1~+1+e3  !A %yy#):EH+H GU[[)Hekk* r+ running_mean running_varmomentum functionalc Ddgttd|jz} tj|j } |} |} |r"tj|j } |j | } tj| | dd\}}tj||z}||z |z}tj|| }tj|| }|!||zd|z |zz} |s|j| |;|j|jdz }tj|| }|||dz z z}||zd|z |zz} |s|j| n|| td|j | d}|} |j | d}|} |}dtj||zz }|j j"d k7r|}|}n"|j%d }|j%d }t'||jdz }t'||jdz }||z |z}|2|j)}t'||jdz }||z}|2|j)}t'||jdz }||z}|j j"d k(r8|j |j }|j |j }|j |j ||| | fS) Nrr#rT)rO correctionrr"z:running_mean and running_var must not be None in eval mode)rrrr)ryrTrOrDr rr8r?var_meanr5rcopy_rrr0rdrrrqrWr@)rrrrGrHrrIrhrJreduction_dimsr:new_running_meannew_running_var input_acc biased_varrrr save_mean save_rstdn squeezed_var unbiased_varinvstds r,native_batch_norm_helperr[sS4a 566N33EKK@#!O!77 DHH#4H5  >> >a  D{{:+,$,$&MM$7 MM$7  #')3q8||6SS ""#34  "  A.A!==^DL'1A;7L&5X8TTO!!/2  ;#6 L $->TJ '!nn+<4nH %ejjs!234 <<   %$II-I-I uyy{Q7"6599;?;$,&( !"6599;?;& ||~ uyy{Q7$ ||E!LLu{{L3 LLu{{L3   $  r+rU save_invstdc >t||||||||d \}} } } } || | fSNFr[ rrrrGrHrrIrhrrUrVrVs r,native_batch_normrars=*B vt\;(CQV*&FIy!Q 9i ''r+c ||tj||||||S| td| td|rtj||||||||Stj|||||||S)Nz`running_mean is None, but running_var is provided. They should both be None or both be provided.z`running_var is None, but running_mean is provided. They should both be None or both be provided.)r_native_batch_norm_legitrx$_native_batch_norm_legit_no_training)rrrrGrHrrIrhs r,native_batch_norm_decompositionres 3,, 648S   <   <  ,, 64{HhPS  88 64{Hc  r+c~|j|}||zdz |z}|dk(rc|dk(r^t|Dcgc]}|}}|||z|z z ||dz <tjjj j |||Stjjjj|||Scc}wNr"r) r2rTr?r6rrrrr)rrrOrrrVrs r,unsafe_chunk_py_implrhs{{3HV#a'F2JQ8q=+0=9az9 9", V0Ch0N"O FQJyy~~55==fkSVWW 99>> & & - -fj# FF:s B:c Ntjj|||||d||Sr^)rrcr)rrrrGrHrIrhs r,rdrds5  ( ( 0 0      r+c >t||||||||d \}} } } } || | fSr^r_r`s r,rcrcs=*B vt\;(CQV*&FIy!Q 9i ''r+c >t|||dd|||d \}}}} } |||fSr^r_) rrrrrIrhrrUrVrVs r,!_native_batch_norm_legit_no_statsrls<*B vtT48S%*&FIy!Q 9i ''r+c vt||||||||d \}} } } } | td| td|| | | | fS)NT#new_running_mean should not be None"new_running_var should not be None)r[r0) rrrrGrHrrIrhrrUrVrQrRs r,#_native_batch_norm_legit_functionalrpsl" ! vt\;(CQU  BCCABB 9i)9? JJr+c Ttjj|||||d|}d}|tjjjk(r tjj ||}tj |tj|j|jS)a Return a reserve tensor for batch norm, used only by cudnn to pass forward state to the backward pass. This is needed for `_batch_norm_with_update` and `_batch_norm_no_update`, which support a variety of backends including cudnn. We create this tensor here to get the correct shape in the traced graph if we detect that will call the cudnn kernel, and rely on DCE to avoid materializing this tensor. Tr)rlayoutr) r?_C_select_batch_norm_backend_BatchNormBackendCudnn(_get_cudnn_batch_norm_reserve_space_sizeemptyuint8rrr) rrrrGrHrhrbackend reserve_sizes r,_get_batch_norm_reserve_tensorr| s hh11 vt\;cGL%((,,222xxHH 8  ;;EKK U\\ r+c dt|||||d||d \}}} } } t||||||d} ||| | fS)NTFrr[r| rrrrGrHrIrhrrUrVrVreserves r,_batch_norm_with_updater- sa*B      *&FIy!Q- vt\;dG 9i 00r+c t|||||d||d \}}} } } t||||||d} | td| td||| | | | fS)NTr~rnro)r[r|r0) rrrrGrHrIrhrrUrVnew_rmnew_rvrs r,"_batch_norm_with_update_functionalrH s ! vt\;hT  - vt\;dG~BCC ~ABB Iy'66 BBr+c dt|||||d||d \}}} } } t||||||d} ||| | fS)NFr~rrs r,_batch_norm_no_updatere sa*B      *&FIy!Q- vt\;eG 9i 00r+c|td|tj|tjj\}}|j |}t j||kj t j}t jd||j|z }|j ||z|z}|j ||fS)Nz=generator must be None for _fused_dropout decomposition, got r|rrgr) r0rDrErrr8r?rvry scalar_tensorr) rrV generatorr:r=rSrrnrxs r,_fused_dropout_decompositionr sKI; W  ',&>&> 5#H#H#P#P'#|*+I OOI & * . .U[[ . AD   +\}}|dk(s ||jkr|j||k(r.|j|}@|Sr\)rHrrrU)r9broadcast_maskr'rs r,_broadcast_batch_norm_backwardr sO/" d 19dQVVm 0E D!A" Hr+rc *t|||||||||| Sr5)native_batch_norm_backward) rrrrGrHrUr\rnrhrrs r,batch_norm_backwardr s/ &     r+c &|j} | |j} n| } tj|j&&fd|||||||fD\} } }}}}}|j}|j }|dkrt d|d}t t|||z }|}|}|r||4t d|| t d|}tj||z}dg|z}||||<g}t|D]}||k7s |j|t||}d|z }tj| |}tj| | |z z|}t||z|}ttj||z||z|} |t||dz}!nt||z|}!|r| |z | z}"| |"z |z |!z}#n| |!z}#| dr||z}$nd}$| dr|}%nd}%|#j| t!|$| t!|%| fS)Nc3HK|]}||jn|ywr5rrs r,rz-native_batch_norm_backward..B s,  $%= a7 s"r#z*rank of the input must be at least 2, got r"z1mean and invstd must not be None in training modezDrunning_mean_cast and running_var_cast must not be None in eval moderg)rrDr rrOr0r*ryr?r5rTrrrrr8r)'rrrrGrHrUr\rnrhrr weight_dtyper#r$r%running_mean_castrunning_var_castsave_mean_castsave_invstd_castr! input_rankr' num_featuresrrZrreduction_axesrrgrad_output_sumdot_p grad_mean proj_scale grad_scaleprojrXr grad_biasr:s' @r,rr' s++K || " 33EKK@          ++KJA~I*VWW D[)*[->>L D F <6> !TU U  $(8(@ V !-34!"j 0N&t,N4 "N : % 9  ! !! $% *$ ?D  Dii ~>O IImzD'89> JE./E~VI/ %$,0J 3FNKcQ 3 [ .  T!Z/$t+y8JF "Z/ 1~fn  1~#    k"K.I|, r+c  t|||||||||| } | | | f}t| D]2\}}| t|||jt |||d4|Sr)rrHrrr)rrrrGrHrUr\rnrhrrrrsrrrXrrJs r,native_batch_norm_backward_outr s"(    Fd#J&!Q1 = jmQWW 5 Q 1 4 PQ r+save_varc Btj|||||||d|gd SNT)TTTrr)rrlrrGrHrUrrs r,miopen_batch_norm_backwardr s5  * *    r+ reserveSpacec Btj|||||||d|gd Srr) rrlrrGrHrUrrrs r,cudnn_batch_norm_backwardr s5  * *    r+c|j|jttjdvfd|jddD]}tj|dk7fd d|dzdk(rxd|dzdk(rjt dt dd|D}t dt dd||D}tjjj|||Sd d fd }|d|d\}}}} |d|d\} } } } |d t|d | f}| s| stj|dSd}||||| d\}}||| | | d\}} d}tt|jdt|jdD]!\}}| |d |dd|f}||d |dd|fz}#||| zz S)NrcdS)Nz9adaptive_avg_pool2d(): Expected 3D or 4D tensor, but got r*rsr,rz%adaptive_avg_pool2d.. sKD6Rr+r^rc"dtdS)Nzjadaptive_avg_pool2d(): Expected input to have non-zero size for non-batch dimensions, but input has shape r#rrsr,rz%adaptive_avg_pool2d.. s99>uaIr+rSc3,K|] \}}||zywr5r*)rrr2s r,rz&adaptive_avg_pool2d.. sG$!QqAvGc3:K|]\}}}||dz |zz ywr"Nr*)rrr2rs r,rz&adaptive_avg_pool2d.. s' '1aAQ! O sc8tj||z|dS)Ntrunc rounding_moder?divrNrOrs r, start_indexz(adaptive_avg_pool2d..start_index syyQ99r+cJtj|dz|z|zdz |dS)Nr"rrrrs r, end_indexz&adaptive_avg_pool2d..end_index s&yy!a%1q1,awGGr+ctj| tj}|||}||zdz}||z}|dk(xs||zdk( }|r|dz }n |dk(r|dz}tj| tj}|jd|z}|rUtj|dz |j |j } tj|| } |||} | |z } n|} || ||fS)Nrr"rrSr)r?rrrUrrrminimum)in_sizeout_sizeorangei0 maxlength in_size_modadaptive range_maxrmaxvali1lengthrrrs r, compute_idxz(adaptive_avg_pool2d..compute_idx s hvU[[I 7 3x'!+ ( #q(GH{,Ba,GH  NI A  NILL6M ll2* ((! 399SZZF--V,C68W5B"WFFFIx//r+.rH)rrSrct|tr||fS|dk\rtd|||jdk\}|dk(r t |d}t j ||d}t || }||fS)Nrz)dim should be negative when masking, got rSr^rHr)r7rr0rUrWr? masked_fill)valsrrrrOrs r, maybe_maskz'adaptive_avg_pool2d..maybe_mask9 s fg &< ax$'PQTPU%VWW 0 0 44Dby(q1$$T45D&vt4F< r+)rrOr)rrrr?rrrnnrJ avg_pool2drWrr rT)rrrrzkernelridxhlength_h range_max_h adaptive_hidxwlength_w range_max_w adaptive_wrrretrjrrrrrs @@@@@r,adaptive_avg_pool2dr s. \\F KKE u:D LL R[[   F I   Ry;r?"a'E"I B,G1,LG#eBCj+*FGG +.uRSz;+O  xx""--eVVDD:H0@/:%)[QS_.U+D(K.9%)[QS_.U+D(K 'a0$6 7D jzz$H--   h jbND(  h jbND( Cdjjn-uTZZ^/DE+1 ;sAq!|$CS!Q\**C + (X% &&r+ ceil_moden_dimc  ddtttzdttdzdttffd }||}|||} ||} ||} tjjdzdzfvfdjdzk(} | sj d tj  d} j}tDcgc]5}tjj| |||| || || ||7}}j}d }|tjurd }n|jr td }ntj|}|j d kr j#tj$n.|j dkrj#tj&|j dkrt|j(dz nt|j(}d}g}t+tD][}| |}||dz | |z||dz | |zzdz}t-d || |d| |zzz | |z}|j/||g]t1j2||}g}tD]}tj4|||tj&j d| |ztj4|||tj&j d | |zz}dgdzz}|||d|z<|||d|zdz<|j7|j9||t:g|}d}tt|}|tDcgc] }|d|zz c}z }|tDcgc] }|d|zzdzc}z }|j=|}t|j d|z} |j8g| d}|j-d\}!}"g}#|"}$tD]9}d}%t|dzD] }&|%||&z}% |#j7|$|%z|$|%z}$;g}'tD]w}dg|zz}|||||z<tj4|||tj&j9|}(|'j7|(| |z|#|| |zz| |z y|'d })tdD]}|)| |z|'|z})|r|!j#|}!| s"|!j?d }!|)j?d })|!|)fScc}wcc}wcc}w)NrrrPct|ttfs|gzS|stj|dud|St |dk(r|zSt|S)Ncy)Nzdefault must be providedr*r*r+r,rz<_max_pool_nd_with_indices.._expand..g rr+r")r7ryrr?rr)rrrs r,_expandz*_max_pool_nd_with_indices.._expandc sV#e}-55=  LL,.P QN!#h!msU{:c:r+)rr"r#c<ddzddzdjdS)Nz Expected r"D or r#z D input, got r?r)rr{sr,rz+_max_pool_nd_with_indices..r s()EAI;eEAI;mDII;aPr+rFrr @TrrrSrr5) ryrr?rrrUrrrT_meta_registrationspooling_output_shaperrrur iinfobitsr8r>rrr_rextendrrrrr*Ellipsisrr)*r{rrzrrrrrksrpadi is_batchedrlrr output_sizesrpromoted fill_valueinfopad_argsrneededrr9 idx_tensorsrrwindowsn_batchperm out_shapevalues local_argmax kernel_pos remainingdivisordd orig_coordsout_pos flat_indicess*` ` r,_max_pool_nd_with_indicesrZ s;T#Y_;tCy4/?;4PS9;  B  $B  B  B LL eai++P eai'J ~~a tzz5&'*+K [[F u   !!66 NBqE2a5"Q%A  L JJEH    6] {{5! 99r>775;;'D YY^775;;'D,0IINU488a<(dhh H eEl #'!uq/A%A."Q%!)r!u1DDqHAvQ!be)!;<=1Eu & ' dHJ/A K 5\ / LLau{{ K U UVW Xe ll2a5u{{CMMaPSUVWSXX Y q5y!#Aa!e a5a!eai3;;u-. /(K()GG g DeEl 3Wq1u_ 33D%, 7QWq1u_q 77Dood#GW]]#4Wu_56Igoo-y-"-G";;2;.FLJI 5\(Au% B r"v G )w./' (K 5\Lw')!_gk,, OF%++ '%.  7RU?Z]RU-BBRUJK Lq>L 1e_F#k!n4{1~E F5! "#++A. < Cr 4 7s":UU9U c &t||||||dS)Nr#rrr{rrzrrrs r,max_pool2d_with_indicesr"  % k67Hiq r+c &t||||||dS)Nrrr r!s r,max_pool3d_with_indicesr% r#r+c j tj dv fdtd D]. tjj dkD fd0j d}j d}|\}}|dk(s|dk(rOt j dd||gz}j|j|tjfS|dk(re|dk(r`tjd d }tjjdd jdjd}||fS||zdk(r(||zdk(r ||z||zg} tj| StS) NrcddS)Nz9adaptive_max_pool2d(): Expected 3D or 4D tensor, but got r?r*rsr,rz%adaptive_max_pool2d.. KD6QRSr+r"rc*djddS)Nzjadaptive_max_pool2d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes  with dimension being emptyrrrsr,rz%adaptive_max_pool2d.. %'',{{m3CA3lTr+r^rSr)r^rSTrr)rr?rrTr2rry new_emptyrrargmaxr@rUrr"NotImplemented) rrh_inw_inh_outw_out output_shaperrrrrs ` @@r,adaptive_max_pool2dr6 sy ::D LL S1d^  JJqMA     ;;r?D ;;r?DLE5 zUaZEKK,-> OOL ) OOL O <   zeqjEx>,,u}}R0b9CCBGQQRTUw e|qTE\Q.u}dem4 ++E;?? r+ct j tj dv fdtd D]. tjj dkD fd0j d}j d}j d}|\}}}|dk(s |dk(s|dk(rPt j dd|||gz}j|j|tj fS|dk(ry|dk(rt|dk(rotjd d } tjjdd jdjdjd} | | fS||zdk(r4||zdk(r,||zdk(r$||z||z||zg} tj| StS)NrHrcddS)Nz9adaptive_max_pool3d(): Expected 4D or 5D tensor, but got r?r*rsr,rz%adaptive_max_pool3d..8 r(r+r"rc*djddS)Nzjadaptive_max_pool3d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes r*r+rr,sr,rz%adaptive_max_pool3d..= r-r+rr^rSr)rr^rSTrr)rr?rrTr2rryr.rrr/r@rUrr%r0)rrd_inr1r2d_outr3r4r5rrrrrs` @@r,adaptive_max_pool3dr=/ s ::D LL S1d^  JJqMA     ;;r?D ;;r?D ;;r?D%E5% zUaZ5A:EKK,-u0EE OOL ) OOL O <   zeqjUaZE|TB LLr* 3 Yr] Yr] Yr]  w e|qTE\Q.4%<13Du}demTU]C ++E;?? r+cDtjddt|D]$\}}tj|dk\||ffd &t |j d t |z}td|Dr|j|Sttj|j d }ttj|}dg|jz} |j d | d |tj||jj!| |zzj#d} |j|} tj%| j#d| g|j#dd j!| j S) N max_unpooling d_forward_outrcdd|d|S)Nr?zSd(): output_size must contain non-negative spatial dimensions, but got output_size[z]=r*)rr2rOs r,rz_max_unpoolnd..t s%u%;;<#RvGr+c3&K|] }|dk( ywrr*)rrs r,rz _max_unpoolnd..z s (a16 (rr"rrSFr()rDalert_not_deterministicrHr?rryranyrqroperatorrrrrrrr*r+) r{rrrOrr2r5nchwindices_nc_shape indices_flatrs ` r, _max_unpoolndrJg su !!M#m"DE[) 4 AIT    5SD)*T+->>L (< ((~~l++  djj3$/ 0B  k *BsTYY"jj3$/Usd$++b+5::;KLrQQ gbk^^L )F  ! !r\NDLL,< "  d6<<r+ctjjtjk(fdtjt dk(fdtjj dvfdtjj j k(fdtdj D].tjjdkDfd 0tdS) Nc"djS)Nz2elements in indices should be type int64 but got: r)rsr,rzmax_unpool2d.. sDW]]OTr+r#c"dtdS)NzMThere should be exactly two elements (height, width) in output_size, but got elements.rrsr,rzmax_unpool2d.. ;'(  4r+rc$djdS)NzLInput to max_unpooling2d should be a 3d or 4d Tensor, but got a tensor with dimensions.rrsr,rzmax_unpool2d.. s%%)YYK| =r+c<djdjSNzBExpected shape of indices to be same as that of the input tensor (z%) but got indices tensor with shape: r)rr{sr,rzmax_unpool2d.. s,PQUQ[Q[P\]229-- Br+r"rc*djddS)NzZmax_unpooling2d(): Expected input to have non-zero size for non-batch dimensions, but got r* being empty.r)rr{sr,rzmax_unpool2d.. s%::,&6qcHr+) r?rrrrrrrTr2rJ)r{rrrs```@r, max_unpool2drX s LL $T LL KA  LL V  LL gmm# 1dii   IIaL1     w Q 77r+c tjjtjk(dtjjdvfdtjt dk(fdtjt dk(fdtjt dk(fdtjj j k(fdtd jD].tjjd kDfd 0tjd d kDxrd d kDxrd d kDfd tdS)Ncy)Nz(elements in indices should be type int64r*r*r+r,rzmax_unpool3d.. rr+r8c$djdS)NzLInput to max_unpooling3d should be a 4d or 5d Tensor, but got a tensor with rSrrsr,rzmax_unpool3d.. s^_d_i_i^jjvwr+rc"dtdS)NzVThere should be exactly three elements (depth, height, width) in output_size, but got rNrOrPsr,rzmax_unpool3d.. rQr+c"dtdS)NzRThere should be exactly three elements (depth, height, width) in stride, but got: rNrOrzsr,rzmax_unpool3d.. sdehioepdqq{|r+c"dtdS)NzSThere should be exactly three elements (depth, height, width) in padding, but got: rNrO)rsr,rzmax_unpool3d.. sefijqfress}~r+c<djdjSrUr)rrsr,rzmax_unpool3d.. s,PQVQ\Q\P]^229-- Br+r"rc*djddS)NzZmax_unpooling3d(): Expected input to have non-zero size for non-batch dimensions, but got r*rWrr,sr,rzmax_unpool3d.. s% ;;-'7s-Ir+r#cdS)Nz5strides should be greater than zero, but got stride: r*r_sr,rzmax_unpool3d.. sGxPr+) r?rrrrrrrTr2rJ)rrrrzrrs`````@r, max_unpool3drd sD LL $&X LL fw LL KA  LL F q| LL G ~ LL w}}$ 1ejj !  JJqMA     LLq A 9&)a-9F1IMP a 88r+)rmrc$t||||d|S)NTinplacerm _index_addr9rOrrrms r, index_add_rk s aeVT GGr+c$t||||d|S)NFrfrhrjs r,rara s aeVU% HHr+rgc tjjtjj j k(xrj j k(fdtjjdkfdjdk(rj dnd jdkDrj nd tj k( fddk7retjj tj tuxstjt  fdzjdk(}|rjdn}dzfz}|rtjntj} | ||d } |rS|r| j!dS| j#S) NcVdjdjdjS)NzVindex_add(): self, index and source expected to be in the same device, but got (self) , (index) , and (source) rrrr9sr,rz_index_add.. 4hhZz%,,@"MM? ,r+r"c$djdSNz(Index should have dimension 1 or 0 (got rCrrsr,rz_index_add.. :5::,aHr+rcdddS)NzNumber of indices (z') should be equal to tensor.size(dim) (z ), for dim=r*)rO index_sizersr,rz_index_add..% s!%j\1XYdXeeqmplrsr+c(dtddS)Nzalpha argument of type z cannot be safely cast to type !)r)rm python_typesr,rz_index_add.., s-d5k]:YZeYffghr+r5Tr()rDcanonicalize_dimsrr?rrr2 dtype_to_typerris_weakly_lesser_typerrUr index_put_ index_putrr)r9rOrrrgrmzero_dimrYrrrrxr{rs```` ` @@@r,riri s  ! !!&&# .C LL ELL >QXX%>  LL aH#(**/AqJ&,kkAo&++c"1K LLz!s z))!''2  4  E**4; D h %vv{H#QB C-5( "C#*I BV 5C!)s{{1~?s~~/??r+ctjt|dkDdtjdk(xsdk(fdt|}|dj}|dd}t d|D}|r||f}n||f}||z}|dj ||} dt|z} t |D]} || j| j } || jdz } dk(rtj| | d| fz|}ntj| | | dfz|}|rtj| |d| } tj| |d| } | S) Nrcy)Nz#received an empty list of sequencesr*r*r+r,rzpad_sequence..> rr+rrcddS)Nz:Expected padding_side to be one of left or right, but got r#r*) padding_sidesr,rzpad_sequence..A sL\NZ[\r+r"c3>K|]}|jdywrrF)rr9s r,rzpad_sequence..F s/!&&)/s)rrrrOr) r?rrr2rrrTr8rrrr) sequences batch_first padding_valuersequences_sizemax_size trailing_dimsmax_lenout_dimsr dim_paddingsrcurrseq pad_amountrows ` r, pad_sequencer; s LLY!#%RS LL9,'"9\^N|  "HQRLM/Y//G"G,^,-'H A,  - 8CC ..L > "@A,// /2w||A. 7 "&&J7C&&Q7C %%c3AQ%?C%%c3AQ%?C@ Jr+c"t||||dS)NTrg _index_copyr9rOrrs r, index_copy_r` s q#ufd ;;r+c"t||||dS)NFrrrs r, index_copyre s q#ufe <.s rrr+r"c$djdSrtrrusr,rz_index_copy..{ rvr+cddS)NzIindex_copy_(): When source is scalar, index should have one element (got rCr*) num_indicessr,rz_index_copy.. s$$/=3r+c>djdjdS)Nztindex_copy_(): When source and destination are not scalars, their dimensionality must match. Source dimensionality (z), destination dimensionality (rCr)rr9sr,rz_index_copy.. s)KK= ?xqJr+r*cddddS)Nz_index_copy_(): Source/destination tensor must have same slice shapes. Destination slice shape: z at dimension z and source slice shape: z at dimension 0.r*)rOtensor_sliced_shapex_sliced_shapesr,rz_index_copy.. s. ((6'7~cUK'':&;;K Mr+c4ddjdS)Nz"index_copy_(): Number of indices (z') should be equal to source.size(dim) (rCrF)rOrrsr,rz_index_copy.. s,4[MB..4kk#.>-?qBr+r5)rwrrDr|rr?rrrrr2rUrrrrr)r9rOrrrgrrrYrrrrrrs```` @@@r,rrk s =  ! !!&&# .C LL ELL >QXX%>  LL aH++-K {{a 1    1 KK166 !  <=66A:QWWTc]QWWS1WY%772N8> a TcV\\#')44R LL~23 {{a 6;;s+ +  vv{H#QB"'**/EOOA uE C-5( "C#*I BV $C!)s{{1~?s~~/??r+c*tj|jd|}tjtj| }|j s |j r|jd}n|}|tj|z |fS)Nr*rM)r?rrqrhrris_xpurL)r{rrjrs r,log_sigmoid_forwardr sn --r*D 1C 599T?"#A ||t{{% Q  ''r+lowhighrc tj|j|jt |t ||j |j |S)N)rzrrrrr)prims_uniform_helperrrzr rr)r9rrrs r,uniformr sF   xxz cN t_ggxx r+c<|jt||||Sr5)rOr)r{rrrs r,uniform_r s ::gdCy9 ::r+c"t|dz }|>tj|dudtjt||k(d|S|tj|dudtjt||k(dg}t|D]Z\}}t ||k(r$|j ||dzt |z8|j t ||dz|z\|Stjddy)Nr#cyNz9Must specify exactly one of output_size and scale_factorsr*r*r+r,rz.upsample_compute_output_size.. rr+cyNr*r*r+r,rz.upsample_compute_output_size.. rr+cyrr*r*r+r,rz.upsample_compute_output_size.. rr+cyrr*r*r+r,rz.upsample_compute_output_size.. rr+Fcyrr*r*r+r,rz.upsample_compute_output_size.. rr+)rr?rrHrrr)r9r scale_factorsspatial_dimensionsrrs r,upsample_compute_output_sizer s Z1, T ! O   S%);;ZH  4  O   S'+==zJ m, CDAq1v{"":a!e#4s1v#=>""7:a!e+>r+cg}t|}|rdnd}t|D]}||}|j| |z} ||||dkDr | | ||zz n| |z } tj|tj |j } | |z| zjtj} t|dz |z D]} | jd} |j| |S)Nrrrrr"rS) rrTrr?rrrr8rrUr)rrrrrnum_spatial_dimsrrrisizernoutput_indices input_indicesrVs r,!_compute_upsample_nearest_indicesr$sG;'SsF # $&A --12 ay$Q UVAY& '  e5==V(61U:>>u{{K '!+a/0 8A)33B7M 8}%7&8 Nr+)preserve_memory_formatrrct|||gSr5rrrrs r,upsample_nearest1drJs UK& ::r+c"t|||gdSrrrs r,upsample_nearest_exact1drVs UK& FFr+scales_hscales_wc t||||gSr5rrrrrs r,upsample_nearest2drds UK(H1E FFr+c$t||||gdSrrrs r,_upsample_nearest_exact2drqs UK(H1ET RRr+scales_dc"t|||||gSr5rrrrrrs r,upsample_nearest3drs UK(Hh1O PPr+c&t|||||gdSrrrs r,_upsample_nearest_exact3drs!  {Xx:$ r+rc`t||||}ddg|z}tj||}|jdvrst j |}|j d}|jdk(r.|jjdk(r|dkrtj}|j|}|S)Nrr8r"rHcudar\) rr _unsafe_indexrrDrrrrr?r`r) rrrrspatial_indicesrrr] n_channelss r,rrs8 {F%OTl_,G   w /F {{f33E: [[^ ;;!  1 1V ; Q!33M"""? Mr+c |r|rd}n |rd}n|rd}nd}t||zdk7rtdt|d|tdt||Dcgc]}t||||zc}Scc}w)NrrHrr#rz len(params)=z is not divisible by group_size=)rr0rTr)params has_biaseshas_projections group_sizers r, gather_paramsrso     6{Z1$3v;-'G | T  49CK3T ./fQZ()  sA2c~|r'|d|z|d|z}}|d|zdz|d|zdz}}n||||}}d\}}||||fS)Nr#r"NNr*)rhiddensr bidirectional cur_params cur_hidden bidir_params bidir_hiddens r,params_hiddensrsk!'AAJ %+AEAI%6A 8Jl !'GAJJ %/" l z< ==r+c||krtd|d|d|j|jd|||z |jdd|S)Nlast_batch_size (z) must be > batch_size (rCr)r0rr3)rlast_batch_size batch_sizers r,update_hidden_for_packedrsc*$00H TU V   NN:$$Q Oj4PQR   Q: ..r+c ||k(r|S||k\rtd|d|dtj||jd|||z fS)Nrz) must be < batch_size (rCr)r0r?concatr3)rrr inp_hiddens r, update_hidden_for_packed_reversersn*$*$00H TU V   <<    a*2N O  r+c X|d}|d}|r|dnd} |r|dnd} g} g} |r|dn|d} |jdd| }tj|t|}|r|ddd}|D]V}|jd}| |k(rn|rt || ||}nt || || }||||| || }|} | j|X|r| jn!| j|| jtj| d}|stj| dn|}||fS)Nrr"r#rrS) r3r?rryrrrrreverser4)inphiddenrr hidden_fn batch_sizesr ih_weight hh_weightih_biashh_bias step_outputrrr split_inprr hidden_outs r,one_layer_rnn_datarsRq Iq I%fQi4G%fQi4GK"$G)0k"ok!nOq!_5J Ck!23IddO ' IIaL a   9OQJ2OQJsJ 7IwW :&#'&z" ))K #C.57A&:J  ?r+cfd}|S)NcDtj||||zSr5rlinearrrr r r r  nonlinearitys r,rKzrnn_cell..inner"s AHHZGDqHIIr+r*rrKs` r,rnn_cellr!sJ Lr+cfd}|S)Ncrtj|||}tj||||zSr5rrs r,rKzrnn_cell_data..inner)s2 HHQ 7 +AHHZGDqHIIr+r*rs` r, rnn_cell_datar(sJ Lr+c x|d}|d}|r|dnd}|r|dnd} tj|||} |r| jdn| } |jd} g} | D] } || | |||| } | j | "|r| j t j| d}|| jdfS)Nrr"r#r) rrfliprUrrr?r4r)rrrrrrr r r r precomputed_inputrrrrs r, one_layer_rnnr 0sq Iq I%fQi4G%fQi4Gi95<)..q1BS!!!$JK 'q*i)WU :&' ))K #C  ""1% %%r+c|d}|d}|r |d}|d}nFtj|j}tj|j}|djd} |djd} g} d} | jd} d}d}d}d}|j }| j } | j } tj j jj|||||| | || | | |||||}|d|d|d}}}||jd|jdffS)Nrr"r#rF) r?rr2rUrr6rmkldnn_rnn_layerrr)rrrrrw0w1w2w3hxcxr mode hidden_size num_layersrrrnoutputsrZhycys r,mkldnn_one_layer_lstmr/FsS B B AY AY [[ # [[ #   Q B   Q BK D''!*KJMK E .. C B Biinn--55         !G$ GAJ 2rA rzz!}bjjm, ,,r+c |r|jddn|}g} t|D]} t||| |\} } }}|r || dz kr|nd}| || | |\}}| j||r!| ||||d\}}| j||r*t j |g|j dz }n|}|dk7s|s| |dz kst j||d}|r|jddn|}|| fS)Nrr"rT)r)rn) transposerTrrr?r4rOrq)rrrrr+rqrnrrlayer_fn final_hiddensrrrrrfwd_inp fwd_hiddenbwd_inp bwd_hiddens r, _rnn_helperr8xs"&1EOOAq !eEM : >=K FA}> : J l$ QU(:'&uj*jQZ( "*|\:t# GZ   , IIw0'++-!2CDEE a,&1EOOAq !eE - r+c |jd} t||d}t|| |||||||ttt t j \} } | t j| dfSNrFr) unbindrr8rr rr?rstack rr'rrr+rqrnrrrrr3s r,rnn_tanh_inputr?tYYq\F 6:u 5F$   %**)=> C  M1- --r+c |jd} t||d}t|| |||||||ttt t j \} } | t j| dfSr:) r<rr8rr rr?rr=r>s r,rnn_relu_inputrBr@r+c |jd} t||d}t|| ||||||dtt|t t j \} } | t j| dfSNrFr r) r<rr8rrrr?rr= datar r'rrr+rqrnrrrr3s r, rnn_relu_datarH{YYq\F 6:u 5F$    ##EJJ/ C  M1- --r+c |jd} t||d}t|| ||||||dtt|t t j \} } | t j| dfSrD) r<rr8rrrr?rr=rFs r, rnn_tanh_datarKrIr+cltj||||z}|jd|}|dj} |dj} |dj } |dj} | |z| | zz} | | j z}||ntj||d}|| fS)NrHrr"r#rrrchunkrr)rr'r(r r  hr_weight chunk_dimgates chunked_gatesin_gate forget_gate cell_gateout_gater.r-s r, lstm_cellrW%s HHRG ,s 2EKK9-MA&&(G"**,Ka %%'IQ'')H r Wy0 1B BGGI B ahhr9d&CB r6Mr+c (|d}|d}|r|dnd}|r|dnd}t|dk(r|dnt|dk(r|dnd} |djd} |djd} tj|||} |r| j dn| } g} | D](}t || | ||| d\} } | j | *|r| jtj| d}|| jd| jdffS)Nrr"r#rrrHrP) rrUrrrrWrrr?r4r)rrrrrr r r r rOr'r(rrrs r,one_layer_lstmrZ3s*q Iq I%fQi4G%fQi4G[A%q F q8H6!9d   Q B   Q Bi95<)..q1BSK 3B 7IQRSB2 ))K #C A 1 . ..r+c |d}|d}|r|dnd}|r|dnd} t|dk(r|dnt|dk(r|dnd} g} g} |r|dn|d} tj|t|}|r|ddd}|d}|d}|j dd| |j dd| }}|D]}|j d}t j|||}|| kra| j|j d|| |z |j d|| |z f|j dd||j dd|}}|| kDrXtj||j d| || z fd}tj||j d| || z fd}t||||| | d\}}|} | j||r| j||f}nZ| j||f| jt| \}}tj|dtj|df}tj| d}||fS) Nrr"r#rrrHrSrY)rr?rryr3rrrrrrWrrr4)rrrrr rr r r r rOrrrrorig_hxorig_cxr'r(rrhidden0hidden1rs r,one_layer_lstm_datar`Nsyq Iq I%fQi4G%fQi4G[A%q F q8H6!9dKG)0k"ok!nO Ck!23IddO QiGQiGq!_-q!_- B  IIaLhhsIw/   NNIIaOa$78IIaOa$78  YYq!Q'1a);B  W^^AO8KLMqBW^^AO8KLMqB3B 7IQRSB236"X Bx =YYw*EIIgq,AA ))K #C  ?r+c4d}||||rtStS)a*Check whether we could use decompose lstm with mkldnn_rnn_layer. All the below conditions need to be met: * ``torch._C._get_mkldnn_enabled()`` returns ``True``. * All the input args are on CPU. * The dtypes of args are either torch.float or torch.bfloat16. * Inference. * ``has_projections`` returns ``False``. Args: * input: the input sequence to LSTM * hx: a tuple of the input hidden state and cell state ``(h_0, c_0)`` to LSTM * params: the weight and bias tensors of LSTM cNtjjsy|gt|ztt j |z}|Dchc]}|j }}t|dk7ry|j}|tj dk7ry|Dchc]}|j}}|D]&}|tjtjfvs&y|jry|djd|djdk7} | ryycc}wcc}w)NFr"rrr#T)r?rs_get_mkldnn_enabledryr from_iterablerrpoprr bfloat16 requires_gradr2) rr'rrtdevicesrdtypesrrs r, use_mkldnnz2select_one_layer_lstm_function..use_mkldnnsxx++-'DH$tE,?,?,G'HH%,-188-- w<1  U\\%( (#*+a!''++ EU[[%..99    Q%**Q-2a5::a=8 ).,s D#D")r/rZ)rr'rrks r,select_one_layer_lstm_functionrls!:%V$$$r+c t|dk7rtdt|t|||djd|djdk7}t t |d|d} t |||} t|| |||||||| \} } t t | } | tj| ddtj| ddfS)Nr#$lstm expects two hidden states, got rr") rr0rr2ryrrlr8r?r=) rr'rrr+rqrnrrrr2rr3s r, lstm_implros 2w!|CCG9MNN 6:r!uzz!}1 1 /M NF #beRU# $F-eR@H$   Cm,-M  M!,a0%++mA>NPQ2R RRr+c t|dk7rtdt|t|||djd|djdk7}t t |d|d} t || ||||||dtt| \} } t t | } | tj| ddtj| ddfS)Nr#rnrr"F)r ) rr0rr2ryrr8rr`r?r=rFs r,lstm_data_implrqs 2w!|CCG9MNN 6:r!uzz!}1 1 /M NF #beRU# $F$   #= Cm,-M  M!,a0%++mA>NPQ2R RRr+c&|jdd}tj|||jdd}|d|dzj}|d|dzj} |d|d|zzj } || z | z| zS)Nrr"r#r)rNrrrr rrr r r r chunked_igateschunked_hgates reset_gate input_gatenew_gates r,gru_cellrysYYq!_NXXj)W=CCAqIN #nQ&77@@BJ #nQ&77@@BJq!^A%6%CDJJLH  !Z /( ::r+cPtj|||jdd}tj|||jdd}|d|dzj}|d|dzj} |d|d|zzj } || z | z| zS)Nrr"rr#rMrss r, gru_cell_datar{sXXc9g6<> 1 1 umWg r+ct|j||}t|d}t|d}tjj j |||||Sr)rr2rr?r6r_upsample_bicubic2d_aars r,upsample_bicubic2d_aa_vecr_V ){M REmQ/GmQ/G 99>> 0 0 umWg r+ct|j||}t|d}t|d}tjj j |||||Sr)rr2rr?r6r_upsample_lanczos2d_aars r,upsample_lanczos2d_aa_vecrkrr+czt|j||}|r|ndgt|z}t||||Sr5)rr2r_upsample_linear)rrrrrrs r,_upsample_linear_vecrws= ){M RE+]$#e*1DF E5- @@r+rc t||||gSr5r)rrrrs r,upsample_linear1drs E; z JJr+c"t|||||gSr5r)rrrrrs r,upsample_bilinear2drs E; (?S TTr+c$t||||||gSr5r)rrrrrrs r,upsample_trilinear3drs!  {MHh+I r+cL|r|dkDr |dz |dz z SdS| |dkDrd|z S||z S)Nr"rgrr*)rrrrns r,_compute_scalersB5=\# (S.1HqH#/EAIsU{U7XCUUr+c&|r||zS||dzzdz SNrr*)rn dst_indexrs r,_compute_source_indexrs$y   C(3..r+weightsweights_precisionctdt||Dd|dz zz}||z }tj|ddj tj S)Nc3K|]F\}}|jtj|jtjzHywr5)r8r?r>)rrrs r,rz%_sum_tensors_uint8..s826!QU[[ADD--sA Ar"r) _sum_tensorsrr?rr8ry)rrrrs r,_sum_tensors_uint8rsd:=c7:K  1$ %'F( (F ;;vq# & ) )%++ 66r+ctj|j}d}tj||j}d|d|dzzzz}|dk\}||j z S)Nrrr"i)r?r=rrrr)r max_weightmax_weight_precision precisionsrrs r,_compute_weight_precisionrsiW%))+J2:;L;LMJ :zA~!67 7F g D $((* ,,r+c Ljd}jdd}t|}tjtjj \}fd}t t|||D cgc]\} \} } }|| | ||dz | z } } } } }tt| \} }}g}tddgg|zD]c}ddgt|Dcgc]}||dk(r||n||c}z}tj|}t|}|j|ett|D]p} | | || z j!ddj#}t|ddd|dddDcgc]!\}}|t%j&||z |z#}}}rt|dk7rt)dt||d}tj*}j,j.d k(r|d krt$j0}t3|t$j4s!t)d t/|j6|j9| }j;s|j=}|Scc}} } } wcc}wcc}}w) Nr"r#r|c|t|| |}tj| jj }t || j d}|j|jdgdg|z}|j tj}|dzj |dz }|||fS)Nrrrrrr"r) rr?rrr8rrr*rr) inp_sizerrnsqueeze scale_factorrx_f32r9xp1rrrs r, get_valuesz$_upsample_linear..get_valuess%h-P  LL%,, 7 : : : G%lA}EKKPSKT ekk!n@sh/?@ HHU[[ !1umm1 m-a}r+rrrgz+Expected vs to have exactly 1 element, got rz$Expected result to be a Tensor, got r\)rrrDrEr INT_TO_FLOATrHrryr rTrrrrr_rr8r?rr0rrrr`r7rr$rruround)rrrrr inp_sizesn_dimsrVrrrrrxs_f32xsxp1svsrNkrvxscalev1v2rr]rs` ` @r,rrsQJ ABI ^F'' !AANNHAu 09  ; /0  +A+(F 8XvvzA~>F CL)FB B 1vh' (TluV}U!qtqybed1g=UU   uc * #Au - !  eFm $ )be#**3477>b1gr!$Q$x0 B 27F+ +    2w!|J3r7)TUU UF//6M ||F"zB// fell +24<3H3H2I J    ]  ;F  " " $ MYV sJ J &J rNrOc4|j|jk(Sr5r)rNrOs r, is_same_sizers 77agg r+c.tj||Sr5)rr)r9rrFs r,_reshape_aliasr$s 99Q r+c.tj||Sr5)rr)r9rs r,rr*s ::a !!r+c2tj||||Sr5)rr)r9rr|r)s r,r+r+/s >>!WeZ 88r+c|D]F}|tj|jtjtjfvdHtj|jtj k(dddlm}||jdk(r.:rr+cyNz*tensors used as masks must be bool tensorsr*r*r+r,rz&_unsafe_masked_index..?rr+rrr"r)r?rrrrrrwrrrmeta_index_TensorrrrTrrr2rrr)r9rrfillrr meta_resultrs r,rr4s   LL  EII66M  LL ejj < Eaggi1n%//AA!WM zz+++T22 3w< ?   q A >GAJ?   a ) 5 5teT BBr+cL|D]F}|tj|jtjtjfvdHtj|jtj k(d|j dk(r|jStt|D]B}||}| |j|j| |j|dz ||<D|j|d}tj|||dS)Ncyrr*r*r+r,rz5_unsafe_masked_index_put_accumulate..Vrr+cyrr*r*r+r,rz5_unsafe_masked_index_put_accumulate..[rr+rr"rTr()r?rrrrrrrrTrrr2rrr+)r9rrrrr masked_values r,rrPs   LL  EII66M  LL ejj <  wwyA~wwy 3w< H   !&&)QGGAJH %%teQ/L  ! !!Wlt ! LLr+c|j}d}|dkrd}|6|dkDr*dg|z}|jd||<|j|}n|}||z}tj||k7|d} | j |} tj ||| j| } tj||k7| d} |tjjk(r|dkDr|jdd} | | fS|lj|j}tj ||| j|} tj||k7| d} | j} n"||k7jj|} |tjjk(r| j} | | fS|tj jk(r| j| z } | | fS)Nr"r#rr*r)rOrrr?rirUgatherrr!r'r|rrrr8r)r()r{rrrr&rr+rwr, safe_target_rr'wsums r,_nll_loss_forwardrjsXXZFK z   A:E"(aE+  E"AAax++f 4fa@K((5Lll4l;CCKP PF [[</ ;FINN(((VaZ}}R- |##  HHTZZ ||A{L9AA+N{{6\14;xxz ,.33588> IMM''' <  inn** * , < r+cpjdkDrjdkstdjdjdkDrtdjdjdk(xrjdk(}|s7tjjdjdk(fdjd}|A|jdk(r|j |k(std |d |jt |||S) Nrr#r>r?r"r@c>djdjdS)NrArBrCr)r{rsr,rz"nll_loss_forward..s 0 Jv||nTUVr+rSz/weight tensor should be defined either for all z7 classes or no classes but got weight tensor of shape: )rOr0r?rrrr)r{rrrr&rE n_classess`` r,nll_loss_forwardrs# HHJNtxxzQDTXXZLPQRSS zz|aOPVPZPZP\~]^ _  88:?8vzz|q'8L  JJqMV\\!_ , V  2I 6::<1#499T=i[I//5||n >  T669l KKr+c t|||||Sr5)r)r{rrrr&s r,nll_loss2d_forwardrs T669l KKr+Ac0|dz|z|dzz |z|zdzS)Nr#rr"r*r9rs r,_upsample_cubic_convolution1rs( UaK1q5 !Q & *Q ..r+c<||zd|zz |zd|zz|zd|zz S)NrrHr*rs r,_upsample_cubic_convolution2rs0 UQU]a !a% '1 ,q1u 44r+rhcd}|jtjdk(rtj|d|z gd}tj|dzd|z gd}t||}t ||}tj |d\}}tj |d\}} ||| |fSt|dz|t ||t d|z |td|z |fS)Ngrrgrrr)rr?r=rrr<) rhrtt1tt2w03w12r#r&r$r%s r, _upsample_get_cubic_coefficientsrs Axx5<<&&kk1cAg,A.kk1s7C!G,!4*32*32cq)Bcq)B2r2~ )S! 4 (A . (q! 4 (q! 4   r+coeffstscPt|}tdt||DS)Nc3,K|] \}}||zywr5r*rrrs r,rz+_upsample_cubic_interp1d..sEHRREr)rrr)rrcoeffs2s r,_upsample_cubic_interp1drs$.r2G EFG0DE EEr+c6ttj|Sr5)rr?r7)rs r,rrs %))R  r+ num_stepsc|dkrtjd||S|s|dz |z nd}tj| ||||S)Nr"rr)stepsrr)r?rlinspace)rrrrrNs r,_linspace_from_neg_onersIA~||AfE::-:)a-9 $A >>1"ayu MMr+thetahrc|j}|j}t||||jd|d}t||||j|dd}t j d||}tj jj|ddd}tj jj|ddd}tj jj|d dd}||z|zS) Nr")r"r"r"r)rr#constantrrr)r|r"r")r#r rrrrr?rrrJr) rrrrrrgrid_xgrid_ygrid_ones r,_make_base_grid_4drs KKE \\F$A}eV D I I!QPQ RF #A}eV D I I!QPQ RFzz)5@HXX $ $VjPQ $ RF XX $ $VjPQ $ RFxx""&&xV*TU&VH F?X %%r+rc|j}|j}t||||jdd|d}t||||jd|dd}t||||j|ddd} t j d||} tj jj|ddd}tj jj|ddd}tj jj| d dd} tj jj| d dd} ||z| z| zS) Nr")r"r"r"r"r)rrrrr)r"r#)r#r")rrr) rrrrrrrrrgrid_zrs r,_make_base_grid_5drs5 KKE \\F #A}eV D I I!QPQST UF #A}eV D I I!QPQST UF #A}eV D I I!QPQST UFzz,eFCHXX $ $VjPQ $ RF XX $ $VjPQ $ RF XX $ $VjPQ $ RFxx""&&xV*TU&VH F?V #h ..r+c|\}}}}t||||}|jddd|jjdzj d}|j|||dS)NrrSrr"r^r#)rrrbrUr) rr2rrWrVrr base_gridgrids r,_affine_grid_generator_4dr "sgJAq!Q"5!QmLI NN2q! $uxx'9'9!'< < A A" ED 99Q1a  r+c|\}}}}}t|||||}|jddd|jjdzj d} | j||||dS)Nr rSrHr"r^r)rrrbrUr) rr2rrWrVrrrr r s r,_affine_grid_generator_5dr,smMAq!Q"5!QOI NN2q! $uxx'9'9!'< < A A" ED 99Q1a ##r+ctjt|dvdt|dk(rt|||St |||S)Nr8cy)NzCaffine_grid_generator needs 4d (spatial) or 5d (volumetric) inputs.r*r*r+r,rz'affine_grid_generator..<rr+rHr )r?rrr r)rr2rs r,affine_grid_generatorr6sJ LL D VU 4yA~(MRR(MRRr+r interpolation_mode padding_mode _expand_gridc  !"#$%&'()*+,-./012tjdvfdtjdvfddtdtdtffd 2dtdtd tdtfd 0dtdtdtf0fd $dtdtdtf$2fd }j\!()|j\}./}|d k7rt d|ddlm} m} tjtjj} jjdvxr| | (| k)| k} | rtjntj+r(|j!d./|j!!./d }dtdtdtf()fd *tj"!jj!ddd"tj"jjddd dtdtdtdt$f!*+./fd #dtdtdtf "#fd &|d} |d}dk(r|| )}||(}|j'|j'c,-,dz-}},-dz}}||}}||z ||z z}||z ||z z}||z ||z z}|,z |-z z}t)&fd,-|f|||f|||f|||ffDSdk(r<|| )}||(}|j+}|j+}&||dS2| )}2|(}|j',|j'-|,z 1|-z }s"1j-d1|j-d}dtdtdtf$&()fd 'd tdtf',-1fd! %t/%fd"t1d#D}t3||S)$N)rr"r#cdS)NzInvalid interpolation mode r*)rsr,rz"_grid_sampler_2d..Us-.@-ABr+cdS)NzInvalid padding mode r*)rsr,rz"_grid_sampler_2d..Xs-B<.+Qr+coordsr2rPcBr|dzdz n|dz}|dzdz }||z|zSrr*)rr2rofsrs r, unnormalizez%_grid_sampler_2d..unnormalize[s8%2tczCs Sj3|c!!r+ twice_low twice_highcP||k(rtj|S|dz }||z dz }||z j}tj||}||z j j tj }tj|dzdk(||z||z|z S)Nr#rr"r)r?r rfmodfloorr8int8ri)rrr coords_min coords_spancoords2extraflipss r,reflect_coordinatesz-_grid_sampler_2d..reflect_coordinatesfs  "##F+ +] !I-2 J&++- 7K0;&--/222D{{ AINEJ. j0H50P  r+cdk(r|Sdk(rtj|d|dz Sr|dd|dz z}n|dd|zdz }tj|d|dz S)Nrr"r#rSr)rr2coords_reflectedrrr's r,compute_coordinatesz-_grid_sampler_2d..compute_coordinatesrsx 1 M Q ;;vq$(3 3#6vq!tax.#Q #6vr1t8a<#P ;;/D1H= =r+c(||}||Sr5r*)rr2 coords_unr*rs r,compute_source_indexz._grid_sampler_2d..compute_source_index~s- "9d33r+r#z4grid last dimension must be 2 (for x,y coords), got r)rsym_and)rxpur"rysc tjd|ktj|ktjd|k|kSrsr?r)rr0iHiWs r,in_bounds_condz(_grid_sampler_2d..in_bounds_condsF  GU&&rBw0A0A!r'2PR70ST  r+rwsc||rndt fd|j |j |fDS)Nr"c3pK|]-}tj|dj/ywr)r?rir)rrhrrroHoWs r,rz1_grid_sampler_2d..clip..s7  KKa # ( (Ar2 6 s36r)rr8) rr0r6rrrrrr5 index_dtyper9r:s @@r,clipz_grid_sampler_2d..clipsQb"% A1 ee+e.K0H"M   r+ixiyc8 |||\}}}||f|zSr5r*) r=r>ridx_xidx_yw_C_idxN_idxrNr<s r, get_summandz%_grid_sampler_2d..get_summands0B?ubue+,r11r+).r).r"c3:K|]\}}}|||ywr5r*)rr=r>rrEs r,rz#_grid_sampler_2d..s( R B " sc<|}|}||dSr\r*)r=r>r9rZr*rEr3r4s r,get_value_boundedz+_grid_sampler_2d..get_value_boundeds*#B+A#B+Aq!Q' 'r+rc|dz z}dz ||dz|dz|f}t|S)Nr"r#)r)riy_ofscsrHix_nwiy_nwtxs r, get_coeffz#_grid_sampler_2d..get_coeffs[cAg&F!%!)V4!%0!%!)V4!%!)V4 B ,B3 3r+c3.K|] }|ywr5r*)rrrOs r,rz#_grid_sampler_2d..s:#y~:rH)r?rrrrr0rwrr.rr>rrrrrrrrr rrrUrrTr)3rNr rrrrr-rVtworr. int32_maxuse_32bit_indicesr9rZr=r>ix_neiy_neix_swiy_swix_seiy_sew_nww_new_sww_se ix_nearest iy_nearesttyrrrCrrDr<r*rOrErHr3r4r5r;rLrMr9r:r'rNrs3` ```` @@@@@@@@@@@@@@@@@@@@r,_grid_sampler_2drbDs" LLi'B LL !#Q"F"#"&"  F  s    PV   >F ># >& >4V4346477LAq"bZZNAr2s axB3% H  N EKK(,,I 8^iy1>"3%++ K yyAr2s+221aRC 6 v & LL188 , 1 1!Q1 =E LL188 , 1 1!Q1 =E    V    4F   22F2&22 V A V AQ !!R ( !!R (xxz288: uqy%ueaiueu urz*U urz* rEz*U rEz* t$t$t$t$      q !!R ( !!R (XXZ XXZ :z155 B  B    %Z %ZaBaB (& (f ( ( (  43 46 4 4:q::'33r+c"t|||||S)N)r rrr)rb)rNr rrrs r,grid_sampler_2drds   -!#  r+c&tjjdk(xrjdk(fdtjjdjdk(fdzj dS)Nr#r"cLdjdjS)Nzmatrix @ vector expected, got rrr{rsr,rzmv..s!0 BswwykJr+rcvdjddjddjddS)Nzsize mismatch, got input (rr9r"z), vec (rCrFrgsr,rzmv..s<,TYYq\N!DIIaL>RURZRZ[\R]Q^^_`r+r)r?rrOr2rrgs``r,rr ss LL  a*CGGINJ LL !  #` 3J    ""r+c|-|dz |zdz}d|z |z|tj|zz }nd|z |ztj|z }|||z}|j|j}t ||Sr\)r logsigmoidr8rr)r{rr pos_weightr log_weightrs r, binary_cross_entropy_with_logitsrms  1n.2 F d"j1<<3E&EFF d"Q\\$%77 f} 776<< D i 00r+tensor1tensor2is_outc  |j|jk\r||fn||f\}}ddlm |jdk\r|jdksy|jr|sy|jdk(ry t |dk(ry|j }|j }dg}t|ddD]}|j||dzt fd t|tt||DS) Nrrrr#FTr"rSc3VK|] \}}}|dk(xs ||k("ywrr*)rrrr2rs r,rzshould_fold..Es9 D% tqy!B^DEM%BBs&)) rrwrrg sym_numelrrzr_rr rry) rnrorpt1t2t1_shape t1_strideexpected_strider2rs @r, should_foldry-s$+<<7<<#?gw gwEWFBD GGqLRWW\ ||qimq()xxH IcO!"&;tob&99:; !$ tH_56"  r+) pass_is_out)rpc ddlm}m}|j}|j}|dk(s|dk(rt d|d||dk(r|dk(rt j ||S|dk(r|dk(rt j||S|dk(rC|dk(r>t jt jt j|d|dS|dk(r|dk(rt j||St|||r>||kD}|r |jn|}|s|n|dk(r|jn|} |j} t| dd} t!t"j$| } | jdk(} | r| j'| jd|j)| | d}| rWt j*j,j/|j| | }|r|jj1S|St j*j,j/|j| | S|dk\rp|dk\rj|dkDr|j3dnd}|j3d}|jdd}|dkDr|j3dn|j3d}|dkDr|j3dnd}g}t5|dz D]"}|j'|j3|$|d k(r|d k(r{||d|dk7rj||ddk(r'|j6rt9|jd|S||ddk(r'|j6rt9||jdStt j:||}|||gz}t=|}|j?|j)|||}|dk(}|r7||gz}|j?|j)||jd}n)|||gz}|j?|j)|||}|} |dkDr| j'||dkDr| j'||r/|jA|jdjC| S|jA|jC| St jDd d y) Nr)r guard_or_truez9matmul does not support 0-dimensional tensors, got dims: z and r"r#rSr^rFcy)Nz/both arguments to matmul need to be at least 1Dr*r*r+r,rzmatmul..rr+)#rwrr|rOr0r?dotrrrrUryrbrhrryrrErrr*r6r _unsafe_viewrr2rTrgrabroadcast_shapesr*rbmmrr)rnrorprr| dim_tensor1 dim_tensor2r1rtrusizes_1r5 folded_dim1 t2_is_matrix t1_foldedrrWm1 batch_tensor1m2rV batch_tensor2rexpand_batch_portiontensor1_expand_sizeexpand_batch_producttensor1_expanded vector_rhstensor2_expand_sizetensor2_expandeds r,raraMsT++-K++-Ka;!+G }TYZeYf g  aK1,yy'**  kQ.xx))  kQ.}}UXXeoogq&A7KQOO  kQ.xx)) Wgv . +- $WZZ'$G+:J799;PW ((GCRL) X\\<8 vvx1}     ,JJ{GBK8 YY^^00b1A<PF-6699'') BF B99>>..y||B/?N N  kQ.!,aGLL Q \\"  cr* !,qW\\" gll26F +aGLL Q#% {Q' 2A  a 1 2 1 q mA.-2BBCmA.!349N9Ngooa0'::mA.!349N9Ngwq'9:: $  " "=- @ 3aW<#$89#>>*=>FF !R !A% "6""= 23-r21  #7"a"@ &~~.ABJJ$b!  , ?    " ?    " #''(89AA"EJJ<X X#''(89>>|L L UUVr+rrcnj\}}t|d||}t|d||}tjtjj \}}t j|djj|} t j|djj|} t|| |} t|| |} | jd} | j} | j}| |z jdd}| | z jdd}| jt j} |jt j}|dz ||dz|d zf}| dz | | dz| d zft|t|}d \}j t j"k(rt%t%|}Dcgc]@}|dzzt j&|d zzjt j(Bc}|Dcgc]@}|d|zzt j&|d zzjt j(B}}fd fd t+fd|D}j t j"k(r| t-dt/|||}nt1dt3||D}tj4}|j7|}|Scc}wcc}w)Nrr"r|rrrSrrgr#rrctj|ddz }tj|ddz }tjdd||g}|Sr)r?rrr)r0ry_idxx_idxrin_hin_wrs r, load_boundedz0upsample_bicubic2d_default..load_bounded sO B4!8, B4!8,   utT5%&@ Ar+ctfdD}jtjk(r t dt |St dt|DS)Nc30K|] }|ywr5r*)rx_ofsrrZs r,rzCupsample_bicubic2d_default..get_x_interp..sBl1e,Bsz4weights_precision_x must not be None for uint8 inputc3,K|] \}}||zywr5r*rs r,rzCupsample_bicubic2d_default..get_x_interp..sJRBGJr)rrr?ryr0rrr)rZsrc_xrixs_ofsrweights_precision_x weights_xs` r, get_x_interpz0upsample_bicubic2d_default..get_x_interpscB'BB ;;%++ %"*$J&eY8KL LJCy4IJJJr+c3.K|] }|ywr5r*)ry_ofsrs r,rz-upsample_bicubic2d_default..s;%,u%;rQz4weights_precision_y must not be None for uint8 inputc3,K|] \}}||zywr5r*rs r,rz-upsample_bicubic2d_default..sL(2rb2gLrr\)rrrDrErrr?rrr8rrUr rrrrryrrint16rr0rrrrr)rrrrrrVh_scale_factorw_scale_factorrrrx_floaty_floatr9rZyscaleriys_ofs weights_yweights_precision_yrsrc_yrr]rrrrrrrs` @@@@@@@r,upsample_bicubic2d_defaultrs{{Aq$$D+a.-QN#D+a.-QN'' 5#H#H#U#UHAu  [^ELL9<<5<IA [^ELL9<<5<IA#NA}EG#NA}EG#G A Ak c *Fk c *F U[[A U[[A1uaQA&G1uaQA&G08I08I/9,, {{ekk!7 B7 B !** +ejjmc.A A E Eekk R  !** +ejjmc.A A E Eekk R   KK ;7; ;E {{ekk!  & !WX X#E96IJLc%6KLL//6M   ]  ;F MI  s9AL-AL2c 6tjt|t|zdk(d|Q| tdt t t t ft dt|jdd|D}|r|nd\}}t|||||S)Nr"cy)Nz:Must specify exactly one of output_size and scale_factors.r*r*r+r,rz(upsample_bicubic2d_vec..4rr+z7scale_factors must not be None when output_size is Nonec3PK|]\}}tt||z ywr5)rr )rrrns r,rz)upsample_bicubic2d_vec..=s*Au ! u,-s$&r#r) r?rrr0r rrrrr)rNrrrrrs r,upsample_bicubic2d_vecr's LL [D//14L   I  #s(O  #AGGABK ?  )6}<GW %amWg VVr+c(fd}t||S)Nctj| ||zj}|dz |dz |jz jz S)Nrr")r?rrrrmiddlerdim_idxrNs r,rz_reflection_pad..idxLsF,,ufunQXXFzVaZ'++-7<<>>>r+_reflection_or_replication_padrNrrs` r,_reflection_padrFs ? *   r+c(fd}t||S)Nctj| ||zj}tj|d|dz S)Nrrr")r?rrrrs r,rz_replication_pad..idx]s6,,ufunQXXF{{7Avz22r+rrs` r,_replication_padrWs 3 *   r+idx_fnc t|dz tj|j dz dzfv fd|j d}|j z }t Dcgc]}|d dz |z z}}t Dcgc]}|d dz |z zdz}}|}t D]E}dg|jz} |||||||| ||z<t j|| }G|jjdvrtj} ntj|} |j| }|Scc}wcc}w)Nr#r"c(dddzddzdS)Nreflection_padz d requires r"rr#zD inputr*rsr,rz0_reflection_or_replication_pad..ps$.[q sQwiwOr+)rmpsr/r\)rr?rrOrrTrrrrr`rDrr) rNrr inp_shapenc_dimr padding_left padding_rightrrr]rOs @r,rrhsd g,! C LL C!GS1W%%OI UUWs]F8=c C1GAq1-.CLC=B3ZHWQ#'A+.23HMH F 3Z1&**,. a)A, a@PQAJ##FC01 xx}}..// 33A6   ]  ;F M1DHs 1EEc t|dz|j dDcgc]}|dz  }}tDcgc]}|ddz |z z}}tDcgc]}|ddz |z zdz}}g}t|jD]c}dg|jz} d| |<|j t j |j||jj| e|d | d} dtDcgc] }| |||z} }tDcgc] }||| |z } }tDcgc]}d||z||z| |z } }tDcgc]}| |d||||z||zf}}tjtjtDcgc] }||c}}tj|| zd}fd}tjtDcgc]}gd c}D]}|t!dgzk(rg}g}tD]t}||dk(r | |}||}n=||dk(r| |}| |d||f}n$||dk(r| |}| |||||z ||dz f}|j |j v||||}|Scc}wcc}wcc}wcc}wcc}wcc}wcc}wcc}wcc}w) Nr#r"rSrcF|\}}}tj||k\||kSr5r2) index_rangerlbubs r,index_range_conditionz7_reflection_pad_backward..index_range_conditions( 2r  b!r'22r+rrc .t D]*}||d||dk}t|ts%|s(|cStjt j |Dcgc] } | c}}t j ||zd}||zScc}w)Nr#r"r)rTr7rrLrrrr) rr index_rangesrupper_less_than_lowerrrgrOrOrlrs r,r)z,_reflection_pad_backward..accumulatess A$0OA$6a9K$K !/6;P      CO PK "; / P   % %k4S# Fax QsB )rSrr")rrrTrrr?rrrrLrrrr itertoolsr r)rlr9rrdhwrrrrrZxyzcenter left_reflect right_reflectrange_crrr)rVareaoutsrrrrOrOrs` @@@r,_reflection_pad_backwardrs` g,! C''3$%. )Q1q5 )C )8=c C1GAq1-.CLC=B3ZHWQ#'A+.23HMHG 166]SS166\  1 u||AGGAJqxx@EEjQRS #A 3$%.C316c ;1c!f|A& ;F ;6;CjALOc!f,ALADI#JOqQQZ,q/1CF:OMO NSSVZHIAs1v Q/-2BBCG    eCjQ0<Q D  $ $[$F C HD !!c #C1J#CD4 5!s# #  s -AAw!|Qi%aj aB"1o"1vq,q/: aA#A&"1vs1v a0@'@#a&1*M KK     , -$l3-40 Kc *CH2<AOR&$Ds5 KK$$K)K.8K3K8K=L & Lrrrcjtj|||}tj|||}||fS)Nr)r?aminr)r{rOrrrs r,aminmaxrs2 ::dW 5D ::dW 5D :r+rctjtjtj|d||||S)Nrr)rrr?riisnan)r{rOrrs r,nansumrs2 88EKK D 11d;S'QV8 WWr+rrrrrrrc Ntjjd|d||||S)Nrr"rrr start_step)rnrrrrrs r,arange_defaultrs/ ;; ! ! 3vf " r+c Ntjj||d||||S)Nr"rr)rmrnrrrrrs r, arange_startr s/ ;; ! ! sAU6&Z " r+cddlm}||i|S)Nr)out_dtype_dense)!torch._higher_order_ops.out_dtyper)rFrGrs r,out_dtype_decomprsA D +F ++r+marginct tjtjjd jd tj|dk(xs|dk(dtjj dk(xr dk7fdtjj dk(xrj k( fdQtjtjj dk(xrj k( fdjdtjd}||z z}|jd}|dk(r|n||z}|z}tj j }tj|k7|d}|tjjk(r|j!S|tj"jk(r |j%|jdz S|j!d S) Nrr"r#cy)Nz only p == 1 and p == 2 supportedr*r*r+r,rz#multi_margin_loss../rr+c"djSNzMExpected non-empty vector or matrix with optional 0-dim batch size, but got: rr\sr,rz#multi_margin_loss..2s_`e`k`k_lmr+c(ddjS)Nz#inconsistent target size, expected r"r)nframersr,rz#multi_margin_loss..6s5fXYv||nUr+c(ddjS)Nz#inconsistent weight size, expected r"r)rOrsr,rz#multi_margin_loss..<s9#i ~Vr+rrr)r? atleast_2d atleast_1drrrrrUrrcrrrir!r(r|rr)r) rrrVrrrurjrrOrs `` ` @@r,multi_margin_lossr s   U #E   f %F [[^F ++a.C LLa!16#MN LL a$C1Hm LL q5V\\^v5U!!&) KK1  63!6 V   a F U0A UA AA !VQA  v  ,,s5<< 0C C6M1a(AINN(((vvx imm)) )uuw##vv!v}r+ is_targetc |j |jtj|}tj|}|jd}tjt dkxr|dk7 fdtjt dkxr k( fdtj ||j }|dk(}tjtj|||dd }||k}tj||d}tj|d| } tj||d} tj|| jd k(d } d | jjd z |z} | jd} | |z } tj| d| } |tjj k(r!| j#d j%} n@|tj&j k(r| j#} n| j#d } | j)|j*j-} | | fS)Nr"r#rcdSrr*)orig_input_shapesr,rz0multilabel_margin_loss_forward..^s_`p_qrr+cddS)Nzinconsistent target size: z for input of size: r*)rorig_target_shapesr,rz0multilabel_margin_loss_forward..bs,->,??STdSefr+rrSTrrrrg)rrS)rr?rrrrrrrirrDrUTrcr!r(r|rrr)r8rr*)rrrrOris_endend_idx target_masktidx0rtidx1rrjrrs @@r,multilabel_margin_loss_forwardrOs{{    U #E   f %F ++a.C LL "/saxr LL !#M(9=M(Mf ,,s6== 1C r\FjjVS#6BMG-K KK VQ /E U%0A KK VR 0E #R!88a@I accmmm##e+A AA CA Iq!$AINN((( EEgE  # # % imm)) ) EEG EEgE  U[[)112CDI i<r+) attn_maskrnquerykey dropout_p is_causalrctjtjfdtjjdk(xr(jdk(xrjdk(fdtjdk(fdtjjdjdk(xrjdjdk(dt j j||d|jdjdk7 \}}|jd d ddjtj jdd d d}||fS) Nc"djS)Nz-query must be FP32, FP64, BF16, FP16 but got r)r sr,rz.s? }Mr+rHcndjdjdjS)Nz,q, k, v must be a 4 dimensional tensor, got rr)r r r|sr,rz.s3>uyy{m2cggi[XZ[`[d[d[fZghr+rcdS)Nz&dropout probability must be zero, got r*)r sr,rz.s$J9+"Vr+rcy)Nz&q, k, v should have the same head sizer*r*r+r,rz.rr+r")rr r  dropout_maskrn enable_gqar#rr\) r?rrurOrr"_scaled_dot_product_attention_mathrr2rrr`) r r r|r r rrnrattns ```` r,*scaled_dot_product_flash_attention_for_cpursV LL &M LL q@SWWY!^@ q0@h LLSV LL A%++a.(KSYYq\U[[^-K8 ::BB   ::a=CHHQK/C LFDF q!Q" %"9"9 : Aq!  4<r+c.t|fd}|S)Nc<|i|}|dj|Srs)rO)rFrGr outplace_ops r, inplace_opz$register_inplace..inplace_ops%4*6*Aw}}S!!r+r)aten_oprrs ` r,register_inplacers"G$"%" r+cB|js&|jst|}t|}tj||}t |t jr|dk7r||z}|dk(r|St |t jr|dk7r||z}||zSrg)rurrr?rr7numbersNumber)r{batch1batch2rcrmrs r,baddbmmr"s  ! ! #DOO,=4yE  YYvv &F eW^^ , % qy dGNN +tqyd{ &=r+c2tj||dS)Nr rr)r{rs r, floor_divider$s 99T5 88r+c`tjtj|jdSr\)rLrrErr)rhs r,rsrss   HLL!''1 55r+rrc|"tjj|g|Stjj|g||S)Nrr&)rr dim_IntList IntList_out)r{rrs r, sum_defaultr*sC {xx##D"E#::xx##D"Es#CCr+c t|tjs|S|@tjj |t t|jStjj ||gSr5) r7r?rrrdimsryrTrO)r{rOs r,squeeze_defaultr-s\ dELL )  {||  tE$((*,='>??||  u--r+c2tfdtt|jD}|jt j k(rt jnd}|jd|d|}|||j|jz z|fS)Nc3.K|] }|k7s |ywr5r*)rrrOs r,rz)_weight_norm_interface..s@1qCxQ@s r#T)rr) rrTrrrr?rfr rr8)rrrOkeep_dim norm_dtypers ` r,_weight_norm_interfacer2st@c!''l 3@@H !5>> 9tJ 66!Xt:6 >D DGGAGG$$ %t ++r+ assume_uniqueinvertct|tjs!tj||j}t|tjs.|rtj ||Stj ||Sddlm}||jdt|jdzkrt|||St||||S)Nrrrg$@g(\?r5r3) r7r?rrrner rwrrr isin_default isin_sorting)elements test_elementsr4r5rs r,isinr=s h -&&x 8L8LM mU\\ 2 88Hm4 488Hm4 4Dm))+dS9I55Q.QQRHmFCC m=  r+)rcB|?tj|jtj|j}n?tj|j|tj|j}||kj |j }|S)Nr)rrr)r?randr2rrr8r)r{rraw_prVs r, bernoullirA1sq  499;emmDKKP IIK--;;   $**%A Hr+r7c|jdk(r%tj|tjS|jdk(r ||k(}|r|S|S|j d|jzz}|j |}ttd|j dz d}||k(j|}|r|S|S)Nrrr"rSr"r) rr? empty_likerrrrrrTrD)r;r<r5rxexpanded_elem_shaper9rOs r,r9r9Ds~~1 ;;Q-'t&3&"..4-2D2D+DD )*A b=---126 7C  " "s " +CC4"s"r+c|j}|j}|rtj||g}tj|d\}}|dd|ddk(} tj| ddgd} |r| j } tj | } | jd|| } | d|jj|jStj|\} } tj| |} tj| | jk| d}| ||k(}|r|j n|}|j|jS)NT)stabler"rSrF) r@r?r4sortr logical_notrDrrr*r searchsortedri)r;r<r4r5 elements_flattest_elements_flat all_elementssorted_elements sorted_orderduplicate_maskrsorted_test_elementsrVrtest_idxcmps r,r:r:QsN$$&M&..0yy-1C!DE (- <(M%(,0DD..~1vuM +779N/q,?A()11(..AA"'**-?"@a  !5}E;;s%9%?%?%AA3J"8, =#)coos{{8>>**r+c.|jd}||SrR)r*)r{r flatteneds r,takerVns R I U r+c|tj}|tjk(r t|}tj ||j |Sr6)r?r`preserve_formatrrresizer)r{rr]s r, resize_asrZusD// ----e4 ;;tU[[ ; FFr+c|tjrtS|jrtS|jj dk(rtS|j d}|j d} |j d} |j d} |jdk(} | s3|jd}|jd}|jd}|j d} |j d}|jtjtjfv}|rtjn |j}tj| |z|| z||j}|j| |z| | z}|j| |z| | z}|r|jtj}|j!d||}|j| ||| }|r|j|j}t#j$|}|j'|}| s|j)d}|S) u Decomposition of max_pool2d_with_indices_backward using scatter_add. This replaces the native implementation with a high-level decomposition that uses scatter_add for gradient accumulation. The scatter-based approach provides automatic optimization opportunities for Inductor and handles all pooling configurations without requiring specialized fallback paths. Algorithm: For each output gradient position, use the corresponding index from the forward pass to scatter the gradient to the input position. When multiple output positions select the same input position as max, scatter_add automatically accumulates their gradients. Complexity: O(B * C * H_out * W_out) Independent of kernel size, unlike traditional O(B * C * H_in * W_in * K²) approaches that iterate over input positions and kernel windows. Known Limitations: - FP16/BF16: Uses FP32 accumulation internally to preserve precision when many gradients accumulate to the same position (overlapping pooling windows). This adds slight overhead but ensures numerical stability. - Deterministic mode: Falls back to native implementation to ensure consistent results across runs Args: grad_output: Gradient w.r.t. pooling output [B, C, H_out, W_out] self: Original input tensor (for shape) [B, C, H_in, W_in] kernel_size: Pooling kernel size stride: Pooling stride padding: Pooling padding dilation: Pooling dilation ceil_mode: Whether to use ceil for output size calculation indices: Indices from forward pass (per-channel linear positions) Returns: Gradient w.r.t. input [B, C, H_in, W_in] rr^rSrHrr"rr\)r?$are_deterministic_algorithms_enabledr0rrrr2rOrUrrrfrrr*r8 scatter_addrDrrr)rlr{rrzrrrr in_heightin_width out_height out_widthrrchannelsuse_fp32_accum accum_dtypegrad_input_flatgrad_output_flatrIrXr]s r, max_pool2d_with_indices_backwardrg~sd 113%' " Iyy}H!!"%J  $IqJ ~~a !++A. ##A&1Jyy|H !&&5==%..*IIN#1%--{7H7HKkkXH!! O#**XzI5??:#8*y:PQL+..u}}=&11!\CSTO!((Xy(SJ]];#4#45 //5M&&]&CJ ''* r+ window_lengthcLtjj|d||||S)ahann_window(window_length, *, dtype=None, layout=None, device=None, pin_memory=False) -> Tensor Returns a Hann window of size :attr:`window_length` with ``periodic=True``. Equivalent to :func:`torch.hann_window` with ``periodic=True``. Args: window_length (int): the size of returned window. Keyword args: dtype (:class:`torch.dtype`, optional): desired dtype. Default: global default. layout (:class:`torch.layout`, optional): desired layout. Default: ``torch.strided``. device (:class:`torch.device`, optional): desired device. Default: current device. pin_memory (bool, optional): if ``True``, pins the returned tensor. Default: ``False``. Tr)r hann_windowperiodic)rhrrrrrs r,rjrjs42    $ $  % r+rkc||ntj}|dk(rtjd||||S|dk(rtjd||||St j |}|r|dzn|}tj |||||}|dtjz|dz z z}tj|}|dzdz}|r|jdd|n|} | j|S) a@hann_window(window_length, periodic=True, *, dtype=None, layout=None, device=None, pin_memory=False) -> Tensor Returns a Hann window of size :attr:`window_length`. .. math:: w[n] = 0.5 - 0.5 \cos\!\left(\frac{2\pi n}{N-1}\right) where :math:`N` is ``window_length + 1`` when ``periodic=True`` (for spectral analysis), or ``window_length`` when ``periodic=False`` (symmetric window). Low-precision dtypes (``bfloat16``, ``float16``) are computed in ``float32`` then cast. Args: window_length (int): the size of returned window. periodic (bool, optional): if ``True``, returns a periodic window for use with STFT. Default: ``True``. Keyword args: dtype (:class:`torch.dtype`, optional): desired dtype. Default: global default. layout (:class:`torch.layout`, optional): desired layout. Default: ``torch.strided``. device (:class:`torch.device`, optional): desired device. Default: current device. pin_memory (bool, optional): if ``True``, pins the returned tensor. Default: ``False``. rrMrr"rCrrr) r?get_default_dtyperxrrDr rpicosr3r8) rhrkrrrrr compute_dtyperWrhwindows r,hann_window_periodicrrsD&EE,C,C,EE{{ fV   zz fV   //6M% =A    A S588^q1u %&A ! A D3A.6QXXaM *AF 99U r+ num_classesc|dk(r*t|jjdz}tjj t j|dk\dtjj t j||kd|jdt j||j|jk(jt jS)NrSr"rz+one_hot: Class values must be non-negative.z7one_hot: Class values must be smaller than num_classes.r)rrrr _assert_asyncmsgr?r rUrrrr8r)r{rss r,one_hotrwrsb$((*//+,q0  $!)5  $$%A r << 4::dkk J Kbor+)FF)noner5)r#)rNNr")rSFFrrMr)r"r"F)Frr)rrgN)rr"N)Fr)NNN)rrFT)rrFr^)rFr)rS(rLrrrEr{collections.abcrr contextlibrenumrrrr r typingr r r?torch._meta_registrations torch._primsrrtorch._prims_common _prims_commonrDtorch.nn.functionalrrJrr rr torch._decomprrrrrrrrtorch._prims_common.wrappersrrrr torch.utilsrrBtorch.utils._pytreerrs DispatchKeyrrystr__annotations___opsr6rr!rrrNrcompute_only_pw_cast_for_opmathpw_cast_for_opmath"pw_cast_for_opmath_non_tensor_argsrpw_cast_for_int_to_realrrWr`rbrkr rwrr8r}rrrrrrrrrrrrrrrrrrrrrrr(r|r _safe_softmaxrrrrrXr r#rr%r.r<rFrIrNrQrSrUrXrkrsslicerrrrrrrrrrrr4r8rgrmrqpy_implCompositeImplicitAutogradAutogradrprrrrrr*rrrrrrrrrrrrr dtype_outrrrrrrrr0r2rBrFr[rare unsafe_chunkrhrdrcno_statsrlrpr|rrr_fused_dropoutrrrr]rlift lift_freshrrrrrrrr_adaptive_avg_pool2drrr"r%r6r=rJrXrdrkrarirrrrrr Generatorrrrrrrrr_upsample_nearest_exact1drrrrrrrrrrrrrr r/r8rnn_tanhrr?rnn_relurBrGrHrKrWrZr`rllstmrorqryr{grur}rrrrrrrrrrrrrrrrrrrrr+rrrrrrrrrrrrrr rrrbrdrrmryraupsample_bicubic2drrreflection_pad1dreflection_pad2dreflection_pad3drreplication_pad1dreplication_pad2dreplication_pad3drrreflection_pad1d_backwardreflection_pad2d_backwardreflection_pad3d_backwardrrrrrrrrrmrrrr+_scaled_dot_product_flash_attention_for_cpurrr"r$rsrr*rrOr-r2r=rAr9r:rVrZrgrjrk periodic_outrraddbmm_addbmmaddmm_addmv_baddbmm_fill_gelu_r hardswish_ hardtanh_hardtanh hardsigmoid___iand____and__ __ilshift__ __lshift__rr index_reduce_ index_reduce__ior____or__ __irshift__ __rshift____ixor____xor__ldexp_ leaky_relu_ leaky_relulogit_logitrelu_rrenorm_renormround_rscatter_r) scatter_add_r]scatter_reduce_scatter_reducesilu_rwr*r+r,rsS ."%$ #,,07 *(hh"" c zz~~   %$) ''99''" 'T#*88@@# uDDLL&-88@@ &" "uDDQQ   c f **+ \2F2v2,2--. \4v4&4/4../ \SS6SS5S0S ))* \           + 2))*+(,())*+"V","(() @f@@*@112 \fF3../ \PP%P05P@EP0P' GFGvG(G//0 F&V1//0 \:F:&:U:1:001 \LL%L7<LNRL2L **+ \))f)3),)<**+5v5f5,5 " &v&&&#&**+ \>v>V>>,> **+66666,6334%% % % 66> %5%667                8 $112 \>3G3G1 1 1-01 1'1../ \11 &1061CF101**+0,0++, ^^)) 1  1  1 1  1- 144<<=  % /5 BE MR > 44??@ R R  R  R R  R  RA R00889  % /5 BE NS : 00445 R R  R  R R  R  R6 R#$#$ #$ #$ TM #$  #$  #$#$ #$L))* \NfNFNNN+N(../ \.. . . TM .  .  .. .0.b001 \$$ $ $ TM $  $  $$ $2$N112 !^^)) 1 1 1 TM1 1  131&::; \ !^^))       TM     <  --. ^^))1 1 11 1/1667 \ ^^))         8  " ))v)%)#),,- &&F&v&.&++,  O Oc O  O  O O  O- O ))* 3? 3? 3? : 3? t 3?  3?+3?l  #d 14t 38_.**+  $U $U $U $U : $U t $U  $U,$UN,,- EEd3iEcERUE.E ../ OO&*3iO9<ODGOORO0O%+:?++334 \  X X!' X.1 X@E  X!5 X778  KK!'K.1K@E K!9K*& $ M McM3iM#Y M I M  M%M` $ ] ]c]c]3i ] #Y ] I ] ]%]@445     v  e  6  ,,- c# c#"3ic#47c#?Bc#JMc# c#.c#L++334;?  % ,1DL   5 & %kCCDk2236eD4K4E&++, VV A& AU A4$; A- A & S'*,6,,,())* FT+,' $        ((556     7B DI  &\  &\ *T&\ tF| # " &\" ""J00$//2E2EFG   &\     $   H   '')C)C)G)GH #   c    f    &\D   ))00155C5c5%PS BT52544<<=67B B $S B03B 63;B>B  ))*//C/c/% :L/+/2--55)) // // &// // 63; ////f++TZZ^^<= fF##>"))4::+?+?@A   +  +  +  +{{ +  +  +B +../            0" f F  s s   # 55f56555$577??@S&S& S& S&  S& D= S&  S& S& S& S&dS& 6D=&4-$ 67S&AS&n77;;<     D=      d ,, ,, ,, 6D=&4-$ 67=8 6D= FTM 77??@QQ Q3iQ  Q  Q TM Q 4-QdQ 6D=&4-$ 67QAQj77;;< 3i     TM  4-d ,, ,, ,, 6D=&4-$ 67=4,,4453 33i3 TM3  3  66> 363l55==>HH H3iH  H TM H d H 6D=&4- '(H?HVR R TMR 4-R4- R $ R  RR RR 6666D=&4- ?@Rj../ UK/ (  ( TM ( 4- (4- ( $ (  ( (  ( 666 !" (00 (4'' (<(<='' (M(MN    TM   4-  4-  $          666 !"  O>  F"";#H#HIG4<GJGAAIIJ  TM 4-       666 !"K*55==> (  ( TM ( 4- ( (  (  ( (  ( 666 !" (? ( 55>>? (  ( TM ( 4- ( (  (  ( 666 !" (@ (@@HHIK K TMK 4-K K  K  KK K 66666 12KJK4  TM 4-      :44<<=1 1 TM1 4-1 1  1  1 1 6666 )*1>14??GGHC C TMC 4-C C  C  C C 666666 9:CIC822::;1 1 TM1 4-1 1  1  1 1 6666 )*1<14++, VV (- ( & !% "&04A A ;; A LL4  A  AA&&-A'ANdiiAB C &&{';';<--. VVVV,   4-4-  $   !&-/=> 00889  TM4-  $  } $  d 66D=&4- /0:677??@ii i TMi4- i $ i } i$i i idi 66D=&4- /0iAiZ77;;<"" " TM"4- " $ " } "$" " "d" ,," ,," ,," 66D=&4- /0"="J778 VVV$  4-  $  } tm%90667 VVV$  4-  $  } tm%82112 d'vd'E#s(Od'3d'NC C cC  IC #Y C 3i C  C  C  66>C L445 UI#   c  I #Y  3i    66> 6 445 UI#   c  I #Y  3i    66> 6 001 UI, ,sCx, 66>,2,^001 UI3 3sC}%3 66>323l   ) 8(!2( % ""(,   e  *u % &" &;'; 8 //334//334//334$$[%J%JK$$[%9%9:$$[%J%JK$$[%9%9:$$[%J%JK$$[%9%9: 3  3cT! 3;% 3 3;L;L;L555 366::;66::;66::;##++K,Q,QR##++K,@,@A##++K,Q,QR##++K,@,@A##++K,Q,QR##++K,@,@A ?  ?cT! ?;% ? ?BSBSBS<<< ?#L0088$:Q:Q:U:UVW  (()N)NO  (()=)=> Dd; ; ;c; DL; ;<?PX; ##++T-K-K-O-OP''// 0U0UV''// 0D0DE Dd; G GcG DLG G<FW G0088$:Q:Q:U:UVW  (()N)NO  (()=)=> Dd;"! G GcGdlGdl G  G<?PXG ##++T-K-K-O-OP''// 0U0UV''// 0D0DE Dd;"! S ScSdlSdl S  S<FW S0088$:Q:Q:U:UVW  (()N)NO  (()=)=> Dd;"!! Q QcQdlQdl Q dl Q  Q<?PXQ ##++T-K-K-O-OP''// 0U0UV''// 0D0DE Dd;"!!   c dl dl  dl   <FW    c      4&>/$FK+\&,/-d& R ++,[BBC[112.3D-.8 ++,[BBC[112.3D-.8 **+KAABK001.2C,.@ **+KAABK001.2C,.@ /6=@/d (>>?--.S/@)S@ ' ==> ,,-S.?(S>;; &{<<={++,.->'.6' ==> ,,-..?(.644889!!))+*O*OP!!))+*>*>?@Q:33778  (()N)NO  (()=)=>?P933778  (()N)NO  (()=)=>?P90044511556##K$I$IJ##K$8$89%%k&K&KL%%k&:&:;&&{'L'LM&&{';';<A=N<M:K76A //779O9O9S9STU  " K KcKKdl K  KVK %%t'?'?'C'CD!!))+*>*>?  "! U UcUUdl U dl U  U@ U &&(A(A(E(EF  "!!   c  dl  dl  dl    V/7 & 7$,V$47IO7 7-'9-f-I IcII   I  IIZ))112Fv$3,,d.?.?@A B++,-"."//019292234C5C6AABCMDM23 3 3  TM3  3  3  66> 3 l--. X~&L L L TML L  L  66> L'/L@//0 X~&L L L TML L  L  66> L'1L/F/u//5F5u55  3E (F%7FVFF !Xf%!&!NN#'N05 NEJ\\N&f&&&T&"/f////T/"!V!49!T!$V$49$T$223 SStCySS4S  p4 p4 p4p4 p4  p4  p4 p4f,,-             .     #! #==> $)..:N:N1?1$ dt@[BBC>>? ',yW@DyWx0088$:Q:Q:U:UVW  (()=)=>  ! T TsCxTTT\ T T\ T  T?XTn//334$$[%J%JK$$[%9%9:  15 W WsCx4'WW&- W  W;L5 W4--.--.--.  v c3h F /// ../../../   sCx V 000 % % 38_% c3_f, -% %P667667667 \T888Tn % UEu& $ X4X%X,,dkkoo>? !% =="&    ;;   LL  LL4    @ **+, !% =="&      ;;   LL  LL4   -  ",#, ../'' (<(<=  ^^)) ) ) ))  ) TM )  ) )>0)X;;<$$,,44[5I5IJ X{#- - -- 66> -$K=-xHHPPQ  A $A A A A A  A}A 4<A 66>ARAH'  ( ))* 9+9'6(6))488<<89!% D  D ;;  D $ D  D: D--t||/?/?@A.&.sTz.B. ==>,?, " 38 # (../)- ,, %  \\  0 $5: #vv vvv?vr))1143C3C3G3GHI !%"&"&"  ;;  LL4   LL4   t  JB))22D4D4D4Q4QRS 6!%"&"&"666 ;;  6 LL4  6 LL4  6t 6 6T6rt{{+djj)djj) -TYY'TYY'$..1/""D$4$45 -!!4??3$..1##T%6%67t{{+!!4??3 -djj)!!4??3djj)TYY't{{+djj) -""D$4$45%%t':':;TYY'kCCD&sFEr+