a h*@s,UddlZddlZddlZddlZddlZddlmZddlmZddlm Z m Z ddlm Z m Z ddl mZmZmZmZmZddlZddlZddlmZddlmZddlmmZddlmZm Z m!Z!ddl"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 l0m1Z2dd l3m4Z4ej5j6Z6gZ7e8e9e:d<ej;jdneej?e@e@dddZAe eAej?jBddZCe eAej?jBdZDe eAej?jBddZEe eAej?jFdZGe!eHe!dddZIe#e=jJe/deDe!e!dddZJe#e=jKe/deDe!e!dd d!ZKe#e=jLe/deDe!e!eMeMd"d#d$ZLe#e=jNe/deDe!eMeMeMe@e!d%d&d'ZNe#e=jOjPgd(d)ZQe#e=jOj!ge!d*d+d,ZRe#e=jSe/eDe!e!d-d.d/ZSe#e=jTe/deDe!e!d0d1d2ZTe#e=jUe/de!e!eMeMd3d4d5ZUe#e=jVe/eDe!e!d-d6d7ZVe#e=jWe/eDe!e!e!d8d9d:ZWe#e=jXe/de!e!eMd;dd?d@ZYe#e=jZe/deDdoe!e!e9dBdCdDZZe#e=j[eDe!e!dEdFdGZ[e#e=j\e/eDe!e!d-dHdIZ\e#e=j]e/deDe!e!e!d8dJdKZ]e#e=j^e!e!e!dLdMdNZ^e#e=j_e!e!e!e`e!e!fdOdPdQZ_e#e=jae/eDe!e!e!eMeMe@e@e!dRdSdTZae#e=jbe/deDe!e!e!e!dUdVdWZbe!eHdXdYdZZcejdd[d\d]Zee#e=jfe/eDe>jgjhfe!e!eHe!d^d_d`Zfe#e=jie/deDe!e!e!eHdadbdcZie#e=jjdpdddeZke#e=jle/eDe>jgjhdffe!e!eHeMdgdhdiZle#e=jmjneDe!e!e!eHeMdjdkdlZme#e=jmjoeDe!e!e!eHeMe!dmdndoZpe#e=jqjneDe!e!e!eHeMdpdqdrZqe#e=jqjreDe!e!e!eHeMe!dsdtduZse!e!e!ee!eHeHe!e!dvdwdxZte#e=jue/deDe!e!eHe!dydzd{Zue#e=jve/de!e!e!ee!eHeHe!e!dvd|d}Zve#e=jwe/de!e!e!ee!eHeHe!e!dvd~dZwe#e=jxe/eDde>jgjhfe!e!ee!eHe!dddZxe#e=jye/deDde>jgjhfe!e!e!ee!eHe!dddZye#e=jze/eDe>jgjhfe!e!eHe!dddZze#e=j{e/deDe>jgjhfe!e!e!eHe!dddZ{e#e=j|e/dqe!e!eMdddZ|e#e=j}e/e!e!e!dddZ}e#e=j~e/e!e8eHeHeHeHeHdddZ~e#e=jj!dre!eHeeHeeHeHdddZe!eHeeHeeHe`eHeHfdddZe#e=je/dse!e!eHeeHeeHeHdddZe#e=je/e!e8eHeHeHdddZe#e=je/e!e8eHeHeHeHdddZe!e!ejddddZe#e=je/deCe!e!eHejddddZe#e=je/eCe!e!eHejddddZddZe#e=je/e!e8eHe8eHe8eHe8eHe!dddZe#e=je/eDe!e8eHe8eHe8eHe8eHe8eHe!dddZe#e=je/e!e!eMdddZe#e=je/e!e8eHeHeHeHe!dddZe#e=jjneDdte!e!eeMe!dddZe#e=je=jjne6je=jjne6je!eMee@dddZe#e=je/ddÃe!eMee@dddńZe#e=je/e!eHe@dƜddȄZe#e=je/ddɍe!eHe@dƜdd˄Ze#e=je/due!e!eHe@e@e!d͜ddτZe#e=je/e!e!eHeHe@dМdd҄Ze8eHdӜddՄZe8e!eHeHe8e!d֜dd؄Ze8e!dٜddۄZe8e!eHdܜddބZe8e!eHeHdߜddZe#e=jjne=jjrgdve8e!eHeHee!e!dddZe#e=jjne=jjrgdwe!e8eHeHee8e!ee8e!dddZe#e=jj!dxe!eHeHe`e!dfdddZe#e=jjndye!e8eHeHe`e!dfdddZe#e=jj!dze!eHeHe`e!dfdddZe=jje6jd{e!e!eHe`e!dfdddZe#e=je/ddɍeDd|e!e!e!eHeHdddZe#e=je/eDd}e!e!e!eHeHe@dddZe#e=je/ddɍeDd~e!e!e!eHeHdddZe#e=jjneDe!e!e!e!ee!eHeHeHeHe8e@e`ee!ee!ee!fd ddZe#e=jjre!e!e!e!ee!eHeHeHeHe8e@ej!ej!ej!e`ee!ee!ee!fdddZee!ee!dddZe#e=jjne!e!e8eHe!e!ee!ee!e8e@e`ee!ee!ee!fd dd Ze#e=jjre!e!e8eHe!e!ee!ee!e8e@ej!ej!ej!e`ee!ee!ee!fd d d Ze!ee!ee!ee!ee!e@eMeMe@e`e!e!e!ee!ee!fd ddZe#e=je/ddde!ee!ee!ee!ee!e@eMeMe`e!e!e!fd ddZe=jjne6je=jjne6je!ee!ee!ee!ee!e@eMeMe`e!e!e!fd ddZe=jjne6jde8e!dddZe#e=jjne!ee!ee!e!e!eMeMe`e!e!e!fdddZe#e=jjne!ee!ee!e!e!e@eMeMe`e!e!e!fd ddZe#e=jje!ee!ee!e@eMeMe`e!e!e!fd d!d"Ze#e=jjne!ee!ee!e!e!e@eMeMe`e!e!e!e!e!fd d#d$Ze!ee!ee!e!e!eMe@e!d%d&d'Ze#e=jjne!ee!ee!e!e!eMeMe`e!e!e!e!fdd(d)Ze#e=jjne!ee!ee!e!e!eMeMe`e!e!e!e!e!e!fdd*d+Ze#e=jjne!ee!ee!e!e!eMeMe`e!e!e!e!fdd,d-Ze#e=je/ddÃeDdd.d/Ze#e=je/ddddddd0ee!e'feejdeeje@e@eejd1d2d3Ze#e=je=je=jge/d4d5Ze=jjne6je#e=jÃe/ddÐd6d7e!e!ee!ee!ee!e@eMeMd8d9d:ZÐd;d<Ze#e=jjne!e!ee!ee!ee!ee!ee!e@eMe8e@e!e`e!ee!ee!fd= d>d?Ze#e=jjne!e!ee!ee!ee!ee!ee!e@eMe8e@e`e!ee!ee!fd@ dAdBZe#e=jjre!e!ee!ee!ee!ee!ee!e@eMe8e@ej!ej!ej!e`e!ee!ee!fdCdDdEZe#e=jȃe/ddÐd6e!e!e!ee!ee!ee!ee!eMdFdGdHZe#e=jɃe/ddÐd6e!e!e!ee!ee!ee!ee!eMe!dI dJdKZe#e=jʃe/eDe!e`eHeHfdLdMdNZe)e)e8eHeHdOdPdQZe#e=j̓e/e)e)e8eHdRdSdTZe#e=j΃e/e)e)e8eHe8eHe8eHdUdVdWZe#e=jσddXe)eHe)e)e'dYdZd[Ze#e=jЃe/ddXe)eHe)e)e'dYd\d]ZddXe)eHe)e)e@e'd^d_d`Ze#e=jjne=jjne6jddbdcZe#e=jӃe)eHe)e)dddedfZe#e=jԃe/e)eHe)e)dddgdhZe)eHe)e)e@didjdkZe#e=jփe/dldmeDe!e`e!e!fd-dndoZe#e=j׃e/de!ee@eHeMfee@eHeMfeejdpdqdrZe#e=jكddsdtZِdudvZڐdwdxZe#e=jj݃e#e=jj݃e#e=jj݃e=jjݠe6je=jjݠe6je=jjݠe6je=jjݠe6je=jjݠe6je=jjݠe6je!ee8eHee8eMe!dydzd{Ze#e=jj݃e#e=jj݃e#e=jj݃e=jjݠe6je=jjݠe6je=jjݠe6je=jjݠe6je=jjݠe6je=jjݠe6je!ee8eHee8eMe!dyd|d}Zdd~dZe#e=jjne=jjrge=jjne6je=jjne6je/dddde!e8eHeeMe!dddZe#e=jjne=jjrge=jjne6je=jjne6je/dddde!e8eHeeMe!dddZe#e=jjne=jjrge=jjne6je=jjne6je/dddde!e8eHeeMeeMe!dddZe#e=jjne=jjrge=jjne6je=jjne6je/dddde!e8eHeeMeeMe!dddZe#e=jjne=jjrge=jjne6je=jjne6je/dddde!e8eHeeMeeMeeMe!dddZe#e=jjne=jjrge=jjne6je=jjne6je/dddde!e8eHeeMeeMeeMe!dddZeDde!e8eHe8eeMe@e!dddZddZddZddZddZdddZddZddZdddZdddZddZe#e=jje=jje6je=jje6jddZe#e=jje=jje6je=jje6jddZe#e=jje=jje6je=jje6jddZe#e=jje=jje6je=jje6jddZddZdddZdddZddZe#e=jje=jje6je=jje6jddZe#e=jje=jje6je=jje6jddZddZddZe#e=jje=jje6je=jje6jddZe#e=jje=jje6je=jje6jdd„Ze#e=jj݃e=jjݠe6je=jjݠe6jdÐdĄZe#e=jj݃e=jjݠe6je=jjݠe6jdŐdƄZ e#e=j j݃e#e=j j݃e=j jݠe6je=j jݠe6je=j jݠe6je=j jݠe6je=j jݠe6je=j jݠe6jdǐdȄZ e#e=j jne=j jrge/de!e8eHe@eeMe!dɜdʐd˄Z e#e=j jne=j jrge=j jne6je/de!e8eHe@eeMeeMe!d̜d͐d΄Z e#e=j jne=j jrge/de!e8eHe@eeMeeMeeMe!dϜdАdфZ ddҐdӄZdԐdՄZee!ee!e!e!d֜dאd؄Ze*e!dٜdڐdۄZeDe!e8eHe@e8eeMe!dܜdݐdބZe#e=jjne!e!e@dߜddZe#e=je=jge/ddZe#e=jgddZe#e=jgdddZe#e=jgddZe#e=jgddZe!e!ee!eHeHe`e!e!fdddZe#e=je/dlde!e!ee!eHeHe`e!e!fdddZe#e=je/dlde!e!ee!eHeHe`e!e!fdddZe!eMe!dddZe!eMe!dddZe!e*dddZe*e!e!dddZ ee!e!dddZ!eHe@ejdejdddZ"e!eHeHe@dddZ#e!eHeHeHe@dd d Z$e!e8eHe@d d d Z%e!e8eHe@d ddZ&e#e=j'e/eDe!e8eHe@d ddZ'de!e!eHeHe@e@e!dddZ(e#e=j)e/eDde!e!eHeHe@e!dddZ)e#e=j*e/ddɍeDddZ*e#e=j+e/dde>jgjhfddZ+ej!ej!e@e@dddZ,e=j-jne6je=j-jre6je/dddd d!d"Z-e#e=j.jne=j.jrge=j.jne6je/eDde!e`eHeHfe@eeMeeMe!d#d$d%Z/e#e=j.j݃e=j.jݠe6je=j.jݠe6je/eDde!ee`eHeHfe@ee`eMeMfe!d&d'd(Z0e#e=j1e#e=j2e#e=j3eDe/e!e`eHdfe!d)d*d+Z4e#e=j5e#e=j6e#e=j7eDe/e!e`eHdfe!d)d,d-Z8e!e`eHdfeeHeHeHge!fe!d.d/d0Z9e#e=j:e#e=j;e#e=j<e/dd1d2Z=e#e=j>e/d3d4ddd5d6d7Z>e#e=j?e/ddd[d8d9Z?e#e=j@jne=j@jrge/dejAddd:e'eejdejBeeje@d;d<d=ZCe#e=j@jDgdejAddd:e'e'eejdejBeeje@d>d?d@ZEe#e%dAdBZFe#e=jGe=jGjne6je/ddde>jgjhfe!e!e'e'ee!eHe!dCdDdEZGe#e=jHe=jHjne6je/dldFe!e!eHe`e!e!fddGdHZHe#e=jIjnddddIe!e!e!eMe@ee!eeMe`e!e!fdJdKdLZJdMdNZKe#e=jLge/ddɍeDddOdPZLe#e=jMe/dQdRZMe#e=jNdSdTZNe#e=jOjne=jOjrgdddUe!eejdee!e!dVdWdXZPe#e=jQjne=jQjRgde!eeHdYdZd[ZSe#ejDs z-type_casts..inner..type_promotion_kindcst|tr|S|SdSNr.rtor1computation_dtyper(r) increase_precNs  z0type_casts..inner..increase_preccst|tr|S|SdSr5r6r8) result_dtyper(r) decrease_precTs  z0type_casts..inner..decrease_prec) rtorchtypesZ_NumberpytreeZarg_tree_leavesutilselementwise_dtypesr)argskwargsZ flat_argsr;r=rr,r*r-r+)r2r:r<r)inner?s     ztype_casts..inner) functoolswraps)r*r+r,r-rGr(rFr) type_casts9s rJT)r+r,)r+)r+r-)r1dimreturncCs$t||D]}|d}q|SN)rangerK unsqueeze)r1rK_r(r(r)_unsqueeze_to_dimvs rR grad_inputout_gradycCs|d||SNr Z conj_physicalrTr(r(r) tanh_backward|srYcCs||d|SrWrXrTr(r(r)sigmoid_backwardsrZ)rUr1beta thresholdcCs.||}t|||k||||dSN?)expr>where)rUr1r[r\zr(r(r)softplus_backwards rb) grad_outputalphascale input_scale is_resultself_or_resultc Csf||}|}|}|r6t|dk||||||St|dk|||t||||SdSNr)r>r`r_) rcrdrerfrgrhZnegcoefZposcoefZ negiptcoefr(r(r) elu_backwards rjcCs t||Sr5)r>Z full_likeselfvaluer(r(r) fill_scalarsrnrmcs(tdkfddt|S)NrcsddS)Nz@fill only supports 0-dimension value tensor but got tensor with z dimensionsrKr(ror(r)zfill_tensor..)r>_checkrKatencopyrkr(ror) fill_tensors   rv)rlrLcCs tjtj|ddddddSNrminmaxr>clamprlr(r(r) hardsigmoidsrrcrlcCst|dk|dk@|ddS)Ngg@gUUUUUU?r>r`rr(r(r)hardsigmoid_backwards rrcrlZmin_valZmax_valcCst||k||kBd|S)Nrrrr(r(r)hardtanh_backwardsrcCs$|tjtj|ddddddSrwr~rr(r(r) hardswishsr)rcrlrLc Cs,t|dkdt|dk||dd|S)Nrrx?rrr(r(r)hardswish_backwards rrcrlr\cCst||kd|Srirrr(r(r)threshold_backwardsrrcrlnegative_slopeself_is_resultcCst|dk|||Srirrr(r(r)leaky_relu_backwardsrnone)gradrl approximatecCsd}d}d}|dkr||d}d}||}||} |||| } t| } d|} d| } d| }d| | }|dd||}| ||}|||S|}||d}ddt||}|t||d }||||SdS) Ng;f?g;f?gmBP ?tanhrgHm?r rxg)r>rerfr_)rrlrZM_SQRT2Z M_SQRT1_2Z M_2_SQRTPIZkBetaZkKappaZx_sqZx_cuberGZ tanh_innerleftrightZleft_derivativeZtanh_derivativeZinner_derivativeZright_derivativeZkAlphaZcdfZpdfr(r(r) gelu_backwards,      r)rcinputcCs:tt|}t|}||d||}|||SrW)r>rFZsoftplussigmoid)rcrZinput_tanh_softplusZ input_sigmoidoutr(r(r) mish_backwards rcCs|t|Sr5)r>rrr(r(r)silu!srcCs,ddt| }||d|d|SrW)r>r_)rcrlrr(r(r) silu_backward(sr)rlweightrLcCst|dk|||Srir)rlrr(r(r) _prelu_kernel0sr)rcrlrrLcCs4t|dk|||}t|dkd||}||fS)Nrrr)rcrlrZ input_gradZ weight_gradr(r(r)_prelu_kernel_backward5sr)rcrlnoiseloweruppertrainingrrLcCs:|r||dkr||S||d}t||||SdS)Ngư>r!)mulrtr)rcrlrrrrrrr(r(r)rrelu_with_noise_backward@s   r)rcrlbufferrLcCsN|dk}t|dd}t|dd}tt| }||||d|S)Nrr rN)r>r`r_abs)rcrlrZ in_negativeZ max_derivsignrar(r(r)log_sigmoid_backwardUs rloss reductioncCs4|tjjkrt|S|tjjkr,t|S|SdSr5)rr&rmr>meanr'sumrr(r(r)apply_loss_reductionbs     rdtypecCs4|tjkrtjS|tjkr tjS|tjkr0tjSdSr5)r>Z complex32float16Z complex64float32Z complex128float64rr(r(r) to_real_dtypeks    r)rltargetrrLcCs||d}t||S)Nr!)r)rlrrrr(r(r)mse_losszs r)rcrrrcCs,|tjjkrd|nd}||||S)N@)rr&rmnumel)rcrrrnormr(r(r)mse_loss_backwardsrcCsFtj|||d}|td}tj||dd}t|}t|||S)N)rKrz-infTrKkeepdim)r>Zsoftmaxeqfloatall zeros_liker`)rlrKrrZmaskedZ masked_rowszerosr(r(r) safe_softmaxs  rr^)rlrrr[cCs<||}t||kd|d||d|}t||S)Nrr!)rr>r`r)rlrrr[rr(r(r)smooth_l1_losss &r)rcrlrrr[c CsZ|tjjkrd|nd}||}t|}||}t||k||||t|Sr])rr&rmrr>rr`r) rcrlrrr[rr1Zabs_xZ norm_gradr(r(r)smooth_l1_loss_backwards   r)rcrlrrr[rScCs*t|||||}t||jt||ddSNTZ copy_fromZcopy_to exact_dtype)rrshaper)rcrlrrr[rSresultr(r(r)smooth_l1_loss_backward_outs  r)rcrlrrdeltac Cs`|tjjkrd|nd}||}t|| k| ||t||k||||||Sr])rr&rmrr>r`)rcrlrrrrr1r(r(r)huber_loss_backwards  r)rcrlrrrrScCs*t|||||}t||jt||ddSr)rrrr)rcrlrrrrSrr(r(r)huber_loss_backward_outs  r)rcrlrrr ignore_index total_weightrLc Cs|dkrdnd}|tjjkr(||}||}t||k|d}t|} t| ||d} | |krzdkrnn ||}|durddt |D} |j d| |<| | }||}t||k|d}| |S)Nr!rr gcSsg|]}dqSr r(r0rQr(r(r)r3rrz&_nll_loss_backward..) rKrr&rmrPr>r`rscatterrOrreshape) rcrlrrrrr channel_dim safe_targetrSZ new_shaper(r(r)_nll_loss_backwards       r)rcrlrKrLc Cs|dksJdt||}||}|ddksNJd|d||d}||d|}||||}t|}d||||} ||}tj|| g|dS)Nrz*glu does not support 0-dimensional tensorsr!z.Halving dimension must be even, but dimension z is size r^rp)rKrAcanonicalize_dimsizenarrowr>rcat) rcrlrKZwrap_dimZnInZ inputSizeZ firstHalfZ secondHalfZgradInputFirstHalfZgradInputSecondHalfr(r(r) glu_backwards  rcCsxd|krdks"nJd|dks6Jd|dkoL|dk}|s|jd|jdksJd|jd|jd|dksJd |jd |d f|dus||jd ksJd |tjjkr8|dkr8|dkr |jd|jdksdJd|jdd|d|jdn,|dkrT|dksdJd|jt|||||||S)Nrr!input tensor should be 1D or 2Dr ;0D or 1D target tensor expected, multi-target not supportedsize mismatch (got input: , target: ):expected total_weight to be a single element tensor, got: z (z elements)rNzmaximumlog1pnew_fulllogr)rlrrrrr(r(r)binary_cross_entropyfs r)rcrlrrrrLcCsRd}|||tj|d||d}|dur6||}|tjjkrN||}|S)Ng-q=r ry)r>rrr&rmr)rcrlrrrZEPSILONrr(r(r)binary_cross_entropy_backward|s "  r)rrrrLcCs tt| |}t||Sr5)r>rr_r)rrrrr(r(r)soft_margin_losssr)rcrlrrrLcCs6||t||d}|tjjkr2||}|SrW)r>rrr&rmr)rcrlrrrSr(r(r)soft_margin_loss_backwards   rr!rotherpcCstj|||dS)N)r)rtrrr(r(r)distsr)x1x2rLc Cs|ddd}tj|tjd}|ddd}tj|tjd}t|d||gd}t|||gd}||j}| d S)Nr!rNT memory_formatr) powrr> ones_likecontiguous_formatrrmatmulmT clamp_minsqrt) rrZx1_normZx1_padZx2_normZx2_padZx1_Zx2_rr(r(r)_euclidean_dists r)rc input_sizesrKstartendstepcCs||}t||||||Sr5) new_zerosr> slice_scatter)rcrrKrrrrSr(r(r)slice_backwards r r )rlrKrrrcCsddlm}m}|}|dkr(tdt||}t|}t| } |dkr`td|durl|nd} |dur||nt j } || dkr| ||7} || dkr| ||7} || dkrd} n|| ||kr||} || t j kr||} n.|| | kr| } n|| ||kr"||} | | | |} | | } | |d|||<| ||9<|j rttdn||| | SdS)Nr)guard_size_obliviousstatically_known_truez,slice() cannot be applied to a 0-dim tensor.zslice step must be positiver z.clamp_wrapr)rint)r1rKrrr r(rr)_normalize_start_ends  r")rsrcrKrrrc CsDt|j|}|j|}t||||\}}t|j}|||d|||<||}|dkrv||krv|dkrv|Sdg|}t j ||j d} | ||||<t j ||j t j d} |dkrt | | |k} ||krt | | |k} |dkr t | | ||dk} dg|} d| |<| | } t| t|| |d|S)Nr rdevicer%rrN)rArrrr"rexpandclonerKr>aranger%onesbool logical_andviewrtr`_unsafe_masked_index) rr#rKrrrrZsrc_sizeindicesidxmaskZ mask_shaper(r(r)r s,      r )rcrrKindexcCs||}t||||Sr5)rr>select_scatter)rcrrKr2rSr(r(r)select_backwardEs r4)rcroffsetdim1dim2cCs||}t|||||Sr5)rr>Zdiagonal_scatter)rcrr5r6r7rSr(r(r)diagonal_backwardLs r8rcrS input_dtypecCs|j|kr||}|Sr5)rr7r9r(r(r)_cast_grad_to_input_dtypeUs  r;)rcoutputrKr:cCs0||}||tj||dd}t|||SNTr)r>rr; contiguous)rcr<rKr:Znew_grad_outputrSr(r(r)_softmax_backward_data]s  r?cCs*|t|tj||dd}t|||Sr=)r>r_rr;)rcr<rKr:rSr(r(r)_log_softmax_backward_dataos r@c CsZ||d||d}ttjtj|d}|d||d}|d|||d} || S)z/Utility function to implement im2col and col2imr!r rr%rrN)rr>r)int64rP) Zinput_dZkernel_dZ dilation_dZ padding_dZstride_dr%Zblocks_dZ arange_kwZblocks_d_indicesZ kernel_gridr(r(r) _im2col_col2im_indices_along_dim{s rC)r kernel_sizedilationpaddingrrLc s(ttdkddttdkddttdkddttdkddddd }|d |d |d d d|d|jt}t|dvotddddDfddtddtddDttddDfdd|dk}|s@|d}|j\}} } } \} } \}}\}}\}}t| |||| |j }t| |||| |j }t |||||f}|dd}|dddd||f}| dddddd}| d}| d}||| ||||}|s$|d}|S) Nr!cSsdS)Nz"im2col(): only 2D kernel supportedr(r(r(r(r)rqrrzim2col..cSsdS)Nz$im2col(): only 2D dilation supportedr(r(r(r(r)rqrrcSsdS)Nz#im2col(): only 2D padding supportedr(r(r(r(r)rqrrcSsdS)Nz"im2col(): only 2D stride supportedr(r(r(r(r)rqrrTcSs<|rtdd|Dntdd|D}t|dddS)Ncss|]}|dkVqdSrNr(r0rr(r(r) rrz1im2col..check_positive..css|]}|dkVqdSrGr(rHr(r(r)rIrrcSsdS)Nz<{param_name} should be greater {'than' zero, but got {param}r(r(r(r(r)rqrrz0im2col..check_positive..rr>rsparam param_namestrictcondr(r(r)check_positives(zim2col..check_positiverDrErFFrNrrxrcss|]}|dkVqdSrGr(r0dr(r(r)rIrrzim2col..rcsdtS)NzmExpected 3D or 4D (batch mode) tensor for input with possible 0 batch size and non-zero dimensions, but got: tupler(rr(r)rqscss<|]4\}}}}}d|d|||dd|VqdS)r r!Nr(r0rpadZdilZkerstr(r(r)rIs rcss|]}|dkVqdSrGr()r0cr(r(r)rIrrc s6dtdddddddd S) Nz!Given an input with spacial 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.rUr(rErD output_sizerFrrr(r)rqsrrrNr rx)T)r>rsrrrrVziprPrCr%rrYpermuterrsqueeze)rrDrErFrrPr batched_inputZ batch_dimrZinput_hZinput_wstride_hstride_w padding_h padding_w dilation_h dilation_wkernel_hkernel_wZblocks_row_indicesZblocks_col_indicesZ padded_inputr<Znum_blocks_rowZnum_blocks_colr(r`r)im2colsd           ro)rrarDrErFrrLc sttdkddttdkddttdkddttdkddttdkddd$d d }|d |d |d dd|d|d|jt}t|dvotddddDfdddd}td|dkfddddtD} | d| dtdkfddtdkfdd|dk} | s|d}|j\} } \} }\}}\}}\}}|dd|g| }|dddd dd!}t | |||| |j }t |d }t | |||||j }d"dtD}| ddt g|}dd||f}tj|||dd#}t|| | | | f}| s|d}|S)%Nr!cSsdS)Nzonly 2D output_size supportedr(r(r(r(r)rqrrzcol2im..cSsdS)Nzonly 2D kernel supportedr(r(r(r(r)rqrrcSsdS)Nzonly 2D dilation supportedr(r(r(r(r)rqrrcSsdS)Nzonly 2D padding supportedr(r(r(r(r)rqrrcSsdS)Nzonly 2D stride supportedr(r(r(r(r)rqrrTcSs<|rtdd|Dntdd|D}t|dddS)Ncss|]}|dkVqdSrGr(rHr(r(r)rIrrz1col2im..check_positive..css|]}|dkVqdSrGr(rHr(r(r)rIrrcSsdS)Nz9{param_name} should be greater than zero, but got {param}r(r(r(r(r)rqrrz0col2im..check_positive..rJrKr(r(r)rPs(zcol2im..check_positiverDrErFFrQrra)r!rxcss|]}|dkVqdSrGr(rSr(r(r)rIrrzcol2im..rcsdtS)NzmExpected 2D or 3D (batch mode) tensor for input with possible 0 batch size and non-zero dimensions, but got: rUr(rWr(r)rqsrr csdddS)Nz|Expected size of input's first non-batch dimension to be divisible by the product of kernel_size, but got input.shape[-2] = rz and kernel_size=r(r()rDrr(r)rqscSs:g|]2\}}}}}d|d|||dd|qSr r!r(rXr(r(r)r3 s zcol2im..rNcs4ddddddddd S NzGiven output_size=r\r]r^r_z , expected input.size(-1) to be but got rN.r(r(LrErDrarFrrr(r)rqscs4ddddddddd Srqr(r(rtr(r)rqsrxrrbcSsg|]\}}|d|qSr!r()r0orr(r(r)r36rr accumulate)T)r>rsrrrrcrPrrdrCr%rRrprodrt_unsafe_index_putrrYre)rrarDrErFrrPrZprod_kernel_sizecolrfZout_hZout_wrgrhrirjrkrlrmrnZ indices_rowZ indices_colZoutput_padded_sizer<r0r(rtr)col2ims            "   r})rcr1recCs$||||jt|d}|SNr)type_asr(rAr)rcr1rerEr(r(r)native_dropout_backwardCs r)r input_size dimensionrrrLc Cst|dkrt|dStt||}tj|||jtjd}|d|| }| d|d ||d}| |}d||f}t j |||ddS)Nrr&rNr r5Trx)rr>Z squeeze_copyrArr)r%int32ZunfoldflattenZmovedimrrtr{r>) rrrrrrKr0rSr2r(r(r)unfold_backwardRs   r)rcrlepsrLc Csx|dur>|}d|}tt||k||k||d|dStt|dk|dk||d||dtdSdS)Nr^rr(nan)r>r`r,rr)rcrlrlohir(r(r)logit_backwardesrrrtraincCs*|r|dkrt|||dS|SdSri)rtnative_dropoutr(rr(r(r)dropoutzs rout0out1cCs|rl|dkrl|dkr.t|tj|tjdfS|jjs>tdt||k}||tdd|}||fS|tj|tjdfSdS)Nrr rz?result type Float can't be cast to the desired output type Longr^) r>rr+ris_floating_pointr rand_likerr)rrrZ bool_maskresr(r(r)rs r)r1rK half_to_floatcCs|}|r|jtjksJtj|tjjd\}}||}| dkrTt |}ntj ||dd}t ||}|tj ||dd}|s||}|SNr4rT)r) r>rr>halfrArBELEMENTWISE_TYPE_PROMOTION_KINDDEFAULTr7rr_amaxr)r1rKrr:r<Z unnormalizedx_maxrr(r(r)_softmaxs     r)rc Cs|}|r|jtjksJtj|tjjd\}}||}| dkrN|}ntj ||dd}||}t tj t ||dd}||}|s||}|Sr)r>rr>rrArBrrr7rrrrr_) r1rKrr:r<ZshiftedrZshifted_logsumexprr(r(r) _log_softmaxs     rrN)rr/ padding_idxscale_grad_by_freqsparserLcCsN|dksJd|jdkrB|d|}|jdkr>|d}|S||SdS)Nr!z'weight' must be 2-Dr r)rKrZ index_selectre)rr/rrrrr(r(r) embeddings     r)rcr/ num_weightsrrc Cstj|tjjd\}}||}t|tj}|rp||f}t |}t j ||g|dd}||} || d}t ||k|j} || d} ||f|j|jd} t j | |g| dd|S)NrTrxrNr)rArBrrr7rr>longrrrtr{rPrRr masked_fillr) rcr/rrrr:r<countsr*Zgrad_weights_scaler1r grad_weightr(r(r)embedding_dense_backwards&       rr8cCsd}|D] }||9}q|SrWr()r1rEir(r(r)rzs rz)tensorsrK num_chunksrLc Csg}|D]}|}|||d||}|||krndgd|j|dd|||g}t||d}|d|t|dg}|||q|S)Nr rr!rN)rrrtconstant_pad_ndr>Sizeappendr) rrKrpadded_tensorstensor tensor_sizeZ pad_along_dimrYZ view_sizer(r(r) _pad_chunks  r)rcCs(|dj}|D]}|j|krdSqdS)NrFTr)rrrr(r(r)have_same_ndimss   r)rrKcCsB|dd|}|D]$}t|d||kddqdS)NrcSsdS)NzG_chunk_cat expects same sizes of 0,...,dim-1 dimensions for all tensorsr(r(r(r(r)rq!rrz+leading_dimension_matches..)rr>rs)rrKZleading_dim_sizesrr(r(r)leading_dimension_matchess r)rrKrcCst|dkddtt|dkdd|dj}|dj}|D]H}t|dkddt|j|kddt|j|kddqDt|rt|d |}n4t|dkd d|D]}t||j kd dqt |||S) Nr cSsdS)Nz&_chunk_cat expects positive num_chunksr(r(r(r(r)rq*rrz._preprocess_chunk_cat_inputs..rcSsdS)Nz0_chunk_cat expects a non-empty input tensor listr(r(r(r(r)rq,rrcSsdS)Nz#_chunk_cat expects non-empty tensorr(r(r(r(r)rq1rrcSsdS)Nz8_chunk_cat expects all input tensors with the same dtyper(r(r(r(r)rq4rrcSsdS)Nz8_chunk_cat expects all inputs tensors on the same devicer(r(r(r(r)rq8rrcSsdS)NzK_chunk_cat expects non-negative dim when input tensors have different ndimsr(r(r(r(r)rq?rrcSsdS)Nz3_chunk_cat expects dim < ndim for all input tensorsr(r(r(r(r)rqDrr) r>rsrrr%rrrArrKrr)rrKrZexpected_dtypeZexpected_devicerr(r(r)_preprocess_chunk_cat_inputs%s:   r)rrKrrrLcCsLt|||}t|||}|dur0t||dStj||d|d|SdS)Nr )r)rrr>r)rrKrrrr(r(r) _chunk_catJs   r)rl split_sizesrKrrLcCs\tj|||d}|dur&dd|DSt||D]"\}}t||jt||ddq0dSdS)NrpcSsg|]}|jtjdqS)r)r(r>r)r0sr(r(r)r3frrz)split_with_sizes_copy..Tr)rtsplit_with_sizesrcrrr)rlrrKrsplitsr<splitr(r(r)split_with_sizes_copy[s  r.)r split_sizerKrLcCstj|||Sr5)rtrr)rrrKr(r(r) unsafe_splitnsr)rrrKrLcCstj|||Sr5)rtrr)rrrKr(r(r)unsafe_split_with_sizesssr)rlrrKrLcs|j}||}dkr,|dks"J|fS|d}ddlm}||}fddt|D}|||d<t|||S)Nrr ) guard_intcsg|]}qSr(r(r0rrr(r)r3rrzsplit..rN)rdetachr rrOr>r)rlrrKrrchunksrrr(rr)rzs   r)rltensor_indices_or_sectionsrKrLcs|jjdksJ|jtjks J|tdkp:dkfdddkrr|}t|t sfJ| ||Sdd|D}| ||SdS)Ncpur rcs ddS)Nz{tensor_split expected tensor_indices_or_sections to be a zero-dimensional or one-dimensional tensor, but got a tensor with z dimsr(r(Z split_dimr(r)rqszAtensor_split_tensor_indices_or_sections_py_impl..cSsg|] }|qSr()itemrr(r(r)r3rrzCtensor_split_tensor_indices_or_sections_py_impl..) r%typerr>rBrKrsrr.r tensor_split)rlrrKsectionsr/r(rr)/tensor_split_tensor_indices_or_sections_py_impls   r)rlmat1mat2r[rdcCsH|s |s t|}t|}|t||}|dkr<|S|||Sri)r is_complexr!r>mm)rlrrr[rdrr(r(r)addmmsr)rlrrr[rduse_gelucCs<t|||||}|r2|jr(tj|ddSt|St|S)Nr)r)ris_cudartgelurelu)rlrrr[rdrrr(r(r)_addmm_activations  r)rlrvecr[rdcCs\|s |s t|}t|}|t||}|dkr<|S|dkrP||S|||Sri)rrr!r>mvr)rlrrr[rdrr(r(r)addmvs r) rcrrrstdgammaNCHxWgroup output_maskrLc  stj|||ddtj||ddtj|ddt|kfddtjfkfddtdupkfddt\} } t| dkfddt|| j d gd } | j d gd } d}d}d}| drvd | }durt|  d |  d }t|  d |  d }t| d  d | }nL| |  d }| |  d }t| d tj d | f|jd}||||||}| |||}| d }t|d}t|d}t| | |t| | ||}| |j|j}| d r| | | |  d | d j dgd  }| d r| j dgd }|||fS)NF)Zallow_cpu_scalar_tensorscsddS)NzExpect input to have z elementsr(r()rrrr(r)rqrrz,native_group_norm_backward..csdddjS)NzExpect mean to have shape (, z , but got rWr()rrrr(r)rqrrcs dddurndS)NzExpect gamma to have z elements but got rN)rr()rrr(r)rqrrrcsddS)NzExpect number of channels z, to be evenly-divisible by number of groups r(r()rrr(r)rq rrr!rpr^rNr r$r)rAZcheck_same_deviceZcheck_same_shaper>rsrrdivmodrr-rrPrr*r%rRr7r)rcrrrrrrrrrZcpgZ_remZdsdbd_inputZd_gammad_biasrZds_valZdb_valc1c2c3r()rrrrrrr)native_group_norm_backwards        ""     $ r)rcrrrrrrrrrrrout2rLc  Csdt|||||||||| } | | | f}t| D]2\}}|dur,t|||jt|||ddq,|Sr)r enumeraterrr)rcrrrrrrrrrrrrrrSrrEr(r(r)native_group_norm_backward_out?s r)r1rLcCs|dur||S|Sr5r7)r1rr(r(r) _maybe_cast\s r) grad_outrnormalized_shaperrrbiasrrLc"sn|j}|} t|jfdd||||fD\} } } } | dusHJ| t|}||d}|d|}g}g}t| D]"}||kr||q|||q|t|}t|}ddl m }||dks||dkr*|dr| |nd|dr| ||dnd|dr$| ||dndfSt || }t || }| ||}| durf| | }n| }||}t ||d}t ||}t ||d}t ||}|||}d}d} d}!|dr|||}|dr| durt|dkr t | ||d} n| |} |drL| durLt|dkrDt | |d}!n| }!t||jt| |jt|!|jfS) Nc3s*|]"}|dur|jtjdn|VqdSr~)r7r>rr/r9r(r)rIqsz-native_layer_norm_backward..rr r r!TF)rrKrAget_computation_dtyperrrOrrzr r rrRr>rrr(r)"rrrrrrrr input_shapeZ input_ndim grad_out_cast input_cast weight_castZ bias_castaxisZ inner_dimsZ outer_dimsZinner_dim_indicesZouter_dim_indicesrrMr Zx_hatZ grad_x_hatabrrrrGrZd_weightrr(r9r)native_layer_norm_backwardcsj                       r) rrrrrrrrrrrrLc Cs`t||||||||} || | f} t| D]2\} }|dur(t| | |jt|| | ddq(| Sr)rrrrr)rrrrrrrrrrrrrSrrEr(r(r)native_layer_norm_backward_outs r) rrr running_mean running_varrmomentumr functionalrLc Csbdgttd|} t|j} |} |} |rt|j} |j| d} tj| | ddd\}}t ||}|||}t || }t || }|dur||d||} |s| | |dur| |j d}t || }|||d}||d||} |s| | n|dur.|dus2J|j| dd}|} |j| dd}|} |}dt||}|jjdkr|}|}n|d }|d }t||d}t||d}|||}|dur|}t||d}||}|dur |}t||d}||}|jjdkrJ|j|jd}|j|jd}|j|jd||| | fS) Nrr!rT)rKZ correctionrr )rrurr)rrOrKrArrr7r>Zvar_meanrsqrtrecopy_rrrr%rrrRr)rrrrrrrrrZreduction_dimsr:new_running_meannew_running_varZ input_accZ biased_varrrr< save_mean save_rstdnZ squeezed_varZ unbiased_varinvstdr(r(r)native_batch_norm_helpersr                  r rr save_invstd) rrrrrrrrrLc Cs,t||||||||d \}} } } } || | fSNFr  rrrrrrrrr<rrrQr(r(r)native_batch_norm#s rc Csz|dur$|dur$t||||||S|dur4td|durDtd|r`t||||||||St|||||||SdS)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.)rt_native_batch_norm_legitr$_native_batch_norm_legit_no_training)rrrrrrrrr(r(r)native_batch_norm_decomposition?s&  rrcs|||}||d|dkrh|dkrhfdd|D}||||d<tjjj|||Stjjj||S)Nr rcsg|]}qSr(r(rrr(r)r3jrrz(unsafe_chunk_py_impl..)rr>opsrtrrrr)rrrKrrr(rr)unsafe_chunk_py_implds r)rrrrrrrrLc Cstj|||||d||Sr )rtrr)rrrrrrrr(r(r)rps rc Cs,t||||||||d \}} } } } || | fSr rrr(r(r)rs r)rrrrrrrLc Cs,t|||dd|||d \}}}} } |||fSr r) rrrrrrr<rrrQr(r(r)!_native_batch_norm_legit_no_statss rc CsPt||||||||d \}} } } } | dus2Jd| dusBJd|| | | | fS)NT#new_running_mean should not be None"new_running_var should not be Noner) rrrrrrrrr<rrrrr(r(r)#_native_batch_norm_legit_functionalsr)rrrrrrrrLc CsPtj|||||d|}d}|tjjjkr8tj||}tj|tj|j|j dS)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>_CZ_select_batch_norm_backendZ_BatchNormBackendZCudnnZ(_get_cudnn_batch_norm_reserve_space_sizeemptyuint8rr%) rrrrrrrbackendZ reserve_sizer(r(r)_get_batch_norm_reserve_tensorsrc CsDt|||||d||d \}}} } } t||||||dd} ||| | fS)NTFrr r rrrrrrrr<rrrQreserver(r(r)_batch_norm_with_updates  r$c Csht|||||d||d \}}} } } t||||||dd} | dusHJd| dusXJd||| | | | fS)NTr rrr!) rrrrrrrr<rrZnew_rmZnew_rvr#r(r(r)"_batch_norm_with_update_functionalsr%c CsDt|||||d||d \}}} } } t||||||dd} ||| | fS)NFr r!r"r(r(r)_batch_norm_no_updates  r&cCsB|dus Jt||kjtjd}|||d|}||fS)Nrr^)r>rr7rr)rr generatorr1rr(r(r)_fused_dropout_decomposition-s r()rrr% pin_memory non_blockingr)r1rr%r)r*rc Cs |r|tjksJd|r"Jdt|tjttttfsstridedr.rr!rr+complexr( scalar_tensorr%r_primsZconvert_element_typeZ device_put) r1rrr%r)r*rZdtype_convertedZx_tensorr(r(r)_to_copy7s,       r/cCs t|Sr5)rtaliasr8r(r(r)nop_decompositionesr1rZout3)rrrrrrexponential_average_factorepsilonc Cs^t||||||||\}} } |r:|| | |jdtjdfS||d|d|jdtjdfS)Nrr)rtrrr>r) rrrrrrr2r3rrr[r(r(r)cudnn_batch_normms"  r4cCs@t|D]2\}}|dkr||jkr0|j||ks||}q|SrW)rrrrP)r1broadcast_maskrr1r(r(r)_broadcast_batch_norm_backwards  r6) rrrrrrr rrrr#rLc Cst|||||||||| Sr5)native_batch_norm_backward) rrrrrrr rrrr#r(r(r)batch_norm_backwardsr8) rrrrrrr rrrrLc &s|j} |dur|j} n| } t|jfdd|||||||fD\} } }}}}}|j}|}|dksrJdd}tt|||}|}|}|r|dur|dusJn&|dur|dusJ|}t||}dg|}||||<g}t |D]}||kr| |qt ||}d|}t | |}t | | ||}t |||}t t |||||} |durt ||d}!nt |||}!|r| || }"| |"||!}#n| |!}#| dr||}$nd}$| dr|}%nd}%|#| t|$| t|%| fS)Nc3s$|]}|dur|n|VqdSr5rr/r9r(r)rIsz-native_batch_norm_backward..r!z$rank of the input must be at least 2r r^)rrArrrKrzrr>rrOrr6rrr7r)&rrrrrrr rrrr:Z weight_dtyperrrZrunning_mean_castZrunning_var_castZsave_mean_castZsave_invstd_castrZ input_rankrZ num_featuresrr r5Zreduction_axesrrZgrad_output_sumZdot_pZ grad_meanZ proj_scaleZ grad_scaleZprojrSrZ grad_biasr(r9r)r7s              r7)rrrrrrr rrrrrrrLc  Csdt|||||||||| } | | | f}t| D]2\}}|dur,t|||jt|||ddq,|Sr)r7rrrr)rrrrrrr rrrrrrrrSrrEr(r(r)native_batch_norm_backward_out s$ r9rrcrrrrsave_varr3c Cs t|||||||d|gd SNT)TTTrtr7r:r(r(r)miopen_batch_norm_backwardA s r> rrcrrrrr;r3Z reserveSpacec Cs t|||||||d|gd Sr<r=r?r(r(r)cudnn_batch_norm_backward[ s r@)rracs|j|jttdvfdd|jddD]}t|dkfddq:d|ddkrԈd|ddkrtdd tdd|D}td d tdd||D}tjj |||Sd d d dfdd}|d|d\}}}} |d|d\} } } } |dt |d| f}| sV| sVtj |ddSdd}||||| dd\}}||| | | dd\}} d}t t |jdt |jdD]B\}}|dur|d|dd|f}n||d|dd|f}q||| S)NrRcs dS)Nz9adaptive_avg_pool2d(): Expected 3D or 4D tensor, but got r(r(rr(r)rq rrz%adaptive_avg_pool2d..rrcsdtdS)Nzjadaptive_avg_pool2d(): Expected input to have non-zero size for non-batch dimensions, but input has shape rsrUr(rWr(r)rq srNcss|]\}}||VqdSr5r()r0rrwr(r(r)rI rrz&adaptive_avg_pool2d..css$|]\}}}||d|VqdSr Nr()r0rrwrr(r(r)rI scSstj|||ddS)NtruncZ rounding_moder>divrrr[r(r(r) start_index sz(adaptive_avg_pool2d..start_indexcSs tj|d||d|ddS)Nr rBrCrDrFr(r(r) end_index sz&adaptive_avg_pool2d..end_indexc stj|tjd}|||}||d}||}|dkpD||dk }|rV|d7}n|dkrf|d8}tj|tjd}|d|}|rtj|d|j|jd} t|| }|||} | |} n|} || ||fS)Nr&r rrNrA)r>r)rBrPr-rr%minimum) in_sizeout_sizeZorangeZi0Z maxlengthZ in_size_modadaptive range_maxr0maxvali1length)r%rHrGr(r) compute_idx s(      z(adaptive_avg_pool2d..compute_idx.r)rrNrpcSsdt|tr||fS|dksJ||dk}|dkr>t|d}t||d}t|| }||fSdS)NrrNrrr)r.rrPrRr>r)valsrPrMrLrKr1r(r(r) maybe_mask s    z'adaptive_avg_pool2d..maybe_mask)rLrKr)r%rrr>rsrVrcnnrZ avg_pool2drRrrrO)rrarTrkernelrQZidxhZlength_hZ range_max_hZ adaptive_hZidxwZlength_wZ range_max_wZ adaptive_wrRrSretrjr()r%rHrrrGr)adaptive_avg_pool2dv sN  (     & rX)rlr/rarKc Cstd|dttj|jd| }ttj|}dg|j}|jd| |d| <|tj||j d || d}| t |jd| t |}tj| d|g| ddd |jS)NZ max_unpoolingZ d_forward_outr r$rNFrx)rAZalert_not_deterministicroperatorrrrrtr)r%r-rrrr{) rlr/rarKZncZhwZindices_nc_shapeZ indices_flatr<r(r(r) _max_unpoolnd s   "rZrlr/racstjtjkfddttdkfddtjdvfddtjjkfddtdjD]$td kfd dq|t dS) Ncs djS)Nz2elements in indices should be type int64 but got: rr()r/r(r)rq rrzmax_unpool2d..r!csdtdS)NzMThere should be exactly two elements (height, width) in output_size, but got elements.rr(rar(r)rq srRcsdjdS)NzLInput to max_unpooling2d should be a 3d or 4d Tensor, but got a tensor with dimensions.rr(rr(r)rq scsdjdjSNzBExpected shape of indices to be same as that of the input tensor (z%) but got indices tensor with shape: rWr()r/rlr(r)rq s r rcsdjddS)NzZmax_unpooling2d(): Expected input to have non-zero size for non-batch dimensions, but got  with dimension being empty.rWr()rrlr(r)rq s ) r>rsrrBrrrrOrrZr[r()rr/rarlr) max_unpool2d s,          rcrr/rarrFcstjtjkddtjdvfddttdkfddttdkfddttdkfddtjjkfd dtd jD]$td kfd dqtd d kod d kod d kfddt dS)NcSsdS)Nz(elements in indices should be type int64r(r(r(r(r)rq0 rrzmax_unpool3d..rrbcsdjdS)NzLInput to max_unpooling3d should be a 4d or 5d Tensor, but got a tensor with r_rr(rr(r)rq4 rrrxcsdtdS)NzVThere should be exactly three elements (depth, height, width) in output_size, but got r\r]r(r^r(r)rq8 scsdtdS)NzRThere should be exactly three elements (depth, height, width) in stride, but got: r\r]r(rr(r)rq? rrcsdtdS)NzSThere should be exactly three elements (depth, height, width) in padding, but got: r\r]r()rFr(r)rqC rrcsdjdjSr`rWr()r/rr(r)rqG s r rcsdjddS)NzZmax_unpooling3d(): Expected input to have non-zero size for non-batch dimensions, but got rarbrWr()rrr(r)rqP s r!cs dS)Nz5strides should be greater than zero, but got stride: r(r(rgr(r)rqY rr) r>rsrrBrrrrOrrZrdr()rr/rrarFrr) max_unpool3d& sB             " rh)rdr1rKr2rrdcCst||||d|dS)NTinplacerd _index_addrir(r(r) index_add__ s rncCst||||d|dS)NFrjrlrir(r(r) index_addk s ro)r1rKr2rrkrdc s*t|jtjdkfddjdkr<dnd|jdkrT|ndtkfdddkrt|jttkpt t fdd|}|jdk}|r| dn|}df}|rt j nt j} | |||dd } |r|S|r| dS| SdS) Nr csdjdSNz(Index should have dimension 1 or 0 (got rrr(r2r(r)rq rrz_index_add..rcsdddS)NzNumber of indices (z') should be equal to tensor.size(dim) (z ), for dim=r(r()rK index_sizerr(r)rq rrcsdtddS)Nzalpha argument of type z cannot be safely cast to type !)rr()rd python_typer(r)rq rrr5Trx)rAcanonicalize_dimsrr>rsrZ dtype_to_typerr+Zis_weakly_lesser_typerrPrt index_put_ index_putrer>) r1rKr2rrkrdzero_dimrr0rwrr()rdrKr2rrrtrr)rmx s6     rmrc Cstt|dkddt|}|d}|dd}tdd|D}|rX||f}n||f}||}|d||}dt|} t|D]V} || } t| | d|| df|} |rtj || d| d}qtj || d| d}q|S) NrcSsdS)Nz#received an empty list of sequencesr(r(r(r(r)rq rrzpad_sequence..r css|]}|dVqdSrGrr/r(r(r)rI rrzpad_sequence..)rrrKr2) r>rsrrr}rrOrtrr3) sequences batch_firstZ padding_valueZsequences_sizemax_sizeZ trailing_dimsmax_lenZout_dimsrZ dim_paddingsrZcurrseqrowr(r(r) pad_sequence s(     rr1rKr2rcCst||||ddS)NTrk _index_copyrr(r(r) index_copy_ srcCst||||ddS)NFrrrr(r(r) index_copy sr)r1rKr2rrkc st|j|}tjdkfdd|jdk}|r@|dn|}jdkrXdnd|f}|rttjntj}||||} |r|S|r| dS| SdS)Nr csdjdSrprr(rqr(r)rq rrz_index_copy..rr5) rArurr>rsrPrtrvrwrer>) r1rKr2rrkrxrr0rwrr(rqr)r s   rr<rcCsRt|d|}tt| }|js0|jr<|d}n|}|t||fS)Nr(r)r>rIrr_rrZis_xpur)rlrzrarr(r(r)log_sigmoid_forward s   rr1lowhighr'cCs$tj|jt|t||j|j|dS)N)rrrr%r')primsZ_uniform_helperrr rr%rr(r(r)uniform srcCs|t||||Sr5)rr)rlrrr'r(r(r)uniform_ srcCst|d}|durDt|duddtt||kdd|S|durt|duddtt||kddg}t|D]J\}}t||kr|||dt|q|t||d|q|StddddS) Nr!cSsdSNz9Must specify exactly one of output_size and scale_factorsr(r(r(r(r)rq rrz.upsample_compute_output_size..cSsdSNr(r(r(r(r)rq rrcSsdSrr(r(r(r(r)rq rrcSsdSrr(r(r(r(r)rq rrFcSsdSrr(r(r(r(r)rq rr)rr>rsrr!rr)rra scale_factorsZspatial_dimensionsrrr(r(r)upsample_compute_output_size s.  rcCs|dur dS||Sr5r()scalesr0r(r(r)get_scale_value sr)rrarrLcCs2t|||}|r|n dgt|}t|||Sr5rrr_upsample_nearestrrarosizerr(r(r)_upsample_nearest_vec# srcCs6t|||}|r|n dgt|}t|||ddSNTexactrrr(r(r)_upsample_nearest_exact_vec8 srcCsg}t|}|rdnd}t|D]}||}|j| |} ||durX| | ||n| |} tj|tj|jd} | || tj} t|d|D]} | d} q| | q |S)NrrrAr rN) rrOrr>r)rr%r7rBrPr)rrarrr/Znum_spatial_dimsr5rTrisizereZoutput_indicesZ input_indicesrQr(r(r)!_compute_upsample_nearest_indicesM s  $  r)Zpreserve_memory_formatr)rrarrLcCst|||gSr5rrrarr(r(r)upsample_nearest1dm s rcCst|||gddSrrrr(r(r)upsample_nearest_exact1dy s r)rrascales_hscales_wrLcCst||||gSr5rrrarrr(r(r)upsample_nearest2d s rcCst||||gddSrrrr(r(r)_upsample_nearest_exact2d s r)rrascales_drrrLcCst|||||gSr5rrrarrrr(r(r)upsample_nearest3d s rcCst|||||gddSrrrr(r(r)_upsample_nearest_exact3d s r)rrarrrLc Cspt||||d}ddg|}t||}|jdkrlt|}|jd}|jjdkr`|dkr`t j }|j |d}|S)Nrrr cudar) rrt _unsafe_indexrrArrr%rr>rr>) rrarrZspatial_indicesr/rr n_channelsr(r(r)r s      rcsb|r|rdn|rdn|r"dndtdksBJtfddtdtDS)Nrbrrxr!rcs g|]}t||qSr(rUrZ group_sizeparamsr(r)r3 sz!gather_params..)rrO)r has_biaseshas_projectionsr(rr) gather_params s rcCsh|rB|d||d|}}|d|d|d|d}}n||||}}d\}}||||fS)Nr!r NNr()rhiddensr bidirectional cur_params cur_hidden bidir_params bidir_hiddenr(r(r)params_hiddens s $rcCs2||ks J||d||||dd|Sri)rr)rlast_batch_size batch_sizerr(r(r)update_hidden_for_packed s rc Cs4||kr |S||ksJt||d|||fSri)r>concatr)rrrZ inp_hiddenr(r(r) update_hidden_for_packed_reverse s rcCs&|d}|d}|r|dnd} |r,|dnd} g} g} |rD|dn|d} |dd| }t|t|}|r||ddd}|D]Z}|jd}| |krn"|rt|| ||}nt|| || }||||| || }|} | |q|r| n| || t | d}|st | dn|}||fS)Nrr r!rxrN) rr>rrrrrrreverser)inphiddenrr hidden_fn batch_sizesr ih_weight hh_weightih_biashh_bias step_outputrrr split_inprr hidden_outr(r(r)one_layer_rnn_data s@     rcsfdd}|S)Ncst||||Sr5rlinearrrrrrr nonlinearityr(r)rG< szrnn_cell..innerr(rrGr(rr)rnn_cell; s rcsfdd}|S)Ncs$t|||}t||||Sr5rrrr(r)rGC szrnn_cell_data..innerr(rr(rr) rnn_cell_dataB s rcCs|d}|d}|r|dnd}|r,|dnd} t|||} |rL| dn| } |d} g} | D] } || | |||| } | | qb|r| t| d}|| dfS)Nrr r!rx) rrfliprPrrr>rre)rrrrrrrrrrprecomputed_inputrrrrr(r(r) one_layer_rnnJ s   rcCs|d}|d}|r&|d}|d}nt|}t|}|dd} |dd} g} d} | d} d}d}d}d}|}| } | } tjjj|||||| | || | | |||||}|d|d|d}}}|| d| dffS)Nrr r!rxF) r>rrrPr>rrtZmkldnn_rnn_layerrre)rrrrrw0w1w2w3hxcxrmodeZ hidden_size num_layersrr|routputsrVhycyr(r(r)mkldnn_one_layer_lstm` sN   rc Cs|r|ddn|}g} t|D]} t||| |\} } }}|rN|| dkrN|nd}| || | |\}}| ||r| ||||dd\}}| ||rt||g|d}n|}|dkr |r | |dkr tj||dd}q |r|ddn|}|| fS)Nrr rT)r)r) transposerOrrr>rrKr)rrrrrrrrr|layer_fn final_hiddensrrrrrZfwd_inpZ fwd_hiddenZbwd_inpZ bwd_hiddenr(r(r) _rnn_helper s*      rc CsR|d} t||d}t|| |||||||ttttjd \} } | t| dfSNrFr) unbindrrrrrr>rstack rrrrrrrrr|rrrr(r(r)rnn_tanh_input s   rc CsR|d} t||d}t|| |||||||ttttjd \} } | t| dfSr) rrrrrrr>rrrr(r(r)rnn_relu_input s   rc CsT|d} t||d}t|| ||||||dtt|ttjd \} } | t| dfSNrFrr) rrrrrrr>rr datarrrrrrrrrrrr(r(r) rnn_relu_data s&  rc CsT|d} t||d}t|| ||||||dtt|ttjd \} } | t| dfSr) rrrrrrr>rrrr(r(r) rnn_tanh_data s&  rcCst||||}|d|}|d} |d} |d} |d} | || | } | | }|durv|n t||d}|| fS)Nrrr r!rxrrchunkrr)rrrrr hr_weight chunk_dimZgatesZ chunked_gatesZin_gateZ forget_gateZ cell_gateZout_gaterrr(r(r) lstm_cell? s      rc Cs|d}|d}|r|dnd}|r,|dnd}t|dkrD|dnt|dkrX|dnd} |dd} |dd} t|||} |r| dn| } g} | D](}t|| | ||| dd\} } | | q|r| t | d}|| d| dffS)Nrr r!rxrbrr) rrPrrrrrrr>rre)rrrrrrrrrrrrrrrr(r(r)one_layer_lstmM s$*  rc Cs|d}|d}|r|dnd}|r,|dnd} t|dkrD|dnt|dkrX|dnd} g} g} |rp|dn|d} t|t|}|r|ddd}|d}|d}|dd| |dd| }}|D]}|jd}t|||}|| kr:| |d|| ||d|| |f|dd||dd|}}|| krt ||d| || fd}t ||d| || fd}t ||||| | dd\}}|} | |q|r| ||f}n:| ||f| t | \}}t |dt |df}t | d}||fS) Nrr r!rxrbrrNr)rr>rrrrrrrrrrrcr)rrrrrrrrrrrrrrrZorig_hxZorig_cxrrrrZhidden0Zhidden1rr(r(r)one_layer_lstm_datah s\*         rcCs dd}||||rtStSdS)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 c SstjsdS|gt|tt|}dd|D}t|dkrHdS|}|tdkrbdSdd|D}|D]}|tj tj fvrtdSqt|j rdS|d d|d dk}|rdSd S) NFcSsh|] }|jqSr(r$r0tr(r(r) rrzEselect_one_layer_lstm_function..use_mkldnn..r rcSsh|] }|jqSr(rrr(r(r)r rrrr!T) r>rZ_get_mkldnn_enabledrr from_iterablerpopr%rbfloat16 requires_gradr) rrrrZdevicesr%Zdtypesrrr(r(r) use_mkldnn s&  z2select_one_layer_lstm_function..use_mkldnnN)rr)rrrrr(r(r)select_one_layer_lstm_function s rc Cst|dksJdt|||dd|ddk}tt|d|d} t|||} t|| |||||||| \} } tt| } | t| ddt| ddfS)Nr!lstm expects two hidden statesrr ) rrrrrcrrr>r) rrrrrrrrr|rrrrr(r(r) lstm_impl s$$  r c Cst|dksJdt|||dd|ddk}tt|d|d} t|| ||||||dtt|d \} } tt| } | t | ddt | ddfS)Nr!rrr F)r) rrrrrcrrrr>rrr(r(r)lstm_data_impl s"$  r c Csr|dd}t|||dd}|d|d}|d|d} |d|d|} || | | S)Nrxr r!r)rrrrr rrrrrrZchunked_igatesZchunked_hgatesZ reset_gateZ input_gateZnew_gater(r(r)gru_cells  r c Cs|t|||dd}t|||dd}|d|d}|d|d} |d|d|} || | | S)Nrxr rr!rr r(r(r) gru_cell_data&s r c CsJt||d}t||d||||||dtt|td \} } | t| dfS)NFrr)rrrrrr r>r) rrrrrrrrrrrr(r(r) gru_impl_data/s   rc CsHt||d}t||d|||||||tttd \} } | t| dfS)NFrr)rrrrrr r>r) rrrrrrrrr|rrr(r(r)gru_implMs   rcCs:t|||}t|d}t|d}tjj|||||SNrr )rrrr>rrt_upsample_bilinear2d_aarra align_cornersrrscale_hscale_wr(r(r)upsample_bilinear2d_aa_vecks    rcCs:t|||}t|d}t|d}tjj|||||Sr)rrrr>rrt_upsample_bicubic2d_aarr(r(r)upsample_bicubic2d_aa_vecws    rcCs4t|||}|r|n dgt|}t||||Sr5)rrr_upsample_linear)rrarrrrr(r(r)_upsample_linear_vecs r)rrarrrLcCst||||gSr5r)rrarrr(r(r)upsample_linear1dsr)rrarrrrLcCst|||||gSr5r)rrarrrr(r(r)upsample_bilinear2ds r)rrarrrrrLcCst||||||gSr5r)rrarrrrr(r(r)upsample_trilinear3ds rcCsD|r |dkr|d|dSdS|dur8|dkr8d|S||SdS)Nr r^rr()rJrKrrer(r(r)_compute_scalesrcCs |r ||S||ddSdSNrr()reZ dst_indexrr(r(r)_compute_source_indexsr!)r#weightsweights_precisionrLcCsBtddt||Dd|d>}||?}t|ddtjS)Ncss*|]"\}}|tj|tjVqdSr5)r7r>r)r0rr[r(r(r)rIsz%_sum_tensors_uint8..r r) _sum_tensorsrcr>rr7r)r#r"r#r<r(r(r)_sum_tensors_uint8s r&)r"rLcCsJt|}d}tj||jd}d|d|d>}|dk}||S)Nr$rr i)r>rr}r)r%r)r"Z max_weightZmax_weight_precisionZ precisionsvaluesr1r(r(r)_compute_weight_precisions r))rrarrrLcsjd}jdd}t|tjtjjd\}fddfddtt|||D}tt|\}g} t ddggD]Fddgfd dt D} t | } t | } | | qtt D]N} || | d d fd dt| ddd| dddD} qt| dksDJ| d} t}jjd krt|dkrttj}t| tjsJ| j|d} s| } | S)Nr r!rc st|||}tj|jdjd}t||jdd}|j|jdgdg|R}|tj }|dj|dd}|||fS)Nr$rrryrr r|) rr>r)r%r7r!rrrrB) inp_sizerKrZnsqueezeZ scale_factorrZx_f32r1Zxp1)rrrr(r) get_valuess z$_upsample_linear..get_valuescs,g|]$\}\}}}|||d|qSrr()r0rr*rKr)r+n_dimsr(r)r3s z$_upsample_linear..rcs(g|] }|dkr|n|qSrr()r0k)rxp1sxsr(r)r3rrrr^cs$g|]\}}|t||qSr()r>r)r0Zv1Zv2)xscaler(r)r3 srr)rrrArBr INT_TO_FLOATrrcrrrOrtrrrreversedrr7rr%rr>rr.rr>rround)rrarrrZ inp_sizesrQr(Zxs_f32vsr0vrrrr() rrrr+rr,r.r/r0r)rsF    "       r)rrrLcCs |j|jkSr5rW)rrr(r(r) is_same_size'sr7cGs t||Sr5)rtr-)r1rrCr(r(r)_reshape_alias,sr8cCs t||Sr5)rtr2)r1r/r(r(r)r2srcCst||||Sr5)rtrw)r1r/rmryr(r(r)r{7sr{cCs|D]*}|durt|jtjtjfvddqt|jtjkddddlm}||dkrtj ||}| |j |St t|D]0}||}|dur|jd||dd||<qt||||S)NcSsdSNz3tensors used as indices must be long or int tensorsr(r(r(r(r)rqBrrz&_unsafe_masked_index..cSsdSNz*tensors used as masks must be bool tensorsr(r(r(r(r)rqGrrrrr r)r>rsrrr!r+r r rZ_meta_registrationsZmeta_index_TensorrrrOrrrrtrr)r1r1r/fillr2r Z meta_resultrr(r(r)r.<s&  r.cCs|D]*}|durt|jtjtjfvddqt|jtjkdd|dkr\|Stt |D]8}||}|durh|j | | | |dd||<qh| |d}t j|||ddS) NcSsdSr9r(r(r(r(r)rq^rrz5_unsafe_masked_index_put_accumulate..cSsdSr:r(r(r(r(r)rqcrrrr rTrx)r>rsrrr!r+rr(rOrrrrrtr{)r1r1r/r(r2rZ masked_valuer(r(r)#_unsafe_masked_index_put_accumulateXs$  &r<)rlrrrrrLcCsV|}d}|dkrd}|durX|dkrLdg|}|jd||<||}n|}||}t||k|d} | |} t||| | } t||k| d} |tj j kr|dkr| dd} | | fS|dur | |j}t||| |} t||k| d} | } n||k |} |tjj kr4| } n|tjj krN| | } | | fS)Nr r!rr(r)rKrr-r>r`rPgatherrerr%rmrr'rr7r'r&)rlrrrrr,rrwrZ safe_target_rrZwsumr(r(r)_nll_loss_forwardrs@         r?rcCs|dkr|dks Jd|dks4Jd|dkoJ|dk}|s~|jd|jdks~Jd|jd|jd|jd }|dus|dkr||ksJd |d |jt|||||S) Nrr!rr rrrrrNz/weight tensor should be defined either for all z7 classes or no classes but got weight tensor of shape: )rKrrr?)rlrrrrrZ n_classesr(r(r)nll_loss_forwards  "r@cCst|||||Sr5)r?)rlrrrrr(r(r)nll_loss2d_forwards rA)r1ArLcCs |d||d||dS)Nr!rxr r(r1rBr(r(r)_upsample_cubic_convolution1srDcCs(||d||d||d|S)Nrbrr(rCr(r(r)_upsample_cubic_convolution2srF)rrLc Csd}|jtdkrtj|d|gdd}tj|dd|gdd}t||}t||}tj|dd\}}tj|dd\}} ||| |fSt|d|t||td||td||fSdS)Ngrr^rrpr)r%r>rrFrDr) rrBZtt1Ztt2Zw03Zw12rrrrr(r(r) _upsample_get_cubic_coefficientss      rG)coeffstsrLcCs t|}tddt||DS)Ncss|]\}}||VqdSr5r(r0rrr(r(r)rIrrz+_upsample_cubic_interp1d..)rGr%rc)rHrIZcoeffs2r(r(r)_upsample_cubic_interp1dsrK)rIrLcCs ttj|Sr5)rr>add)rIr(r(r)r%sr%) num_stepsrrr%cCsB|dkrtjd||dS|s(|d|nd}tj| ||||dS)Nr rr&)Zstepsr%r)r>rZlinspace)rMrrr%rr(r(r)_linspace_from_neg_onesrN)thetahr>rc Cs|j}|j}t||||d|d}t|||||dd}tjd||d}tjjj|dddd}tjjj|dddd}tjjj|d ddd}|||S) Nr )r r r rA)rr!constantrrYrrm)r r )r!r rr%rNr-r>r*rTrrY) rOrPr>rrr%grid_xgrid_ygrid_oner(r(r)_make_base_grid_4dsrW)rOrTrPr>rc Cs|j}|j}t||||dd|d}t||||d|dd}t|||||ddd} tjd||d} tjjj|dddd}tjjj|dddd}tjjj| d ddd} tjjj| d ddd} ||| | S) Nr )r r r r rA)rrxrQrrRrpr!r )rxrrS) rOrTrPr>rrr%rTrUZgrid_zrVr(r(r)_make_base_grid_5dsrYrOrrc CsL|\}}}}t||||d}|ddd|jdd}||||dS)NrrNrxr rr!)rWr-rrPr) rOrrr rQrPr> base_gridgridr(r(r)_affine_grid_generator_4d$s  r^c CsR|\}}}}}t|||||d}|ddd|jdd} | ||||dS)Nr[rNrr rrx)rYr-rrPr) rOrrr rQrTrPr>r\r]r(r(r)_affine_grid_generator_5d.s r_cCsDtt|dvddt|dkr2t|||dSt|||dSdS)NrecSsdS)NzCaffine_grid_generator needs 4d (spatial) or 5d (volumetric) inputs.r(r(r(r(r)rq>rrz'affine_grid_generator..rr[)r>rsrr^r_rZr(r(r)affine_grid_generator8s  r`)rr]interpolation_mode padding_moder _expand_gridrLcs(tdvfddtdvfddtttdfdd ttttdd d tttdfd d tttdfd d }j\ |j\}}|dksJr|d|d}tttd fdd tjjddddtjjddddtttt dfdd tttdfdd |d} |d} dkrB|| } || } | | d} }d}}| |}}|| || }| ||| }| | | |}| | }t fdd|f| ||f|||f|||ffDSdkr||| } || } | }| } ||dS| } | } | | | | }sʈ d| d}tttd fd d! ttd" fd#d$ t fd%dtd&D}t||SdS)'N)rr r!cs dS)NzInvalid interpolation mode r(r()rar(r)rqWrrz"_grid_sampler_2d..cs dS)NzInvalid padding mode r(r()rbr(r)rqZrr)coordsrrLcs0r|ddn|d}|dd}|||Sr r()rdrrofsr[r(r) unnormalize]s z%_grid_sampler_2d..unnormalize)rd twice_low twice_highrLcSsv||krt|S|d}||d}||}t||}||jtjd}t|d@dk|||||S)Nr!rr r)r>rrfmodfloorr7Zint8r`)rdrgrhZ coords_minZ coords_spanZcoords2extraZflipsr(r(r)reflect_coordinateshs    z-_grid_sampler_2d..reflect_coordinatescsjdkr |Sdkr&t|d|dSr@|dd|d}n|dd|d}t|d|dSdS)Nrr r!rNr~)rdrZcoords_reflected)rrbrlr(r)compute_coordinatestsz-_grid_sampler_2d..compute_coordinatescs||}||Sr5r()rdrZ coords_un)rmrfr(r)compute_source_indexs z._grid_sampler_2d..compute_source_indexr!r )r/ysrLc s,td|kt|ktd|k|kSrir>r,)r/ro)iHiWr(r)in_bounds_conds$z(_grid_sampler_2d..in_bounds_condr$)r/rowsrLcsN||rndtfdd|jtjd|jtjd|fDS)Nr c3s(|] }t|dVqdSrG)r>r`r-r)rr[rOoHoWr(r)rIsz1_grid_sampler_2d..clip..r)rVr7r>rB)r/rort)rrrcrsrurv)r[rOr)clips   z_grid_sampler_2d..clip)ixiyrLcs&|||\}}}||f|Sr5r()rxryr>Zidx_xZidx_yZw_)C_idxN_idxrrwr(r) get_summandsz%_grid_sampler_2d..get_summand).r).r rc3s |]\}}}|||VqdSr5r()r0rxryr>)r|r(r)rIsz#_grid_sampler_2d..cs |}|}||dSrWr()rxryr1rV)rmr|rqrrr(r)get_value_boundeds  z+_grid_sampler_2d..get_value_bounded)rerLcsF|d}d||d|d|f}t|S)Nr r!)rK)reZiy_ofscs)r}ix_nwiy_nwtxr(r) get_coeffs    z#_grid_sampler_2d..get_coeffc3s|]}|VqdSr5r()r0re)rr(r)rIrrr)r>rsrr!rr-r'r)r%rrjr%r4rPrVrOrK)rr]rarbrrcrnrQZtwor1rVrxryZix_neZiy_neZix_swZiy_swZix_seZiy_seZw_nwZw_neZw_swZw_seZ ix_nearestZ iy_nearesttyrHr()rrzrr{rcrrrwrmrr|r}rqrrrsrarrrurvrbrlrrfr)_grid_sampler_2dFsx      "               r)rr]rarbrrLcCst|||||dS)N)r]rarbr)r)rr]rarbrr(r(r)grid_sampler_2ds rcs`tdkodkfddtddkfddjddS)Nr!r csddS)Nzmatrix @ vector expected, got rrpr(rlrr(r)rqrrzmv..rcs*dddddddS)Nzsize mismatch, got input (rr1r z), vec (rryr(rr(r)rq rrrp)r>rsrKrrrr(rr)rs  rcCsd|dur4|d|d}d|||t|}nd||t|}|durZ||}t||SrW)rZ logsigmoidr)rlrrZ pos_weightrZ log_weightrr(r(r) binary_cross_entropy_with_logitssr)tensor1tensor2is_outrLc s|j|jkr||fn||f\}}ddlm|jdkr@|jdksDdS|jrR|sRdS|jdkr`dS|dkrtdS|j}|}dg}t|ddD]}|||dqt fd d t |t t||DS) Nrrrxr!FTr rNc3s(|] \}}}|dkp||kVqdSrAr()r0rrrrr(r)rI7szshould_fold..) rr r rrrrr3rrrcr) rrrt1t2Zt1_shapeZ t1_strideZexpected_striderr(rr) should_folds(     r)Z pass_is_out)rcCs|}|}|dkr |dks$J|dkr@|dkr@t||S|dkr\|dkr\t||S|dkr|dkrttt|d|dS|dkr|dkrt||St|||r||k}|r|jn|}|s|n|dkr| n|}|j }t |dd} t t j| } |dk} | r(| |j d|| |d} | rhtjj| || } |rd| jS| Stjj| || Sn|dkrz|dkrz|dkr|dnd}|d}|j dd}|dkr|dn|d}|dkr|dnd}g}t|dD]}|||q |dkr|dkr|d|dkr|ddkrp|jrpt|d|S|ddkr|jrt||dSt t||}|||g}t|}|||||}|dk}|r||g}||||d}n |||g}|||||}|} |dkr8| ||dkrL| ||rh||d| S||| Sntddd dS) Nrr r!rNrrxFcSsdS)Nz/both arguments to matmul need to be at least 1Dr(r(r(r(r)rqrrzmatmul..)rKr>dotrrerrPrrrrrrrYrrrrrt _unsafe_viewr>rrOrrZbroadcast_shapesrzr'bmmr-rs)rrrZ dim_tensor1Z dim_tensor2rrrZsizes_1Z output_shapeZ folded_dim1Z t2_is_matrixZ t1_foldedr<r m1Z batch_tensor1m2rZ batch_tensor2rZexpand_batch_portionZtensor1_expand_sizeZexpand_batch_productZtensor1_expandedZ vector_rhsZtensor2_expand_sizeZtensor2_expandedr(r(r)r?s                  r)rrarrrrLcs j\}}t|d||}t|d||}tjtjjd\}}tj|djdj |d} tj|djdj |d} t || |} t || |} | d} | } | }| | dd}| |  dd}| tj} | tj}|d||d|d f}| d| | d| d ft|t|}d \jtjkrttt|fd d Dfd d |D}fddfddtfdd|D}jtjkr؈dusJt||}ntddt||D}t}|j|d}|S)Nrr rr$rrNrr^r!rcs.g|]&}|d>t|dtjqSr rr>rr7int16r0r>)weights_precision_xr(r)r3sz.upsample_bicubic2d_default..cs.g|]&}|d>t|dtjqSrrr)weights_precision_yr(r)r3scs<t|dd}t|dd}tdd||g}|Sr)r>rrtr)ror/Zy_idxZx_idxr6)in_hin_wrr(r) load_boundedsz0upsample_bicubic2d_default..load_boundedcsTtfddD}jtjkr<dus0Jt|Stddt|DS)Nc3s|]}|VqdSr5r()r0Zx_ofs)rrVr(r)rIrrzCupsample_bicubic2d_default..get_x_interp..css|]\}}||VqdSr5r(rJr(r(r)rIrr)rVrr>rr&r%rc)rVZsrc_x)rixs_ofsrr weights_x)rVr) get_x_interps    z0upsample_bicubic2d_default..get_x_interpc3s|]}|VqdSr5r()r0Zy_ofs)rr(r)rIrrz-upsample_bicubic2d_default..css|]\}}||VqdSr5r(rJr(r(r)rIrrr)rrrArBrr2r>r)r%r7r!rPrjrrBrGrrr)rVr&r%rcrr>)rrarrrrQZh_scale_factorZw_scale_factorrrrWZx_floatZy_floatr1rVZyscaler0Ziys_ofsZ weights_yZsrc_yrrr() rrrrrrrrrr)upsample_bicubic2d_defaultsR           r)rrarrrLcCstt|t|dkdd|durd|dus4Jttttftddt|jdd|D}|rl|nd\}}t|||||S)Nr cSsdS)Nz:Must specify exactly one of output_size and scale_factors.r(r(r(r(r)rqrrz(upsample_bicubic2d_vec..css"|]\}}tt||VqdSr5)rr )r0r>rer(r(r)rI#sz)upsample_bicubic2d_vec..r!r) r>rsr+r rVr!rcrr)rrarrrrr(r(r)upsample_bicubic2d_vecs   r)rrFrLcsfdd}t||S)Ncs4tj| ||jd}|d|d|S)Nr$r )r>r)r%rrZmiddlerZdim_idxrr(r)r02sz_reflection_pad..idx_reflection_or_replication_padrrFr0r(rr)_reflection_pad,s  rcsfdd}t||S)Ncs*tj| ||jd}t|d|dS)Nr$rr )r>r)r%rrrr(r)r0Csz_replication_pad..idxrrr(rr)_replication_pad=s  r)rrFidx_fnrLc stdt|ddfvfdd|j d}|}fddtD}fddtD}|}tD]>}dg|} |||||||| ||<t|| }qt |} |j | d}|S) Nr!r cs ddddddS)NZreflection_padz d requires r zD or r!zD inputr(r(rpr(r)rqVrrz0_reflection_or_replication_pad..cs g|]}dd|qSrXr(rrKrFr(r)r3[rrz2_reflection_or_replication_pad..cs$g|]}dd|dqSrXr(rrr(r)r3\rrr) rr>rsrKrrOrtrrArr>) rrFrZ inp_shapeZnc_dim padding_left padding_rightrrr0rr(rr)rNs"       rcshtddd|j dDfddtDfddtDg}t|jD]:}dg|j}d||<|tj|j||jd|qd|d | d d d  fd dtD fd dtD} fd dtD}fddtD t t j  fddtD}t |d} fdd} tjddtDD]} | tdgkrqg} g} tD]}| |dkrڈ|} |}nb| |dkr||} |d|f}n8| |dkr<||} ||||df}| || |q| | | | } q| S)Nr!cSsg|] }|dqSrr()r0rPr(r(r)r3qrrz,_reflection_pad_backward..cs g|]}dd|qSrXr(rrr(r)r3srrcs$g|]}dd|dqSrXr(rrr(r)r3trrr rNr$cSs|\}}}t||k||kSr5rp) index_rangerZlbZubr(r(r)index_range_conditions z7_reflection_pad_backward..index_range_conditioncsg|]}||qSr(r(rrxyzr(r)r3rrcsg|]}||qSr(r(rrr(r)r3rrcs(g|] }d|||qSrvr(r)dhwrrr(r)r3rrcs.g|]&}|d|||fqSrr(r)centerrrrr(r)r3scsg|]}|qSr(r(r)rrange_cr(r)r3rrrcsttD]2}||d||dk}t|tr|r|Sqttjfdd|D}t||d}||S)Nr!r csg|] }|qSr(r()r0r)rr(r)r3rrz@_reflection_pad_backward..accumulate..r)rOr.r+rHrrtr,r.)rr index_rangesrZupper_less_than_lowerrOg)rrKrcrr(r)rys  z,_reflection_pad_backward..accumulatecSsg|] }gdqS))rNrr r(rr(r(r)r3rrr)rrrOrrr>r)r%r-rHrrtr,r. itertoolsrrV)rcr1rFr/rZ view_shapeZ left_reflectZ right_reflectrOrryareaZoutsrrrr() rrrrKrcrrFrrrrr)_reflection_pad_backwardjsT  $  " rrzr}rcCs(tj|||d}tj|||d}||fS)Nr)r>aminr)rlrKrrrr(r(r)aminmaxsrcCs"tjtt|d||||dS)Nrr)rtrr>r`isnan)rlrKrrr(r(r)nansumsrrrr%r)rrrr%r)c Cstjjd|d||||dS)Nrr rrtr)Z start_steprr(r(r)arange_defaults rrrrrr%r)c Cstjj||d||||dS)Nr rrrr(r(r) arange_starts rcOsddlm}||i|S)Nr)out_dtype_dense)!torch._higher_order_ops.out_dtyper)rCrDrr(r(r)out_dtype_decomps r)rrrmarginrrrLc sttjdjdt|dkp:|dkddtjdkoXdkfddtjdko~kfdddurttjdkokfdddtjdd }||}| d}|dkr|n||}dur$|}tj j d }t |k|d}|t jjkr\|S|t jjkr|||jdS|jdd SdS) Nrr r!cSsdS)Nz only p == 1 and p == 2 supportedr(r(r(r(r)rqrrz#multi_margin_loss..cs djSNzMExpected non-empty vector or matrix with optional 0-dim batch size, but got: rWr(rfr(r)rq rrcsddjS)Nz#inconsistent target size, expected rrrWr()nframerr(r)rqrrcsddjS)Nz#inconsistent weight size, expected rrrWr()rKrr(r)rqrrrzr$rp)r> atleast_2dZ atleast_1drrsrrrPr=rr)r%r`rr&rmrr'r) rrrrrrurar0r()rKrrrrr)multi_margin_losssB             r is_targetc s|j|jt|}t|}|jd}ttdko@|dkfddttdkodkfddtj||jd}|dk}tjt|||dd d }||k}t||d}tj |d|d } t||d} tj || j dd kdd } d | j j dd |} | d} | |} t| d| } |tjjkrJ| jdd } n$|tjjkrb| } n | jdd } | |j} | | fS)Nr r!rcs dSrr(r()orig_input_shaper(r)rq6rrz0multilabel_margin_loss_forward..csddS)Nzinconsistent target size: z for input of size: r(r(rZorig_target_shaper(r)rq:rrr$rNTrrzrpr^)rrN)rr>rrsrr)r%rr`r=anyrPTrrr&rmrrr'r7rr) rrrrKr0Zis_endZend_idxZ target_maskZtidx0rZtidx1rrar(rr)multilabel_margin_loss_forward's@        r) attn_maskre)querykeyrm dropout_p is_causalrrerLc sttfddtdko@dko@dkfddtdkfddtjdjdkojdjdkddtjj||d|d \}}|d d d dj tj d d d d d}||fS)Ncs djS)Nz-query must be FP32, FP64, BF16, FP16 but got rr()rr(r)rqsrrz.rcs"dddS)Nz,q, k, v must be a 4 dimensional tensor, got rrpr()rrrmr(r)rqwrrrcs dS)Nz&dropout probability must be zero, got r(r()rr(r)rqzrrrxcSsdS)Nz&q, k, v should have the same head sizer(r(r(r(r)rq~rr)rrrZ dropout_maskrer!rr r) r>rsrrKrrtZ"_scaled_dot_product_attention_mathrrdr>r) rrrmrrrrer<Zattnr()rrrrmr)*scaled_dot_product_flash_attention_for_cpufs@  "& "rcst|fdd}|S)Ncs|i|}|d|Sri)r)rCrDr outplace_opr(r) inplace_opsz$register_inplace..inplace_opr)Zaten_oprrr(rr)register_inplacesrcCsx|s |s t|}t|}t||}t|tjr@|dkrH||}|dkrT|St|tjrh|dkrp||}||S)Nr r)rrr!r>rr.numbersNumber)rlZbatch1Zbatch2r[rdrr(r(r)baddbmms rcCstj||ddS)NrjrCrD)rlrr(r(r) floor_dividesrcCsttj|jdSrW)rHrrYrr)rr(r(r) sym_numelsrrr)rlrrrLcCs2|durtjj|g|dStjj|g||dSdS)Nrr)rtrZ dim_IntListZ IntList_out)rlrrr(r(r) sum_defaultsrrlrKcCsFt|tjs|S|dur2tj|tt|Stj||gSdSr5) r.r>rrtredimsrrOrKrr(r(r)squeeze_defaults  rcs`tfddtt|jD}|jtjkr2tjnd}|jd|d|d}||| |j|fS)Nc3s|]}|kr|VqdSr5r(rrpr(r)rIrrz)_weight_norm_interface..r!T)rr) rVrOrrrr>rrrr7)r6rrKZkeep_dimZ norm_dtyperr(rpr)_weight_norm_interfaces r assume_uniqueinvertcCst|tjstj||jd}t|tjsD|r8t||St||S|dt|dkrlt |||dSt ||||dSdS)Nr$g$@g(\?rr) r.r>rrr%nerrr isin_default isin_sorting)elements test_elementsrrr(r(r)isins    r)r')rlr'rLcCsP|dur"tj|tj|jd}ntj||tj|jd}||k|j}|S)NrA)r'rr%)r>Zrandrrr%r7r)rlr'Zraw_prr(r(r) bernoullisrrcCsl|dkrtj|tjdS|jd|j}||}ttd|j dd}||kj |d}|rh|S|S)NrrrrNr rp) rr> empty_liker+rrr-rVrOr)rrrZexpanded_elem_shaper1rKrr(r(r)rs  rcCs|}|}|rt||g}tj|dd\}}|dd|ddk} t| ddgd} |rj| } t| } | d|| } | d|St|\} } t | |} t | | k| d}| ||k}|r|n|}| |j SdS)NT)Zstabler rNrF) rr>rsortrZ logical_notrrrZ searchsortedr`rr)rrrrZ elements_flatZtest_elements_flatZ all_elementsZsorted_elementsZ sorted_orderZduplicate_maskr1Zsorted_test_elementsrQr0Ztest_idxcmpr(r(r)r#s$   rcCs|d}||SrM)r)rlr2Z flattenedr(r(r)take@s rcCs2|durtj}|tjkr t|}tj||j|dSr~)r>rZpreserve_formatrrtresizer)rlrrr(r(r) resize_asGs  r)FF)r)N)r!)rNNr )rNNr )N)rNFF)N)rN)r)r)r)r)r r )r r F)r r )r)N)Fr)rr^N)rr N)F)N)N)NN)NN)NNN)NNN)F)F)F)F)F)F)N)NN)NNN)N)F)rrFT)rrF)NN)N)NF)rF)r r )N)r)N(rHrrrYrcollections.abcrenumrrrrrtypingrr r r r r>Ztorch._meta_registrationsZ torch._primsr.rZtorch._prims_commonZ _prims_commonrAZtorch.nn.functionalrTrrr rrZ torch._decomprrrrrrrrZtorch._prims_common.wrappersrrrrZ torch.utilsrr@Ztorch.utils._pytreerrZ DispatchKeyrrstr__annotations__Z_opsrrtrrr+rJrZcompute_only_pw_cast_for_opmathZpw_cast_for_opmathZ"pw_cast_for_opmath_non_tensor_argsr2Zpw_cast_for_int_to_realr!rRrYrZrbrrjr;ZScalarrnrvrrrrrrrrrrrrrrVrrrrrrr&rmrZ _safe_softmaxrrrrrSrrrrrrrrrrrrrrr slicerr"r r4r8r;r?r@rCror}rrrrZpy_implZCompositeImplicitAutogradZAutogradrrrrrrzrrrrrrrrrrrrrrrrrrrrr rrZ unsafe_chunkrrrZno_statsrrrr$r%r&Z_fused_dropoutr(r/r%rrZliftZ lift_freshr1r4r6r8r7r9r>r@Z_adaptive_avg_pool2drXrZrcrhrnrormrrrrrr GeneratorrrrrrrrrZ_upsample_nearest_exact1drrrrrrrrrrrrrrrrZrnn_tanhrrZrnn_relurrrrrrrrZlstmr r r r Zgrurrrrrrrrrrrr!r&r)rr7r8rrr{r.r<r?r@rArDrFrGrKr%rNrWrYr^r_r`rrrrrrZupsample_bicubic2drrZreflection_pad1dZreflection_pad2dZreflection_pad3drZreplication_pad1dZreplication_pad2dZreplication_pad3drrZreflection_pad1d_backwardZreflection_pad2d_backwardZreflection_pad3d_backwardrrrr)r+rrrrrrrZ+_scaled_dot_product_flash_attention_for_cpurrrrrrrrerKrrrrrrrrZaddbmm_ZaddbmmZaddmm_Zaddmv_Zbaddbmm_Zfill_Zgelu_rZ hardswish_Z hardtanh_ZhardtanhZ hardsigmoid___iand____and__ __ilshift__ __lshift__rvrwZ index_reduce_Z index_reduce__ior____or__ __irshift__ __rshift____ixor____xor__Z leaky_relu_Z leaky_reluZlogit_ZlogitZrelu_rZrenorm_ZrenormZround_r4Zscatter_rZ scatter_add_Z scatter_addZscatter_reduce_Zscatter_reduceZsilu_r(r(r(r)s            *              *!    : '    P`      %   "    "    W  R R  #           ,   h %  $g -7 (   "          $     $                .  2 )     @ 2               I      5           $'      w S    ,   ,     W     ,   <  C    "