o i -@s8ddlZddlmZmZddlmZddlmZddlm Z ddl m Z ddl Z ddl mZddl mZmZmZddlmZmZmZmZdd lmZdd lmZmZdd l mZmZmZm Z m!Z!m"Z"m#Z#m$Z$m%Z%m&Z&m'Z'dd l(m)Z)m*Z*m+Z+m,Z,m-Z-dd l.m/Z/m0Z0ddl1m2Z3ddl4m5Z5m6Z6ddl7m8Z9e dZ:e dZ;e j?dddZ@eAd\ZBZCZDddZEddZFdeee;e:fgee;e:fffddZGde!fdd ZHd!d"ZId#d$ZJeGe=jKe=jLge-ddde jMd%d%fd&d'ZNeGe=jOjPe=jOjQge-d(d)ZReGe=jSjPe=jSjQge-d*d+d,d-ZSeGe=jTe-d.d/ZTeGe=jUjPe=jUjQe=jVjPe=jVjQge-d0d1d2d3ZWeGe=jXjPe=jXjQge-d4d5ZXd6d7ZYdvd8ed9eZe[d:e\fd;d<Z]eGe=j^jPe=j^jQge-d=d>Z_dZ`d9eZe[fd?d@ZaeGe=jbjPe=jbjQge-dAdBZceGe=jdjeddCdDdEZfeGe=jdjPe jgddddFdGdHZheGe=jijPe=jijQge-e jgddddFdIdJZjeGe=jijke=jijlge-e jgddddFdKdLZmeGe=jnjPe=jnjQge-dddddFdMdNZoeGe=jpjPe=jpjQge-d8ed9eZe[dOe[dPe[fdQdRZqeGe=jrjPdvdSdTZsdUdVZteGe=jujPdWdXZveGe=jw   dwdYedZed[ed\edBd]exdBd^e jydBf d_d`ZzeGe=j{ dxdaedbedced^e jydBfdddeZ|eGe=j}dfdfddgdYedaedbedced^e jydBf dhdiZ~eGe=j    % f *dydje jdke jd\edBdledBd^e jydBdme\dne[doe[dpe[fdqdrZeGe=jjPdsdtd8ed9e[duedve jdwexdxe\defdydzZeGe=jjPdsdtd8ed9e[duedve jdwexdxe\defd{d|Ze-eGe=jjPd}d~ZeGe=jjPddddd%dddedwexdedBd1edBdedBde[de\defddZeGe=jjPe=jjge-ddZeGe=jjdvddZeGe=jjPe=jjge-ddZeGe=jjdvddZeGe=jjPddZeGe=jjQddZeGe=jjPddZeGe=jjddZeGe=jjPddZeGe=jjPddddddddZeGe=jjPdzddZeGe=jjPdwddZeGe=jjPdzddZeGe=jjPddZeGe=jjddZd8edexfddZd8ededexfddZ sd{dedexde\fddZd|dedexdexfddZdedede\dexfddZ d}dexdedYedexfddZdexfddZeGe=jjPe=jjge-ddd~dedexde\fddńZeGe=jjPe=jjQge-dYedefddDŽZeGe=jge-dddYefddɄZdedefdd̄ZeGe=je-d8edede\defddτZeGe=je-dvd8edede\defddфZeGe=je-dvd8ede\defddӄZeGe=je-dvd8ede\defddՄZeGe=jjPddede\de\fdd؄ZeGe=jjPe=jjQge-dYededefddۄZeGe=jjPdvdede\fdd݄ZeGe=jjPe=jjQge-dddd%d%dd8ede\de\deeeeffddZeGe=jjPe=jjQge-d%ddededede\def ddZeGe=jjPe=jjQge-ddddsddede\deeeeffddZeGe=jjPe=jjQge-ddddsd%ddede\de\deeeeffddZeGe=jjPe=jjQge-dsd%ddededede\de\def ddZeGe=jǃe-ddd s sddedede\de\deeeeff ddZdexdee\e\ffddZeGe=jjPe=jjQge-ddddedexdeeeffddZeGe=jjPe=jjge-dddd߃dedeeeeeffddZeGe=jjP % s ddede\de\dexdBfd d Zϐd ed edeeZe[eZe[ffd dZАd ed edexdBdeeeffddZdYedede\fddZeGe=jӃdsd%ddddddedede\de\dedBdedBdedBdedBdeeeeeffddZeGe=jjPe=jjQgdsd%dddedede\de\de\dedBdefddZeGe=jփe-dddsd s % %dd8edede\de\de\deeeff d d!ZeGe=jjPd"d#ZeGe=jڃe- s %ddYededede\de\def d$d%Zڐd&d'Zېd(d)ZeGe=j݃e-d*d+ZeGe=j߃e-d,d-Zd.d/ZeGe=je-d0d1d2ZeGe=je-d0d3d4Zd5d6ZeGe=je-d7d8ZeGe=je-d9d:ZeGe=jjPd;d<ZeGe=jjPe=jje=jjPe=jjge-d0d=d>Zd?d@ZeGe=je-dAdBZeGe=je-dCdDZeGe=jjPe=jje=jjPe=jjge-d0dEdFZeGe=je-dd8edHedefdIdJZeGe=je-dKed8edHedLedef dMdNZeGe=jjPe=jjQge-dsddfdfdOdPdQZeGe=jjPe=jjQge-ddCdRdSZeGe=jjddUdVZeGe=jjddWdXZeGe=jjPe=jjQge-dxdYdZZeGe=jjP % %dd[d\Z eGe=j e-dsddxd^e jydBfd]d^Z d_d`Z ddbdcZ dxdde jdZe jdeeZe[e[BdfeZe[e[BdgeZe[e[Bdhe\die[djeZe[e[BdBfdkdlZdmdnZeGe=jjPdde jdZe jd\e jdBdoe jdBdpe jdBdqe\dredsefdtduZeGe=jjPdde jdZe jd\e jdeeZe[dfeZe[dgeZe[dhe\djeZe[die[fdvdwZe jj re j>?dxddZeGe j<jjjPdydzZeGe j<jjjPd{d|Ze jj re j>?d}ddZeGe j<jjd~dZ e j>?dddZ!eGe j<j"j#jPeGe j<j"j$jPeGe j<j"j$j%ddZ&eGe j<j"j#j'eGe j<j"j#j(ddZ)eGe j<j"j*jPeGe j<j"j*j%ddZ+eGe j<j"j*j'eGe j<j"j*j(ddZ,eGe j<j"j-jPeGe j<j"j.jPddZ/e j>?dddZ0eGe j<j1j2   %dddZ3eGe j<j1j4ddZ5ddZ6eGe=j7jP  % s dddZ8ddZ9eGe=j:jPddZ;eGe=j<e-  % s dddZ=eGe=j>e-d0ddZ?eGe=j@jPddZAeGe=jBjPddZCeGe=jDjPddZEeGe=jFe-d0ddZGdedexfddZHeGe=jIe-dd1ddZJeGe=jKe-d0ddZLeGe=jMe-dd1ddZNeGe=jOe-d0ddZPeGe=jQjdxddZReGe=jSjPe=jSjQge-ddZTeGe=jUjPe=jUjQge-d*dde[fddZUeGe je[d?e[d@e[dAe[dBe[dCe[dDe[dEe[dFe[dexdGe\f,dHdIZ̐dJdKZdYeded4e[d5e[d6e[d7e[d8e[d9e[d:e[d;e[djPddZ>eGe=j?jPddZ?dd Z@d d ZAeGe=jBjPe=jCjPgdd d ZDeGe=jEjPdddZEeGe=jFjPdddZGeGe=jHe- dddZIeGe=jJjPe=jJjge-d0d1dddZKe jLZMddZNeGe=jOjPddZOeGe=jPjPddZPeGe=jQjPddZReGe=jSjPddZSeGe=jTje=jTjUge-d%d%d d!d"ZVeGe=jWge-dd$d%ZXeGe=jYjPe=jZjPg  dzd&d'Z[eGe=j\jPg  dzd(d)Z]eGe=j^jPd*d+Z^eGe=j_jPe=j_jQge-dwd,d-Z_eGe jZleGe=jme-d%d%d?d@dAZneGe=joe-dBe[d8edefdCdDZpeGe=jqd8efdEdFZreGe=jse-dsdd8edefdGdHZseGe=jte-d8edefdIdJZtdKdLZu     %ddMedNede jdBde jdBdOedBd\edBde jdBd^e jydBde\fdPdQZveGe=jwe-   dwdMedNedOedBd\edBd^e jydBdef dRdSZxeGe=jyg     %ddMe jdNe jde jde jdOe jdBd\e jdBde jdBd^e jydBde\fdTdUZzeGe=j{e-dVed9e[dWe\defdXdYZ|eGe=j}e-ddZd[Z~eGe=je- * % %ddZed1ed\e[d]e\d^e\def d_d`ZeGe=jjP dd0edeZedaeZe[dbefdcddZdedfZeGe=jjP  s s ddgdhZdidjZee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jee=jdkdlZeGe=je-dmdnZeGe=je-dfdodpdqZeGe=je-dfdodrdsZee=jZee=jZee=jZddl.Z ddlZ ddlZ dtduZedS(N)CallableSequence)Enum)wraps)TypeVar) ParamSpec)SymBoolSymFloatTensor)_add_op_to_registry_convert_out_paramsglobal_decomposition_table meta_table) OpOverload)_prim_elementwise_meta$ELEMENTWISE_PRIM_TYPE_PROMOTION_KIND) BoolLikecorresponding_complex_dtypecorresponding_real_dtypeelementwise_dtypesELEMENTWISE_TYPE_PROMOTION_KIND FloatLikeIntLikemake_contiguous_strides_forNumbersuggest_memory_format TensorLike)_maybe_convert_to_dtype_maybe_resize_out_resize_output_check_safe_copy_out out_wrapper)_broadcast_shapes_maybe_broadcast)_config) ScalingType SwizzleType)_pytree_T_PatenIMPLMetacCs||d|SNabr0r0[/sda-disk/www/egybert/egybert_env/lib/python3.10/site-packages/torch/_meta_registrations.pyceil_div7r5cC||d||S)z$Rounds up x to nearest multiple of yr/r0xyr0r0r4round_up;sr;returncsfdd}|S)Ncs$tfdd}t|S)Ncstt|dSN)r ropfnr0r4registerDr6z0register_meta..wrapper..register)r pytree tree_map_)rArBr>r@r4wrapperAs  zregister_meta..wrapperr0)r?rEr0r>r4 register_meta@s  rFtype_promotioncs>tj|d|i\}fdd|D}t|}t|dtjiS)Ntype_promotion_kindcsg|]}t|qSr0)r.0r9 result_dtyper0r4 Vz$elementwise_meta..rG)utilsrr#rrDEFAULT)rGargs_r0rKr4elementwise_metaMs rScCs(tjtjtjtjtjtji}|||Sr=)torch complex32halfcfloatfloatcdoubledoubleget)dtype from_complexr0r0r4toRealValueTypeas  r^cs2ttg|RtkfdddS)NcddS)Nzoutput with shape z# doesn't match the broadcast shape r0r0broadcasted_shape self_shaper0r4nz)check_inplace_broadcast..)tupler"rT_check)rb args_shaper0r`r4check_inplace_broadcastjs  rhFc sttjrtdkddttjr$tdkddtddfDrMttdur>ntt fddnpRtttj s[Jt tt fddtt sqJtdkd dtj f|d ||d S) NrcSdSNz:linspace only supports 0-dimensional start and end tensorsr0r0r0r0r4rcz(meta_linspace_logspace..cSrirjr0r0r0r0r4rcrkcss|]}t|tVqdSr=) isinstancecomplex)rJargr0r0r4 sz)meta_linspace_logspace..cddS)Nzlinspace(): inferred dtype z& can't be safely cast to passed dtype r0r0)default_complex_dtyper\r0r4rcrdcs*dtjdtjdtjdS)Nz4received an invalid combination of arguments - got (, ))type__name__r0)endstartstepsr0r4rcscSri)Nz$number of steps must be non-negativer0r0r0r0r4rcrkmetar\layoutdevice pin_memory requires_grad)rlrTr rfdimanyrOrget_default_dtypeis_complex_dtyper\ _check_typerempty) rwrvrxbaser\r|r{r}r~r0)rqr\rvrwrxr4meta_linspace_logspacersH      rcsNtjtjkfddt|dkodk dd|jS)Nc djS)Nz2take(): Expected a long tensor for index, but got r\r0indexr0r4rc zmeta_take..rcSri)Nz*take(): tried to take from an empty tensorr0r0r0r0r4rcrk)rTrfr\long _check_indexnumel new_emptyshape)selfrr0rr4 meta_takes   rrcshj}j}t||kddtdkodkfddtjj}|S)NcSri)Nz=linalg.cross: inputs must have the same number of dimensions.r0r0r0r0r4rcrkzlinalg_cross..r-cs"dddS)Nzlinalg.cross: inputs dimension z must have length 3. Got  and sizer0rotherrr0r4rcs )ndimrTrfrr"rr)rrrx_dy_d out_shaper0rr4 linalg_crosss rcCs$t|dt|dtj|tjdS)Nzlinalg.matrix_exp memory_format)squareCheckInputscheckFloatingOrComplexrT empty_likecontiguous_formatrr0r0r4linalg_matrix_exps  rvaluesindicescCsVtj|j|j|jd}tj|j|jtjd}|dkr'|jdkr't||j||fS)Nr|r\r) rTrrr|r\int64rrmaybe_wrap_dim)rrrrr0r0r4 cummaxmins  rcCst||jtj|tjdSNr)rrrTrr)rrr0r0r4 logcumsumexps rcs|j}t|}||}tt|}ddt|D} |D]} d| | <qgg} } |D]} | | s6| | q*| | q*| | }t| } ||d| }|jfdddd||| d}||}dgt|j|d}| |}| d}||d<t|}tt|D] }|||||d <q|j |t j d d dt|D}d }|d }|dkr||d|||<||||9}|d 8}|dkst||D]}|d |||||<q|||||S) NcSg|]}dqSFr0rJrRr0r0r4rMrdz_exec_fft..Tc|Sr=r0r9 self_stridesr0r4rc z_exec_fft..keyreverserrr/rcSrrr0rr0r0r4rMrd)rlenlistrangeappendstridesortpermuterreshaperresize_rTr as_strided_storage_offset)outr out_sizesrforwardr signal_ndim batch_dims dim_permuteis_transformed_dimdleftright batch_endtmpinput batched_sizes batch_sizebatched_out_sizesi out_strides batch_numelr0rr4 _exec_fftsN         rrr exclude_lastcs<t|}||dt|t|jfddd|S)Ncrr=r0rrr0r4rc,rz_sort_dims..)r)rrrintr)rrr sorted_dimsr0rr4 _sort_dims(s  rcCsHt|jj|s |St||}||}t|||||dS)Nr) rTrfr\ is_complexclonerrrr)rr normalizationrrrr0r0r4 meta_fft_c2c3s  rcCs8t|tkst|dkr|ddkr|ddkrdSdS)Nrr/FT)rcufft_max_ndimrr0r0r4use_optimized_cufft_pathBs0rcst|jjt|}t|}|d}||dd}t|}|||<|r+|||<t|dks7t|dkr|j|t |jd} |} t|dkrXt |rXt | | ||ddngt |dkr`|n|} t | | | |gddt |dkr}|j|t |jd} |dd} | r| | } } | | jfd d dd ttt | } | t | | d}t | | ||dd| dt | | } | s|s| |||kr| j|tjd | } | S|j|t |jdS) Nrrr/cudaxpurTrcrr=r0rstridesr0r4rcprzmeta_fft_r2c..rr)rTrfr\is_floating_pointrr device_hintrrOrrrrrrminrrr)rrronesided input_sizesrlast_dimlast_dim_halfsizeonesided_sizesoutputworking_tensor target_sizesrmax_dims last_dimsr0rr4 meta_fft_r2cIsX          r) generatorcCst|t|gSr=)rrTSize)nrrr0r0r4 meta_randpermsrr\r{r|r}cCtj|||||dSNrrTr)rr\r{r|r}r0r0r4meta_randperm_defaults rcs2dtkfddtj|||||dS)Nrcr_Nz:random_ expects 'from' to be less than 'to', but got from=z >= to=r0r0highlowr0r4rcrdzmeta_randint..rrTrfr)rrr\r{r|r}r0rr4 meta_randints   rcs.tkfddtj|||||dS)Ncr_rr0r0rr0r4rcrdz"meta_randint_low..rr)rrrr\r{r|r}r0rr4meta_randint_lows   rcCrrr)rr\r{r|r}r0r0r4meta_rand_defaults r rlastdimc Cs*t|jjt|dkrZt|}|||d<|j|t|jd}t |r5t ||j tj d||ddSt |dkrGt||ddd|}n|j tj d}t ||||dgddS|}t |dkrv|dd}t|||dd}|dd}t|}|||d<|j|t|jd} t | |||ddS) NrrrrFrr/r)rTrfr\rrrrrr^rrrrrr) rrrr rrtemprc2c_dimsrr0r0r4 meta_fft_c2rs4          r cCsfddlm}||st|dkrtdt|tr1|||}||kr1t j |||S)Nr)free_unbacked_symbolsr/zQmore than one element of the written-to tensor refers to a single memory location) %torch.fx.experimental.symbolic_shapesrrT_debug_has_internal_overlap RuntimeErrorrlr torr* expand_copydefault)rsrc non_blockingr intermediater0r0r4 meta_copy_s   rcCsXt|}t|}||krdn||||}||d|||||fSr.)rrrrinsert)tensorr result_sizesresult_strides new_strider0r0r4inferUnsqueezeGeometry s     rcCs0t||d}t||\}}||||Sr.)rrrr)rrg_sizes g_stridesr0r0r4meta_unsqueeze_s r!rweight_metabias_activation_opt out_dtypec Cst|j}|dur|d|dksJd|d|ddks%J|d|d<t|jdks7Jdd|df}|durQ|jtjkrM|tjksQJd|j||dur[|jn|d ||}|S) Nrzoutput size mismatchr/rrz*we can only handle the squashed input case9out_dtype is only supported for i8i8->i32 linear operatorr) rrrrr\rTint8int32r as_strided) rr"r#r$r%r& output_sizestransposed_stridesrr0r0r4meta_sparse_structured_linear!s$ r-mat1 mat1_metamat2cCst|jdks Jt|jdksJt|jdksJ|d|ddks)J|d|dg}|durF|jtjkrB|tjksFJd|j||durP|jn|d}|S)Nrr/rr'rrrrr\rTr(r)r)r.r/r0r&r+rr0r0r4meta_sparse_structured_mmCsr2r/)alphabetar&c Cst|jdks Jdt|jdksJt|jdksJt|jdks&J|d|dks4Jd|d|ddksBJ|d|dg}|dur_|jtjkr[|tjks_Jd|j||duri|jn|d}|S)Nr/zEonly input broadcasted to columns of mat1 * mat2 product is supportedrrr'rr1) rr.r/r0r3r4r&r+rr0r0r4meta_sparse_structured_addmm\s( r5 compressed_Adense_Br3transpose_resultalg_idsplit_k split_k_modec Cs |jtjtjtjtjtjhvsJd|j|jksJdt|jdks(Jd|jtjtjfv} | r;| r;Jd| d} | d} |durR| | dksRJ|durt| rd|tjtjtj tjhvstJd|jd |jd |d |rz| | fn| | f} |j | |d S) Nz;_cslt_sparse_mm only supports fp16, bf16, int8, and fp8e4m3zinputs must have the same dtyperz'_cslt_sparse_mm only supports 2d inputsz.dense input must be transposed for 8bit dtypesr/rzout_dtype is not supported for z x z -> z matmul!r) r\rTfloat32float16bfloat16r( float8_e4m3fnrr is_contiguousrr)r) r6r7r$r3r&r8r9r:r;is_8bit_input_typerm output_shaper0r0r4meta__cslt_sparse_mms>    rDT) include_selfrsourcereducerEcCtj|tjdSrrTrrrrrrFrGrEr0r0r4meta_index_reduces rKcC|Sr=r0rJr0r0r4meta_index_reduce_s rMcCs.t|}|dkr|||<||SNr)rrrrr)rrr result_sizer0r0r4meta_index_selects    rP)lengthsroffsetsaxisunsafeinitialdatarQrRrSrTc sf|durtdfdd}|dur||jS|dur/|jdd|jddf} || Std)Nz?segment_reduce(): indices based reduction is not supported yet.cs(tj|jddjdtjdS)Nr/ryr\r|r)rTrrr\r) lengths_shaperSrVr0r4segment_reduce_lengths_tensors z:meta_segment_reduce..segment_reduce_lengths_tensorrr/zt||kdS||||ddS)Nr)_constrain_range_for_sizerr) rrrTrfrlr rrrtr)rrrrr0r0r4sym_constrain_range_for_sizeSs  rcCrr)r*rrr0r0r4'functional_sym_constrain_range_for_sizegrrcCs|Sr=r0)rvryrr0r0r4functional_assert_async_metamrxrf_namecCsX|dks J|d|d|dks*J|d|dd|dddS)Nrz3: The input tensor must have at least 2 dimensions.rz5: A must be batches of square matrices, but they are  by matrices)rr)rrr0r0r4rss  rAnamecstjjkfddtjjkfddtddkfddtddkfdddS)NcdjdjdS)Nz:Expected b and A to be on the same device, but found b on z and A on instead.r~r0rrr0r4rc z(linearSolveCheckInputs..cr)Nz=Expected b and A to have the same dtype, but found b of type z and A of type rrr0rr0r4rcrrrcsdddddS)Nz3A must be batches of square matrices, but they are rrrrrr0rr0r4rcs c s:ddddddddd S)NzIncompatible matrix sizes for z: each A matrix is rrz but each b matrix is rrr0rrrr0r4rcs)rTrfr|r\r)rrrr0rr4linearSolveCheckInputss      rtallow_low_precision_dtypescs^|jt|p |fdd|s-ttjtjtjtjfvfdddSdS)NcdS)Nz<: Expected a floating point or complex tensor as input. Got r0r0r\rr0r4rcz(checkFloatingOrComplex..cr)Nz*: Low precision dtypes not supported. Got r0r0rr0r4rcr) r\rTrfrrrXrZrWrY)rrrr0rr4rs  rarg_namecs"t|dkfdddS)NrcddS)Nz: The input tensor z! must have at least 2 dimensions.r0r0rrr0r4rcrdzcheckIsMatrix..)rTrfr)rrrr0rr4 checkIsMatrixs  rBrcsZtttrddkn ddkfdddS)Nrrc sHdrdndddddddddd S) Nz2: Incompatible shapes of A and B for the equation zAX = BzXA = Bz (rr9rrrsrr0rrrrr0r4rcs z#checkInputsSolver..)rrrTrfr)rrrrr0rr4checkInputsSolvers  *rresultfn_name result_namecs&tjjkfdddS)Nc s$dddjdj S)Nz : Expected z5 and input tensors to be on the same device, but got z on z and input on r~r0rrrrr0r4rcs z!checkSameDevice..)rTrfr|)rrrrr0rr4checkSameDevices rUPLOcs8}ttdko|dkp|dkfdddS)Nr/ULc dS)Nz1Expected UPLO argument to be 'L' or 'U', but got r0r0rr0r4rc zcheckUplo..)upperrTrfr)rUPLO_uppercaser0rr4 checkUplos  r eigenvalues eigenvectorsr compute_vcCspt|dt|t|j}|r ||}||t|ddn|dg}||j|t|j d}||fS)Nz linalg.eighF row_majorrr) rrrrrrrpopr^r\)rrrrvecsvalsr0r0r4meta__linalg_eighs    rcCs@t|dt|jr|jnt|j}|j|jdd|dS)Nzlinalg.eigvalsrr)rrOrr\rrr)r complex_dtyper0r0r4meta__linalg_eigvalss   rcCs|t|dt|jr|jnt|j}|j|jdd|d}|j|j|d}t|dk}||jt |j|d||fS)Nz linalg.eigrrrr) rrOrr\rrrrrr)rrrvectorsis_cudar0r0r4meta_linalg_eigs    rrcCs|jjtjdddS)Nrrr)mTrrTr transpose)rr0r0r4cloneBatchedColumnMajorrrcCst|Sr=)r)rrrr0r0r4_cholesky_solve_helpersrcsPtjdkfddtjdkfddtd\}}t|||S)NrcdjdS)Nz-b should have at least 2 dimensions, but has  dimensions insteadrr0rr0r4rc rz cholesky_solve..cr)Nz-u should have at least 2 dimensions, but has rrr0rr0r4rc$rcholesky_solve)rTrfr!_linalg_broadcast_batch_dims_namer)rrrself_broadcasted A_broadcastedr0rr4rs   rcCs.|dkrtj|tjdSt|dt|S)Nrrcholesky)rrTrlegacy_contiguous_formatrrrrr0r0r4r,s  rcCst|dt|S)Ncholesky_inverse)rrrr0r0r4r5s r check_errorscCsft|dt|d|j}t|}t|d}||}||||j|d|dtjd}||fS)Nzlinalg.choleskyFrrr) rrrrrrrrTr))rrrA_shaper L_stridesrinfosr0r0r4linalg_cholesky_ex=s     rtaucstjdkddtddkddtddkddtjjdkfd djdkr[jdd}jddt|kfd dtjjkfd dtd d tjjtjddjj dS)NrcSri)NzHtorch.linalg.householder_product: input must have at least 2 dimensions.r0r0r0r0r4rcVrkz,linalg_householder_product..rrcSri)Nzbtorch.linalg.householder_product: input.shape[-2] must be greater than or equal to input.shape[-1]r0r0r0r0r4rcZrkcSri)Nz`torch.linalg.householder_product: input.shape[-1] must be greater than or equal to tau.shape[-1]r0r0r0r0r4rc^rkr/cdjdjS)Nzptorch.linalg.householder_product: Expected tau to have one dimension less than input, but got tau.ndim equal to  and input.ndim is equal to rr0rrr0r4rcc cr)Nzltorch.linalg.householder_product: Expected batch dimensions of tau to be equal to input.shape[:-2], but got r0r0actual_batch_tau_shaper0r4rcmcr)Nz,torch.linalg.householder_product: tau dtype z does not match input dtype rr0rr0r4rcus z torch.linalg.householder_productrFrrrr\r|) rTrfrrrr\r empty_stridedrr|)rrexpected_batch_tau_shaper0)rrrr4linalg_householder_productOsD      rcCs^t|dt|ddd||j}||jt|jdd|j|jddtjd}||fS)Nz linalg.inv_exF)rrrrrrrrrrrTr))rrrrr0r0r4linalg_inv_ex_metas  rLDpivotsinfo) hermitianrrcCstt|dt|dtj|jt|jdd|j|jd}|j|jddtj d}|j|jddtj d}|||fS)Nztorch.linalg.ldl_factor_exFrrrrr) rrrTrrrr\r|rr)rrrrrrr0r0r4linalg_ldl_factor_ex_metas    r)rcstdtdtdtjdkfddjdd}t|jkfddttj fddtj j kfddt \}}tj |t |d d j j d S) Nztorch.linalg.ldl_solvercr)NzMtorch.linalg.ldl_solve: Expected B to have at least 2 dimensions, but it has rrr0)rr0r4rcz'linalg_ldl_solve_meta..rcr)Nzjtorch.linalg.ldl_solve: Expected LD.shape[:-1] and pivots.shape to be the same, but got pivots with shape  insteadrr0rr0r4rcrcr)Nz.rrr)rTrfrrrrr) rrsizesrBrkrrrr0rr4linalg_lu_metas$        rLU)rrc stjdkfddtj}|d}|d}tj|t|ddjjd}| t |||d<j |tj d }| j |tj d }|||fS) Nrcr)NzFtorch.lu_factor: Expected tensor with 2 or more dimensions. Got size: rrr0rr0r4rcrz*linalg_lu_factor_ex_meta..rrFrrr) rTrfrrrrrr\r|rrrr) rrrrrBrrrrr0rr4linalg_lu_factor_ex_metas&    r)radjointrcstdtjjkfddtjtjkddtdt|dtddkddtjddjkfddt \}}tj |t || d jj d }| d kru|su|ru|}|S) Nztorch.linalg.lu_solvecr)NzPlinalg.lu_solve: Expected LU and B to have the same dtype, but found LU of type  and B of type rrr0)rrr0r4rc rz&linalg_lu_solve_meta..cSri)NzElinalg.lu_solve: pivots should be a Tensor of scalar type torch.int32r0r0r0r0r4rc'rkzlinalg.lu_solvercSri)NzYlinalg.lu_solve: Number of pivots per batch should be same as the dimension of the matrixr0r0r0r0r4rc/rkcr)Nzclinalg.lu_solve: Expected LU.shape[:-1] and pivots.shape to be the same, but got pivots with shape rrr0rr0r4rc5rrrr)rrTrfr\rrrrrrrrr|rrconj)rrrrrrrRrr0)rrrr4linalg_lu_solve_metas<       r unpack_data unpack_pivotsc stjdkfdd|rt|jtjkddtj}|d}|d}t||}||d<|r9|}ndg}|rX||d<|} ||d<||d<|} n dg} dg} || | fS)Nrcr)NzFtorch.lu_unpack: Expected tensor with 2 or more dimensions. Got size: rrr0rr0r4rcUrz lu_unpack_meta..cSri)Nztorch.lu_unpack: LU_pivots is expected to be a contiguous tensor of torch.int32 dtype. Note: this function is intended to be used with the output produced by torch.linalg.lu_factorr0r0r0r0r4rcZrrr) rTrfrr\r)rrrr) rrrrrrBrrrrrr0rr4lu_unpack_metaKs4           rmodecsddkr d}d}||fSdkrd}d}||fSdkr$d}d}||fStdfdd||fS)NreducedTcompleteFrc ddS)Nzqr received unrecognized mode 'z=' but expected one of 'reduced' (default), 'r', or 'complete'r0r0rr0r4rc z _parse_qr_mode..rTrf)r compute_qrr0r r4_parse_qr_modeus"   rQRrc Cst|dt|dt|\}}|jd}|jd}t||}|r>t|j}|r*|n||d<||}||t|ddn|dg}t|j} |sM|sO|n|| d<|| } | | t| dd|| fS)Nz linalg.qrrrFrr) rrrrrrrrr) rrr reduced_moderBrrQ_shaperR_shaperr0r0r4linalg_qr_metas"           rsign logabsdetcCst|dt|dd|j}||dd}|j|ddt|jd}tj|t|d|j|j d}|j|ddtj d}||||fS)Nzlinalg.slogdetFrrrr) rrrrr^r\rTrrr|r))rrrrrrr0r0r4_linalg_slogdets   r full_matrices compute_uvdrivercCst|dt|dt|jdd}|jd}|jd}t||}|r]|||r*|n|g}||} | |t|dd||rB|n||g} || } t|dk} | | t| | dn |dg} |dg} |j||gt |j d} | | | fS) Nz linalg.svdrrFrrrr) rrrrrrrrrr^r\)rrrrrrBrrU_shaperV_shapeVrSr0r0r4_linalg_svd_metas$           r!arg1arg2cCsn|jdd}|jdd}t||}t|}||d|dg7}t|}||d|dg7}||fS)Nrr)rr"rr)r"r#arg1_batch_sizesarg2_batch_sizesexpand_batch_portionarg1_expand_sizearg2_expand_sizer0r0r4rs rcCsV|rt|||t||\}}||jkr|n||}||jkr"|n||}||fSr=)rrrexpand)r"r#rr'r(arg1_broadcastedarg2_broadcastedr0r0r4rs rrcCs6|jdd}|jdkp|jd|jko|j|k}|S)Nrr/)rr)rrexpected_batched_rhs_shape vector_caser0r0r4linalg_solve_is_vector_rhss  r.)rrrrrrcshtdtjjkfddt}|r dn} t| |dt| \} } t|p6| dd|rC| ddn| } tj| t | | jj d} j }tj|t |djj d}j |ddtj d}j |dd tj d}||||f}| |||f}td d |Drt||D]\}}t||j ||j |t||dd q|S) Nz linalg.solvecdjdjdS)NzKlinalg.solve: Expected A and B to have the same dtype, but found A of type rrrr0rrr0r4rcrz"_linalg_solve_ex..rcSri)Nzlinalg.solve: Vector broadcasting of the left hand side is not supported for left=False. In this case linalg.solve is equivalent to B / A.squeeze(-1)r0r0r0r0r4rc'rrFrrcss|]}|duVqdSr=r0rIr0r0r4ro>z#_linalg_solve_ex..) copy_fromcopy_to exact_dtype)rrTrfr\r. unsqueezerrrrr|rrr)allziprrrr )rrrrrrrrr-B_ B_broad_shaperR result_shaperesult_rLU_pivots_info_rresr or0r0r4_linalg_solve_ex sJ        rA)r unitriangularrrBrc Cs|dur |dg}t|tsJt|||dt||d\}}|ddo+|}|r6t||j }|St ||j rL| |ddj | dd|S)Nrzlinalg.solve_triangularrr) rrlrrrrr@is_conjrrrr transpose_) rrrrrBrr8A_ avoid_copy_Ar0r0r4linalg_solve_triangular_metaHs    rGXM)r4rc stjdkfddtjdkfddtdjtjkrOt\}}tj|t|ddj j d}tj|t|ddj j d}||fSjtj ks[jtj krjt }d g}||fStdd d||fS) Nrcr)NzMtorch.triangular_solve: Expected b to have at least 2 dimensions, but it has rrr0rr0r4rcmrz'triangular_solve_meta..cr)NzMtorch.triangular_solve: Expected A to have at least 2 dimensions, but it has rrr0rr0r4rctrtriangular_solveFrrrcSri)Nz+triangular_solve: Got an unexpected layout.r0r0r0r0r4rcrk)rTrfrrr{stridedrrrr\r| sparse_csr sparse_bsrrr) rrrrrBself_broadcast_sizeA_broadcast_sizesolutioncloned_coefficientr0rr4triangular_solve_metabs<          rRcCspt|dt|d||jdd}||j}||jt|jdd|j|jddtjd}|||fS)Nz linalg.detrFrrrr)rdetrrr0r0r4_linalg_det_metas    rTcstjdkddtjdkdd|rdndtjjdkfddtjjdkfddtjdjdkd dtjjd kfd dtjjkfd djdkrjdd}jddt|kfd djddt|kfddtjjkfddtjjkfddtddtddtjjtjddjjdS)NrcSri)Nz3torch.ormqr: input must have at least 2 dimensions.r0r0r0r0r4rcrkzormqr..cSri)Nz3torch.ormqr: other must have at least 2 dimensions.r0r0r0r0r4rcrkrrcr )Ntorch.ormqr: other.shape[z0] must be greater than or equal to tau.shape[-1]r0r0left_size_conditionr0r4rcrcr )NrUz"] must be equal to input.shape[-2]r0r0rVr0r4rcrcSri)NzHtorch.ormqr: tau.shape[-1] must be less than or equal to input.shape[-1]r0r0r0r0r4rcrkr/cr)Nz[torch.ormqr: Expected tau to have one dimension less than input, but got tau.ndim equal to rrr0rr0r4rcrcr)Nzhtorch.ormqr: Expected other to have the same number of dimensions as input, but got other.ndim equal to rrr0rrr0r4rcrcr)NzWtorch.ormqr: Expected batch dimensions of tau to be equal to input.shape[:-2], but got r0r0rr0r4rcrcr)NzYtorch.ormqr: Expected batch dimensions of other to be equal to input.shape[:-2], but got r0r0)actual_batch_other_shaper0r4rcrcdjdjS)NzPtorch.ormqr: Expected input and tau to have the same dtype, but input has dtype z and tau has dtype rr0rr0r4rcrcrZ)NzRtorch.ormqr: Expected input and other to have the same dtype, but input has dtype z and other has dtype rr0rXr0r4rcrz torch.ormqrrrFrr) rTrfrrr\rrrr|)rrrrrexpected_batch_shaper0)rYrrrWrrr4ormqrsn               r\csttdkfddj}|dk}|}| }|r3td|D] }|o0|dk}q&nt|D] }|oA|dk}q7t|pH|fdddS)NrcsdddtS)Nzpadding size is expected to be rz , but got: rr0)rpaddingr0r4rcz,_padding_check_valid_input..r/rcs dddddjS)N Expected r/zD or rzcD (batch mode) tensor with possibly 0 batch size and other non-zero dimensions for input, but got: rr0)rrr0r4rcs)rTrfrrrr)rr^r input_dim is_batch_modevalid_batch_modevalid_non_batch_moderr0)rrr^r4_padding_check_valid_inputs$    recsd}dd}jdkrd}d7|d7}t|dd|\|}|rHtko>kfddtdkfddjdkra|fS||fS) Nrr/r-rcddddjSNzcArgument #4: Padding size should be less than the corresponding input dimension, but got: padding (rr) at dimension of input rr0dim_wrpad_lpad_rr0r4rc1z_pad1d_common..crp)Nz input (W: z%) is too small. Calculated output W: r0r0)input_woutput_wr0r4rc9rdr)rrrerTrfr)rr^ is_reflection dim_planenbatchnplaner0)rkrrorprlrmr4 _pad1d_commons0       rucCt||ddSNTrq)rurr^r0r0r4meta_reflection_pad1dBrzc*tjtjkfddt|ddS)NcdjdS)Nz)"replication_pad1d" not implemented for ''r\__str__r0rr0r4rcMz(meta_replication_pad1d..Frx)rTrfr\boolruryr0rr4meta_replication_pad1dH   rcsd|stt|dkddjdkrd7|\}||r=t|ko3|kfddtkfddjS)Nr/rcSri)Nz padding size is expected to be 2r0r0r0r0r4rcUrkz(_pad1d_backward_common..r-crfrgrr0rjr0r4rcbrncddSNz(grad_output width unexpected. Expected: , Got: rr0rk grad_outputrpr0r4rcjrNrTrfrrrrr)rrr^rqror0)rkrrrprlrmr4_pad1d_backward_commonRs$     r grad_inputcCt|||ddSrwrrrr^r0r0r4meta_reflection_pad1d_backwardprcCr)NFrxrrr0r0r4meta_replication_pad1d_backwardvrrcs2ddd}d}t|ddj}|dkr'd}d7d7|d7}|\ |}  |rptkoS k fddt kofk fddtdkpydkfd djd kr|fS||fS) Nrr/rrcrfrgrr0rjr0r4rcrnz_pad2d_common..cddddjSNzcArgument #6: Padding size should be less than the corresponding input dimension, but got: padding (rrrhrirr0dim_hrpad_bpad_tr0r4rcrncsddddS)Nz input (H:  W: z%) is too small. Calculated output H: r0r0)input_hrooutput_hrpr0r4rcs r-rerrrTrfr)rr^rq dim_slicesrsrrtr0) rrkrrrorrprrlrmrr4 _pad2d_common|sB        rcCrvrw)rryr0r0r4meta_reflection_pad2dr{rcr|)Ncr})Nz)"replication_pad2d" not implemented for 'r~rr0rr0r4rcrz(meta_replication_pad2d..Frx)rTrfr\rrryr0rr4meta_replication_pad2drrcCst|}t|}||fSr=rTr)grad_wsaved_vsaved_g saved_normsrgrad_vgrad_gr0r0r4meta_weight_norm_backwards  rc sddd}|j}|dkrd7d7|d7}|\}}}}|} |} | ||| ||tkfddtkfdd||jS)Nrr/rrcrrrr0rr0r4rcrNz%meta_pad2d_backward..crNz)grad_output height unexpected. Expected: rrr0rrrr0r4rcrN)rrrTrfrr) rrr^rrrbrlrmrrrror0)rrkrrrpr4meta_pad2d_backwards,        rc sdddd}t|ddjdk}|r+d}d7d7d7|d7}|\  |}     |rt kodk fddtkow k fd dt ko k fd dt dkpdkpdk fd d|r|| fS| fS) Nr-rr/rrcrfrgrr0rjr0r4rc rnz_pad3d_common..crrrr0rr0r4rcrncr)NzcArgument #8: Padding size should be less than the corresponding input dimension, but got: padding (rrrhrirr0)dim_drpad_bkpad_fr0r4rcrnc s(dddddd S)Nz input (D:  H: rz%) is too small. Calculated output D: r0r0)input_drrooutput_drrpr0r4rc#sr)rr^rqrr batch_modersrtr0)rrrkrrrrorrrprrrrlrmrr4 _pad3d_commonsP         rcCrvrw)rryr0r0r4meta_reflection_pad3d/r{rcr|)Ncr})Nz)"replication_pad3d" not implemented for 'r~rr0rr0r4rc:rz(meta_replication_pad3d..Frx)rTrfr\rrryr0rr4meta_replication_pad3d5rrc s(tt|dkdd|jdksJj|jksJddd|jdkr2d7d7d7|\}}}}}}|} |} |} | ||| ||| ||tkfddtkfd dtkfd d||jS) NcSri)Nz padding size is expected to be 6r0r0r0r0r4rcIrkz%meta_pad3d_backward..r-rr/rcrrrr0rr0r4rcarNcrrrr0rr0r4rcerNcr)Nz(grad_output depth unexpected. Expected: rrr0)rrrr0r4rcirNr) rrr^rlrmrrrrrrror0)rrrkrrrrpr4meta_pad3d_backward?s<            rrpcCs^t|dd|d}|dkr|dgjtjdS|||ddfjtjdS)NcSri)Nz(_pdist_forward requires contiguous inputr0r0r0r0r4rcsrkz%meta__pdist_forward..rr/rr)rTrfr@rrrr)rrrr0r0r4meta__pdist_forwardos  rgradpdistcCs8t|ddt|ddtj|tjdS)NcSri)Nz._pdist_backward requires self to be contiguousr0r0r0r0r4rcrkz&meta__pdist_backward..cSri)Nz/_pdist_backward requires pdist to be contiguousr0r0r0r0r4rcrkr)rTrfr@rr)rrrrr0r0r4meta__pdist_backward~s  r)r4r3c sddlm}m}d}d}d} |t|j||| fr-||| ft dkddt dkddt j satj j koVj knfd dj} j| d| dtdko|dkfd d S) Nr) guard_or_truesym_eqr/rr-cSriNzbatch1 must be a 3D tensorr0r0r0r0r4rcrkzmeta_baddbmm..cSriNzbatch2 must be a 3D tensorr0r0r0r0r4rcrkcsdjdjdjS)Nz+Input dtypes must be the same, got: input: z , batch1: z , batch2: rr0)batch1batch2rr0r4rcc &ddddddd SNz@Expected size for first two dimensions of batch2 tensor to be: [rrz ] but got: [rr/].r0r0 batch2_sizesbscontraction_sizer0r4rcs)rrrrrTsym_notrr)rfr exp_config&skip_dtype_check_in_meta_registrationsr\r) rrrr4r3rrdim1dim2dim3 batch1_sizesr0)rrrrrrr4 meta_baddbmms,   rcCrHrrIrrr0r0r4meta_bernoullisr?cCrLr=r0rrrr0r0r4meta_bernoulli_rxrcCrHrrIrr0r0r4meta_bernoulli_prrcC t|Sr=rrr0r0r4 meta_poissonrarc Cs6t| |kddtj|tjd} t|| fS)NcSri)NzJError in fused_moving_avg_obs_fake_quant_cpu: ch_axis must be < self.dim()r0r0r0r0r4rcrkz6meta__fused_moving_avg_obs_fq_helper..r)rTrfrrr)r observer_on fake_quant_on running_min running_maxscale zero_pointaveraging_const quant_min quant_maxch_axisper_row_fake_quantsymmetric_quantmaskr0r0r4$meta__fused_moving_avg_obs_fq_helpers  rcst|dkddt|dkdd|j\|j\tkfdd|durNt||jkpI|tjkoI|jtjtjfvdd|durU|jn|}|jf|dS)NrcSri)Nz a must be 2Dr0r0r0r0r4rcrkzmeta_mm..cSri)Nz b must be 2Dr0r0r0r0r4rcrkc sddddd S)Nz/a and b must have same reduction dim, but got [rrz] X [rr0r0M1M2Nrr0r4rcscSri)NzFout_dtype must be the same as input dtype or fp32 for fp16/bf16 inputsr0r0r0r0r4rcrkr) rTrfrrr\r<r=r>r)r2r3r&rLr0rr4meta_mms"    rcs0|rtfddtjDStjS)Nc3s&|]}|vr j|ndVqdS)r/NrrJrdimsrr0r4ros$z+_compute_reduction_shape..)rerrrOcompute_reduction_output_shaper)rrrhr0rr4rfsrfstrcCsDt|tjjr |jjSt|dr t|jdr |jjdkr |jjSdS)Nr|rtryr)rlrT _subclasses FakeTensor fake_devicerthasattrr|)rr0r0r4rs  r input_tensorrr^dilation is_transposedgroupsoutput_paddingcsLdtdtdtdtdtdtf dd}dtdtdtdtdtd tdtfd d } |jd d} |jd d|r<||jd } n|jd} |jd ||jd krQtd|jd| gt|tre|gt}nt|d krt|dgt}t|tr|gt}nt|d kr|dgt}t|tr|gt}nt|d kr|dgt}d} |rt|tr|gt} nt|d kr|dgt} n|} ttD]2} | r| | || || | | || | | qՈ|| || || | | || qddlm }t |ddd dDfddS)Nlnrrrr{r<cSs$|d|||dd|dS)a Formula to apply to calculate the length of some dimension of the output See: https://pytorch.org/docs/stable/generated/torch.nn.Conv2d.html Args: ln: length of the dimension p: padding in that dim d: dilation in that dim k: kernel size in that dim s: stride in that dim Returns: The output length rr/r0)rrrrr{r0r0r4_formula s$z+calc_conv_nd_return_shape.._formular?cSs(|d|d|||d|dS)a Formula to apply to calculate the length of some dimension of the output if transposed convolution is used. See: https://pytorch.org/docs/stable/generated/torch.nn.ConvTranspose2d.html Args: ln: length of the dimension p: padding in that dim d: dilation in that dim k: kernel size in that dim s: stride in that dim op: output padding in that dim Returns: The output length r/rr0)rrrrr{r?r0r0r4_formula_transposed$ s(z6calc_conv_nd_return_shape.._formula_transposedrr/rzInvalid channel dimensions)sym_orcSsg|]}|dkqSrr0rIr0r0r4rMu rz-calc_conv_nd_return_shape..csdtddddS)NzGiven input size per channel: z&. Calculated output size per channel: rz. Output size is too small)rr0r ret_shaper0r4rcv s z+calc_conv_nd_return_shape..) rrrrlrrrrrrrTrf)rr"rr^rrrrrr kernel_size out_channelsoutput_padding_listrrr0rr4calc_conv_nd_return_shape sd" &         "  rcCtj|tjkSr=rT _prims_commonr channels_lasttenr0r0r4is_channels_last~ r running_mean running_vartrainingexponential_average_factorepsiloncsj}|dur |jn|j} |dur|jn|j} fdd} |j| d} |r4| } | }n d} d}| | |fS)Ncs(trtjSjtjdrtjStjSr)rrTr r@rr0rr0r4pick_memory_format s z2meta_miopen_batch_norm..pick_memory_formatrr)rrr)rr"r$rrrrrrsave_mean_shapesave_var_shaperr save_meansave_varr0rr4meta_miopen_batch_norm s       rc CsJt||||||||r |nd} d} d} || dkrd| | <|| } | SNr/r)rrr) rr"r$rr^rrrr shape_outinput_channels_dimoutput_channels_dimrr0r0r4 meta_conv s     r!mkldnnc CsHt|||||d|g} || } tj} |dkrtj} | j| d} | S)NFrr)rrrTr rchannels_last_3dr) rr"r$r^rrrattrscalars algorithmrrout_memory_formatr0r0r4meta_mkldnn_convolution_default s    r(cCs$|g|jdd|jdRSNrrrr)rr"r$r$r%r&r0r0r4meta_linear_pointwise_default s$r+mklcCs$|g|jdd|jdRSr)r*)r packed_weight orig_weightr$rr0r0r4meta_mkl_linear sr/onednnc Cst||||| d| d}| dur|j} | tjtjtjtjtjfvs"J|j|| d}t |dvs3Jdtj tj tj dt |}|j |d}|S)NFr)r-rrz-Expect output to be 3d/4d/5d for conv1d/2d/3dr)rr\rTr<r>uint8r(r?rrrr r#r)r9x_scalex_zpww_scalew_zpr$rr^rr output_scaleoutput_zero_point output_dtyper$r%r&rrformatr0r0r4meta_qconv_pointwise s>   r;cCs|dksJ|S)Nsumr0)r9r2r3r4r5r6accumr$rr^rrr7r8r9 accum_scaleaccum_zero_pointbinary_op_namer3 unary_op_name unary_op_argsunary_op_algorithmr0r0r4meta_qconv2d_pointwise_binary. s rDc CsJt|j} |jd| d<| tjtjtjtjtjfvsJ|j| | d}|S)Nr/rr) rrrTr<r>r(r1r?r)r9r2r3r4r5r6r$r7r8r9 post_op_name post_op_argspost_op_algorithmrCrr0r0r4meta_qlinear_pointwiseK s  rHcCsV| dkr|St|j}|jd|d<| tjtjtjtjtjfvs"J|j|| d}|S)Nr<r/rr) rrrTr<r>r1r(r?r)r9r2r3r4r5r6x_2r$r7r8r9x2_scalex2_zpr@r3rArBrCrCrr0r0r4meta_qlinear_pointwise_binaryi s  rLcCs&t|j}|jd|d<||}|S)Nr/r)rrr)r9r4r$rCrr0r0r4meta_linear_dynamic_fp16 s  rM quantizedr0rr/c Csrt||||||\}}}|dkr|dnd} tj} |dkr(|||g} n| |||g} tj| |j|j| dSNrr/r-rW)#max_pool2d_checks_and_compute_shaperrrTr rr\r| rrrr^r ceil_mode nInputPlane outputHeight outputWidthrsrrr0r0r4meta_quantized_max_pool2d s$    rXcstdkfddtdkfddtjtjtjtjfvfddtjtjkfddtjtjkfddtjjkfddj  d  d jd S) NrcddS)Nzx must be a 2D tensor, got Drr0rr0r4rc rdz/meta_int4mm_packed_weight_cpu..crY)Nzw must be a 2D tensor, got rZrr0r4r0r4rc rdcrNz#expected x to be f32/f16/bf16, got rr0rr0r4rc rcrNzexpected w to be uint8, got rr0r[r0r4rc rcr)Nz q_group_size must be int64, got rr0) q_group_sizer0r4rc rcr)Nz5q_scale_and_zeros must have the same dtype as x, got rr0)q_scale_and_zerosr0r4rc rrr) rTrfrr\r<r=r>r1rrrr9r4r^r_r0)r^r_r4r9r4meta_int4mm_packed_weight_cpu s$     racs4tkojkfdddS)Ncs8dddddddjS)NzExpected a tensor of dimension z and tensor.size[z] == rrzbut got : dimension z] = rrr0rdim_sizerrr0r4rc sz check_dim_size..)rTrfrr)rrrdrr0rcr4check_dim_size srecsdd}|d|\}} tt|dvddtjtjtjtjtjfvfddt|dkr8|| } } nt|d krH|d|d} } n|d |\} } |d |\} } t|dup_|dkd dd kro dnd } d} d} d}t ||| | d |}t || | | d |}t }t || | | | | d d ||||||dkr|||g}n||||g}tj|jj|dS)NcDtt|dvfdd|d}t|dkr|n|d}||fS)Nr/rcr )Nz avg_pool2d: 4 must either be a single int, or a tuple of two intsr0r0rr0r4rc rz1meta_avg_pool2d..unpack..rr/rTrfrrrvHWr0rir4unpack   zmeta_avg_pool2d..unpackrrr/rcSriNzOavg_pool2d: stride must either be omitted, a single int, or a tuple of two intsr0r0r0r0r4rc rkz!meta_avg_pool2d..cr})Nz""avg_pool2d" not implemented for 'r~rr0rr0r4rc rrr/rr^cSriNzdivisor must be not zeror0r0r0r0r4rc rkrrQrrr-rW)rTrfrr\r1uint16uint32uint64rrpooling_output_shaperOrpool2d_shape_checkrr|)rrrr^rTcount_include_paddivisor_overridernkHkWdHdWpadHpadWrsrU inputHeight inputWidthrVrWrrr0rr4meta_avg_pool2d sj              rcCsjt|||||||dd| | | | | ||}| }t|||d|t|||d| t|||d| dS)Nr/r-r)rxrre)r gradOutputrsr{r|r}r~rrrUrrrVrW mem_formatr nOutputPlaner0r0r4avg_pool2d_backward_shape_check- s,rcCstt|dkp t|dkdd|d}t|dkr|n|d} tt|dkp5t|dkp5t|dkddt|dkrB|n|d} t|dkrN| n t|dkrV| n|d} tt|dkpgt|dkdd|d} t|dkrx| n|d} t|dup|dkdd|j}|d kr|d nd}|d }|d }|d }t||| | d|}t|| | | d|}t|}t||||| | | | | ||||||tj ||j |j |dS)Nr/rcSri)NzKavg_pool2d: kernel_size must either be a single int, or a tuple of two intsr0r0r0r0r4rcg rkz*meta_avg_pool2d_backward..rcSrirqr0r0r0r0r4rcm rkcSri)NzGavg_pool2d: padding must either be a single int, or a tuple of two intsr0r0r0r0r4rcs rkcSrirrr0r0r0r0r4rcz rkrrQrsrrrW) rTrfrrrrwrOrrrr\r|) gradOutput_rrrr^rTryrzr{r|r}r~rr input_sizersrUrrrVrWrr0r0r4meta_avg_pool2d_backwardY sj "( rcs6tt|dvdd|d}t|dkr|n|d}t|dkr$|n|d} t| p2t|dvddtjtjtjtjtjfvfdd|sP|n|d} |sX|n t|dkr`| n|d} |sh| n t|dkrp| n|d} tt|dvd d|d} t|dkr| n|d}t|dkr| n|d}tjd vd dt| p|dkd d d} d } d} d} d}t ||| | d|}t |||| d|}t || || d|}t |||| | | | | ||ddd||||||dddjdkr ||||fS |||||fS)Nr/r-cSriNzFavg_pool3d: kernel_size must be a single int, or a tuple of three intsr0r0r0r0r4rc rkz!meta_avg_pool3d..rr/rcSriNzJavg_pool3d: stride must be omitted, a single int, or a tuple of three intsr0r0r0r0r4rc rkcr})Nz""avg_pool3d" not implemented for 'r~rr0rr0r4rc rcSriNzBavg_pool3d: padding must be a single int, or a tuple of three intsr0r0r0r0r4rc rkrrcSriNz9non-empty 4D or 5D (batch mode) tensor expected for inputr0r0r0r0r4rc rkcSrirrr0r0r0r0r4rc rkrQrsrrz avg_pool3d()T)check_input_sizer) rTrfrr\r1rtrurvrrrwpool3d_shape_checkr)rrrr^rTryrzkTr{r|dTr}r~padTrrrsnslicesitimeiheightiwidthotimeoheightowidthr0rr4meta_avg_pool3d s            rcCstt|dvdd|d}t|dkr|n|d} t|dkr$|n|d} t| p2t|dvdd|s;|n|d} |sC| n t|dkrK| n|d} |sS| n t|dkr[| n|d} tt|dvdd|d}t|dkrw|n|d}t|dkr|n|d}t|jd vd dt| p|dkd d|d }|d }|d}|d}t|||| d|}t|| || d|}t|| || d|}t||||| | | | | |||||||||d||jS)NrcSrirr0r0r0r0r4rc rkz*meta_avg_pool3d_backward..rr/rcSrirr0r0r0r0r4rc rkcSrirr0r0r0r0r4rc rkrcSrirr0r0r0r0r4rc" rkcSrirrr0r0r0r0r4rc' rkrQrsrrzavg_pool3d_backward()) rTrfrrrrwavg_pool3d_backward_shape_checkrr)rrrrr^rTryrzrr{r|rr}r~rrrrrrrotime_for_shape_checkoheight_for_shape_checkowidth_for_shape_checkr0r0r4meta_avg_pool3d_backward st          rcsZtjdkp jdkfddjddt|}t}tj|jj |dS)Nr-rcr)Nz"Expected 3D or 4D tensor, but got rr0rr0r4rcP rz*meta_adaptive_avg_pool2d..rrW) rTrfrrrerOrrr\r|)r output_sizerCrr0rr4meta_adaptive_avg_pool2dL s  rcs@tjdkp jdkfddjddt|S)Nrrcr)Nz"Expected 4D or 5D tensor, but got rr0rr0r4rcb rz*meta_adaptive_avg_pool3d..rs)rTrfrrrre)rrr0rr4meta_adaptive_avg_pool3d^ s  rcsj}td|D]tdkfddqt|dkp$|dkfddtjjkfddtj}trDtj} j j |d S) Nr/rcsdjddS)Nz{adaptive_avg_pool2d_backward(): Expected grad_output to have non-zero size for non-batch dimensions,  with dimension being emptyrr0)grad_outrr0r4rcm s z4meta__adaptive_avg_pool2d_backward..r-rcr)NzBadaptive_avg_pool2d_backward(): Expected 3D or 4D tensor, but got rr0rr0r4rcr rcrNexpected dtype z! for `grad_output` but got dtype rr0)rrr0r4rcv rr) rrrTrfrr\rrr rrr)rrrrr0)rrrr4"meta__adaptive_avg_pool2d_backwardg s$     rcCst|dtj|tjdS)Nadaptive_avg_pool3d_backwardr)!_adaptive_pool_empty_output_checkrTrrrrr0r0r4"meta__adaptive_avg_pool3d_backward~ s rrcs<j}td|D]tdkfddqdS)Nr/rcsdjddS)Nzc(): Expected grad_output to have non-zero size for non-batch dimensions, but grad_output has sizes rrrr0rrrr0r4rc s z3_adaptive_pool_empty_output_check..)rrrTrfr)rrrr0rr4r s rc s"j}t|dvfddtd|D]tdkfddqtt|dkddd}d}d}jd krGd}|d7}|d}|\}}jd krm|||f}|} j|tjd } | | fS||||f}t } |j | d } j|tjd j | d } | | fS) Nr-rcr)Nz:adaptive_max_pool2d(): Expected 3D or 4D tensor, but got: rr0rr0r4rc rz*meta_adaptive_max_pool2d..r/rcdjddS)Nzjadaptive_max_pool2d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes rrrr0rrr0r4rc rcSri)NzCadaptive_max_pool2d(): internal error: output_size.size() must be 2r0r0r0r0r4rc rkrr-rr) rrTrfrrrrrrOrr) rrrdimHsizeBsizeDosizeHosizeWrrrrr0rr4meta_adaptive_max_pool2d sD           rcsdj}t|dvfddtdtjjkfddt}jj |dS)Nrcr)NzKadaptive_max_pooling2d_backward(): Expected 3D or 4D grad_output, but got: rr0rr0r4rc rz3meta_adaptive_max_pool2d_backward..adaptive_max_pool2d_backwardcrrrr0rrr0r4rc rr) rrTrfrr\rOrrrr)rrrrrr0rr4!meta_adaptive_max_pool2d_backward s     rc sj}t|dvfddtd|D]tdkfddqtt|dkddd}d}d}|d krFd}|d7}|}|\}}}|d kr[||||f} n|||||f} | } j| tjd } | | fS) Nrcr)Nz:adaptive_max_pool3d(): Expected 4D or 5D tensor, but got: rr0rr0r4rc rz*meta_adaptive_max_pool3d..r/rcr)Nzjadaptive_max_pool3d(): Expected input to have non-zero size for non-batch dimensions, but input has sizes rrrr0rr0r4rc rr-cSri)NzCadaptive_max_pool3d(): internal error: output_size.size() must be 3r0r0r0r0r4rc rkrrr)rrTrfrrrrr) rrrdimDrrosizeTrrrrrr0rr4meta_adaptive_max_pool3d s8        rcCst|d||jS)Nadaptive_max_pool3d_backward)rrr)rrrr0r0r4!meta_adaptive_max_pool3d_backward s  rcCs|durtd||S)Nz:cannot repeat_interleave a meta tensor without output_size)rr)repeatsrr0r0r4meta_repeat_interleave_Tensor s rcCsD|jjsJ|jjs Jt|t|j|t|jtjd}|SNrG)r\rrSrrrrP)realimagrr0r0r4 meta_complex s  r) fill_valuercCs|j||ftjdSrc)rrrTr)rrrr0r0r4nonzero_static srcCs<ttjddtj||fd|ftj|jdS)NcSri)NaYThe register_meta function for torch.nonzero() raises unimplemented by default, as a correct data-independent implementation does not exist. This implementation returns a fake value, assuming all elements of the tensor are non-zero. To enable this registration, please set 'torch.fx.experimental._config.meta_nonzero_assume_all_nonzero' to True.r0r0r0r0r4rc' rkznonzero..r/r\r|) rT_check_not_implementedrmeta_nonzero_assume_all_nonzerorrrrr|rr0r0r4nonzero" s rc sDttddg}tD]q\dur|tjtjtjtjtjfvddjtjtjfvrv}t |t j j kfddt j D]#t j j kfdd||dqQq|q|q|tt j kfddddlm}t|jt j krdt j ksd}d }D]|dkrLjdurd}q|dkr҈durd }qdurnqd }|sg}g}tD]\dur||qtD]\dur||q||gggtD]&\} durBr8j | q"j | q"tj q"fd d } } ddlm} | dkrk| S| } t| \}}t|tt t |krt| j |}t|}t|t|}| | |} | S)NcSri)Nz#at least one index must be providedr0r0r0r0r4rc7 rkz#meta_index_Tensor..cSri)Nz?tensors used as indices must be long, int, byte or bool tensorsr0r0r0r0r4rc? rkcr)N)too many indices for tensor of dimension rr0rr0r4rcF rc s$djddjdS)NzThe shape of the mask at index z0 does not match the shape of the indexed tensor rr0)rrjrrr0r4rcK s r/csdjdtdS)Nrz (got rs)rrr0)rrr0r4rcV r_rFrTcsL}t|}dgt|tt|jt<|||S)zI This follows restride_src in TensorAdvancedIndexing.cpp r)rrrrr*)rrr) after_shape before_shapereplacement_shaper0r4 _restride_src s   z(meta_index_Tensor.._restride_srcguard_or_false) rTrfr enumerater\rrr(rrrrrrrselect torch._refs_refsrr#rrrrrrO3compute_elementwise_output_logical_to_physical_perm apply_permr invert_permr*r)rrrrrefsstatehas_contiguous_subspacertransposed_indicesrrrrrestrided_selfpermrR perm_shaperr0) rrrrrrrrrr4meta_index_Tensor5 s                rc CsTd} d} d} | dr||} | dr||} | dr%||} | | | fSNrr/rrr) grad_output_input_weight_bias_sizes_optrr^r transposedrr output_maskbackend_grad_inputbackend_grad_weightbackend_grad_biasr0r0r4meta_convolution_backward s  rcsd}d}|||f}tdkddtdkddtddkfddtddkfd dt|d|ko^|d|kd d||S) Nr/rr-cSrirr0r0r0r0r4rc rkzmeta_addbmm..cSrirr0r0r0r0r4rc rkrcddddS)Nz8batch1 and batch2 must have same number of batches, got rrrr0rrr0r4rc rc 6ddddddddd S)Nz#Incompatible matrix sizes for bmm (r/r9rrrsrr0rr0r4rc cSri)Nz.self tensor does not match matmul output shaper0r0r0r0r4rc rk)rr)rTrfrr)rrrr4r3rrr0rr4 meta_addbmm s$    rcKs||Sr=r)rrkwargsr0r0r4meta_randint_like sr) grad_scale found_infc s4||||||fD]tttfddqdS)NcdtSNz'exponent must be a tensor list but got rtr0lr0r4rc rz#meta__fused_adam_..rTrfrlr)rgradsexp_avgs exp_avg_sqsmax_exp_avg_sqs state_stepslrbeta1beta2 weight_decayepsamsgradmaximizerrr0rr4meta__fused_adam_ s  rc sZ||||||fD]tttfddqdd}||||||||||fS)Ncrrrr0rr0r4rc&rz"meta__fused_adam..cSsdd|DS)NcSsg|]}t|qSr0r)rJrr0r0r4rM*rNz=meta__fused_adam..empty_like_list..r0) tensor_listr0r0r4empty_like_list)sz)meta__fused_adam..empty_like_listr)rrrrrrr r r r r rrrrrr0rr4meta__fused_adams rcstdkddtdkddtjtjufddtjtjufddtddkfd djddftjd S) NrcSri)Nza must be a 2D tensorr0r0r0r0r4rc8rkzmeta__int_mm..cSri)Nzb must be a 2D tensorr0r0r0r0r4rc9rkcr)Nzexpected self to be int8, got rr0)r2r0r4rc<rcr)Nzexpected mat2 to be int8, got rr0)r3r0r4rc@rr/rc r)Nz'Incompatible matrix sizes for _int_mm (rr9r/rrsrr0r1r0r4rcDrr)rTrfrr\r(rrr)r1r0r1r4 meta__int_mm5s      rcsttdkddtjtjufddd}dd}j|d||dd |dftjd S) NrcSriNzw must be a 2D tensorr0r0r0r0r4rcNrkz2meta__convert_weight_to_int4pack..crr]rr0r[r0r4rcQrrr/ r)rTrfrr\r1rrr)r4 inner_k_tilesrrr0r[r4 meta__convert_weight_to_int4packLs    rcs`tdkddtjtjufddd}d}j||dftjdS)NrcSrirr0r0r0r0r4rcbrkz:meta__convert_weight_to_int4pack_for_cpu..crNzexpected w to be int32, got rr0r[r0r4rcerrr/r)rTrfrr\r)rrr1rr0r[r4(meta__convert_weight_to_int4pack_for_cpu`s     rcstdkddtdkddtjtjtjtjfvfddtjtjufddj d dd jd S) NrcSriNzx must be a 2D tensorr0r0r0r0r4rcqrkz*meta__weight_int4pack_mm..rcSri)Nzw must be a 4D tensorr0r0r0r0r4rcrrkcrr\rr0rr0r4rcurcrrrr0r[r0r4rcyrrrr rTrfrr\r<r=r>r)rrr`r0r4r9r4meta__weight_int4pack_mmos   "r!ctdkddtdkddtjtjtjtjfvfddtjtjufddj d djdS) NrcSrirr0r0r0r0r4rcrkz2meta__weight_int4pack_mm_for_cpu..cSrirr0r0r0r0r4rcrkcrr\rr0rr0r4rcrcrr]rr0r[r0r4rcrrr) rTrfrr\r<r=r>r1rrr`r0r r4 meta__weight_int4pack_mm_for_cpu~   r#cr") NrcSrirr0r0r0r0r4rcrkz;_weight_int4pack_mm_with_scales_and_zeros..cSrirr0r0r0r0r4rcrkcrr\rr0rr0r4rcrcrrrr0r[r0r4rcrrrr)r9r4r^qScaleqZerosr0r r4)_weight_int4pack_mm_with_scales_and_zerosr$r'r2r3cCr7r.r0r1r0r0r4 kai_roundupsr(c s|dkrv||kr/d}d}d}d dddd fddfd d }||||||S|d d krx||d krzd}d}d}d ddd fd d} fddddfddfdd |||||||SdSdSdS)NrrrrcSst||d}t||S)Nrr()rkrsrkr_sr_roundedup4r0r0r4kai_k_roundedups z3get_kai_packed_weight_size..kai_k_roundedupcs8|||}|ddksJd||dS)Nrrzk_internal must be evenr0)rnrr*r+ k_internal)r-kai_num_bytes_biaskai_num_bytes_multiplier_rhskai_num_bytes_sum_rhsr0r49kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0s z]get_kai_packed_weight_size..kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0cs t|||}|||||Sr=r))rrr.r*r+num_rows)r3r0r47kai_get_rhs_packed_size_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0sz[get_kai_packed_weight_size..kai_get_rhs_packed_size_rhs_pack_nxk_qsi4cxp_qsu4cxs1s0rrcsR||dksJ|dksJ|dksJt|||}||||||SrNr))rrr.r*r+blr4)kai_bl_multiple_of;kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0kai_nr_multiple_ofr0r49kai_get_rhs_packed_size_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0s z]get_kai_packed_weight_size..kai_get_rhs_packed_size_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0cs^||dksJ|dksJ|dksJ}||}||}|||SrNr0)rr.r*r+r6num_bytes_multiplier_rhsnum_blocks_per_rownum_bytes_per_block)r7#kai_get_bf16_datatype_size_in_bytesr9kai_num_blocks_per_rowr0kai_num_bytes_per_blockr2r0r4r8s z_get_kai_packed_weight_size..kai_get_rhs_packed_stride_rhs_pack_nxk_qsi4c32p_qsu4c32s1s0cSri)Nrr0r0r0r0r4r>rzGget_kai_packed_weight_size..kai_get_bf16_datatype_size_in_bytescs|dksJt|||SrNr))rr6r7r0r4r?sz:get_kai_packed_weight_size..kai_num_blocks_per_rowcs|dksJ|d|S)Nrrr0)r6r;rAr0r4r@s z;get_kai_packed_weight_size..kai_num_bytes_per_blockr0) n_bitsrK groupsizekai_nrkai_krkai_srr5r:r0) r7r>r8r3r-r9r?r0r1r@r2r4get_kai_packed_weight_sizes@     .rHcstjtjufddtjjrE||kr|jtjks4||krE|ddkrE||dkrE|jtjkrEt d|||}j t |tjdS | }|durW|| 7}j |tjdS)Ncrr]rr0weightsr0r4rcrz2meta__dyn_quant_pack_4bit_weight..rrrr) rTrfr\r1backendskleidiai is_availablerXr>rHrrr)rJ scales_zerosr$ block_size in_features out_featurespacked_weight_sizer0rIr4 meta__dyn_quant_pack_4bit_weights$       rScshtdkddtjtjkpjtjkokfddd}j||jdS)NrcSri)Nzinput must be a 2D tensorr0r0r0r0r4rc1rkz-meta__dyn_quant_matmul_4bit..csdjddS)NzPexpected input to be f32 or bf16 (bf16 requires block_size == in_features), got z with block_size=z and in_features=rr0rOrPinpr0r4rc5srr)rTrfrr\r<r>rr)rUpacked_weightsrOrPrQrIr0rTr4meta__dyn_quant_matmul_4bit)s  rWcstdkddtjtjtjtjfvfddtdkddtjtjufddj d djdS) NrcSrirr0r0r0r0r4rc@rkz*meta__weight_int8pack_mm..crr\rr0rr0r4rcCrcSrirr0r0r0r0r4rcErkcr)Nzexpected w to be int8, got rr0r[r0r4rcHrrr) rTrfrr\r<r=r>r(rr)r9r4q_scalesr0r r4meta__weight_int8pack_mm>s   rYc stdkfddtdkfddtddkfddttjfddttjfddt|d kd dtd vfd dd }d }jdd }jdd }tt ||}| ||g |S)NrcrY)Nz1cdist only supports at least 2D tensors, X1 got: rZrr0x1r0r4rcQrdz$meta_cdist_forward..crY)Nz1cdist only supports at least 2D tensors, X2 got: rZrr0x2r0r4rcUrdrcr)Nz4X1 and X2 must have the same number of columns. X1: rz X2: rr0)r[r]r0r4rcYrcr)Nz3cdist only supports floating-point dtypes, X1 got: rr0rZr0r4rc]rcr)Nz3cdist only supports floating-point dtypes, X2 got: rr0r\r0r4rcarrcSri)Nz)cdist only supports non-negative p valuesr0r0r0r0r4rccrkrcr)Nz(possible modes: None, 0, 1, 2, but was: r0r0) compute_moder0r4rcfrr) rTrfrrrOis_float_dtyper\rrbroadcast_shapesextendr) r[r]rr^r1r2 batch_tensor1 batch_tensor2rCr0)r^r[r]r4meta_cdist_forwardMs@             rfc Cs|jd}|jd}|jd}|jdd}|jdd} tt|| } | } | ||gt| } |dksE|dksE|dksE| dkrJt|S| t|jkrV| | }tj |tj dS)Nrrrr) rrrTr`copyramathprod zeros_liker)rr) rr[r]rcdistc1rbrcrdrer&tensor1_expand_size batch_productr0r0r4meta_cdist_backwardqs       roc s tjtjtjfvfddtjtjtjfvfddttjfddd} |rEt| dkdd| d8} | d} durzt|t kddtj dkfd dt kfd dfd d d dfdd} t dkrd} } |t kr| d}nRd}nL| | |}|tt fvs|s̈d} nd} | } jd}|t kr|rt|dkdd|d8}|jd}n| }| | | |fS)Ncr)Nz(expected indices to be long or int, got rr0)rr0r4rcrz$meta_embedding_bag..cr)Nz(expected offsets to be long or int, got rr0)rRr0r4rcrcr)Nz/expected weight to be floating point type, got rr0)r"r0r4rcrrr/cSriNz1include_last_offset: numBags should be at least 1r0r0r0r0r4rcrkcSri)Nz@embedding_bag: per_sample_weights only supported with mode='sum'r0r0r0r0r4rcrkcr)Nz1expected per_sample_weights to be 1D tensor, got rZrr0)per_sample_weightsr0r4rcrcsdddS)Nz%expected per_sample_weights.numel() (z$ to be the same as indices.numel() (rsrr0)rrqr0r4rcs cs|||o |ddkSNrr/rrrr padding_idx)is_fast_path_index_selectr0r4is_fast_path_index_select_scalesz;meta_embedding_bag..is_fast_path_index_select_scalecSs<|jtjks |jtjko|ddko|ddko|dkSr)r\rTrXrVr)rrrvr0r0r4rws  z5meta_embedding_bag..is_fast_path_index_selectcs"|dur ||||S|||Sr=r0ru)rwrxr0r4 is_fast_paths z(meta_embedding_bag..is_fast_pathcpucSrirpr0r0r0r0r4rcrk)rTrfr\rrrOr_rrMODE_SUMrrrMODE_MAX MODE_MEANr)r"rrRscale_grad_by_freqrsparserqinclude_last_offsetrvnum_bagsrry offset2bagbag_size max_indices fast_path_sumnumBagsr0)rrwrxrRrqr"r4meta_embedding_bagst               rcGsBt|||g|R\}}}}t|dkr||}||||fS)Nrz)rrrr)r"rrRrQrrrrr0r0r4meta_embedding_bag_forward_onlys  rcCs.|r|S|jjs |jjr|jS|rtjS|jSr=)r\rrrTr)rr\promote_int_to_longr0r0r4_get_reduction_dtypesrrcCs6t||dd}t|j|}t|||}|j||dS)NT)rr)rrOrerrfr)rrrhr\r9rCr0r0r4 meta_nansums rcCs$t|jtt|}||Sr=)rOrrrerrr)rrCr0r0r4 meta_median s rcCsLt|dkr tdt|j|f}t|||}|||j|tjdfS)Nrzmedian CUDA with indices outputr) rrOalert_not_deterministicrerrfrrTr)rrrhrCr0r0r4meta_median_mode_dims  rcCrLr=r0rr0r0r4meta_logical_not_*rxrcstt|kddtD]\tdkfddqt|}d|t|jfddttD}||S)NcSri)NzZNumber of dimensions of repeat dims can not be smaller than number of dimensions of tensorr0r0r0r0r4rc3rkzmeta_repeat..rcr_)Nz"Repeats cannot be negative, found rr0r0)rrepr0r4rc8rdrOcsg|] }||qSr0r0r) padded_sizerr0r4rM?rzmeta_repeat..) rTrfrrrrerrr)rrnum_new_dimensions target_sizer0)rrrrr4 meta_repeat/s  rcCrLr=r0rr0r0r4 meta_zero_CrxrcCst|tjr t|j|j|Sr=)rlrTr rhrrrr0r0r4meta_binop_inplaceHs rcCsfdd}dd}dd}||r||rtd||r$||s$tdt|tjr1t|j|j|S) a* Some checks for inplace ops. Checks for promotion rules for some dtypes. int.add/sub_(float) and bool.add/sub_(others) are rejected. Promoting in these in-place operations would require reallocating and copying over elements, hence not allowed. Checks for alpha param. cS t|tr t|jSt|tSr=)rlrrOrr\rrnr0r0r4 is_integerick   z.meta_binop_inplace_alpha..is_integericcSrr=)rlrrOr_r\rrr0r0r4 is_floaticqrz,meta_binop_inplace_alpha..is_floaticcSrr=)rlrrOis_boolean_dtyper\rrr0r0r4 is_booleanicwrz.meta_binop_inplace_alpha..is_booleanicz]Promotion of int.add/sub_(float) in in-place ops are not possible due to element size change.z_Promotion of book.add/sub_(others) in in-place ops are not possible due to element size change.)rrlrTr rhr)rrr3rrrr0r0r4meta_binop_inplace_alphaYs rcCt||tjdSrrSrrPrrr3r0r0r4meta_binop_alphasrcKt|tjdSrr)rrr0r0r4 meta_roundsrcslttjfddttjr&ttjfdddStttfdddS)NcdjS)Nz7: Expected input tensor to have an integral dtype. Got rr0)rrr0r4rcrdz#shift_dtype_check..cr)Nz6: Expected shift value to have an integral dtype. Got rr0rrvr0r4rcrdcsdS)Nz): Expected shift value to be an int. Got r0r0rr0r4rcr)rTrfrOrr\rlr rrrrvr0rr4shift_dtype_checks      rcCtd||t||tjdS)NrshiftrrrSrrPrr0r0r4 meta_rshifts rcCr)Nlshiftrrrr0r0r4 meta_lshiftsrrcCs ||jSr=r*rr0r0r4 meta_zero rcCrLr=r0rrvr0r0r4 meta_fill_rxrcCrr=rrr0r0r4 meta_fill rcCrLr=r0rr0r0r4 meta_relu_rxrcCrrrrr0r0r4meta__add_relur?UUUUUU?cCrr=rrnoiselowerrrrr0r0r4meta_rrelu_with_noises rcCst|t|fSr=rrr0r0r4 meta_rrelu_with_noise_functionalsrcCrLr=r0)rrrrrr0r0r4meta_rrelu_with_noise_srcCrr=rrrr accumulater0r0r4meta_index_putrrcCst|j|j|Sr=rhr)rrvaluer0r0r4meta_masked_fill_srcCs ||jt|d}|Sr)rrrrOr)rrr masked_scaler0r0r4meta__masked_scalesrcs@t|jtjtjfvddtjjkfddS)NcSri)NzMask must be bool or uint8r0r0r0r0r4rc rkz&meta_masked_scatter_..crZ)NzEmasked_scatter: expected self and source to have same dtypes but got rrr0rrFr0r4rcs )rTrfr\rr1)rrrFr0rr4meta_masked_scatter_s  rcCs*t||\}}tj|tjd}t|||Sr)r#rTrrr)rrrFrr0r0r4meta_masked_scatters rcCs ||Sr=r_)rrrr0r0r4meta_masked_scatter_backwardrrcCrLr=r0rr0r0r4meta_index_put_!rxrc sPddlm}m}t|dkddt|dkdd|}||d|d|d}d} || ft||d|dfd d|r|jtjkpc|jtj koh|tj k} t||jkpq| d d|  |} n| } |sdurtdkd dt|fd d| S) Nr)sym_andrr-cSrirr0r0r0r0r4rc)rkz)common_meta_baddbmm_bmm..cSrirr0r0r0r0r4rc*rkrr/c rrr0r0rr0r4rc7scSri)Nzfout_dtype only supported for torch.float32 output with float16/bfloat16 inputs or same as input dtypesr0r0r0r0r4rc@rkcSri)Nzself must be a 3D tensorr0r0r0r0r4rcHrkcsddS)Nz*Expected an input tensor shape with shape z but got shape: rr0)r self_baddbmmr0r4rcKr) rrrrTrfrrr\r=r>r<rr) rris_bmmrr&rrrres_rowsres_colssupported_out_dtyperr0)rrrrrr4common_meta_baddbmm_bmm&s@      rcCs t||dS)NTr)rr0r0r0r4meta_bmmQrrcCst||d|dS)NT)r&r)rr0r&r0r0r4meta_bmm_dtypeVsrcCs<||}||}|dkrt|dkt|dkkr|d8}|Srs)r)r9r:qr r0r0r4div_rtn[s  rcCsZt|||||dd|r|dnd|d}|r+|d|||kr+|d8}|Sr)r) inputSize kernelSizerlrmrrrT outputSizer0r0r4pooling_output_shape_pad_lres*    rcslt|dkddtdkfddtdddkfddt|||S)NrcSri)Nzstride should not be zeror0r0r0r0r4rcrkz&pooling_output_shape..cr)Nz'pad must be non-negative, but got pad: r0r0padr0r4rcrr/rcsdddS)NzApad should be at most half of effective kernel size, but got pad=z, kernel_size=z and dilation=r0r0rrrr0r4rcs )rTrfr)rrrrrrTr0rr4rwsrwc sd} tdkodkfddtdko dkfddtdko1dkfddddkoFddk}|tjkrbt|dkoZ|oZd dkfd dn"t|d krqddkrq|p}|dko}|o}d dkfd dtdkod k fd dt dko dk fd ddS)Nrcrp)Nz5kernel size should be greater than zero, but got kH: , kW: r0r0)r{r|r0r4rcrdz$pool2d_shape_check..crp)Nz0stride should be greater than zero, but got dH: , dW: r0r0)r}r~r0r4rcrdcrp)Nz9dilation should be greater than zero, but got dilationH: , dilationW: r0r0) dilationH dilationWr0r4rcrdr/rrr-cdS)NzExpected 4D (batch mode) tensor expected for input with channels_last layout with optional 0 dim batch size for input, but got: rr0rr0r4rcscr)NzYExpected 3D or 4D (batch mode) tensor with optional 0 dim batch size for input, but got: rr0rr0r4rcrcsddddS)NzKpad should be smaller than or equal to half of kernel size, but got padW = z , padH = z, kW = z, kH = r0r0)r{r|rrr0r4rcsc s*ddddddd SNzGiven input size: (r9z). Calculated output size: (z). Output size is too smallr0r0)rrrUrrVrWr0r4rcs)rrTrfrr )rr{r|r}r~rrrrrUrrrVrWrr valid_dimsr0)r}r~rrrrrr{r|rUrrVrWrrr4rxsB      rxrrr{r|rr}r~pTpHpW dilationTrrrrrrrrrc s j}t dkodko dk fddtdko&dko&dkfddtdko<dko<dkfddt|dv fddt|D]|dkradkraqVt dk fd dqV|rt ko ko k fd dt d kod ko d k fd dtd kod kod k fdddS)NrcdddS)Nz5kernel size should be greater than zero, but got kT: z, kH: rr0r0)r{rr|r0r4rcz$pool3d_shape_check..cr)Nz0stride should be greater than zero, but got dT: z, dH: rr0r0)r}rr~r0r4rcscr)Nz9dilation should be greater than zero, but got dilationT: z , dilationH: rr0r0)rrrr0r4rcrrcr)Nz/: Expected 4D or 5D tensor for input, but got: rr0)rrr0r4rcrdrcsdjddS)NzZ: Expected input's non-batch dimensions to have positive length, but input has a shape of z and non-batch dimension z has length zero!)rrr0)rrrr0r4rcs c s*ddddddd S)Nzinput image (T: rrz ) smaller than kernel size (kT:  kH:  kW: rsr0r0)rrrr{rr|r0r4rcsrc s(dddddd S)NzHpad should be smaller than or equal to half of kernel size, but got kT: rrz padT: z padW: z padH: r0r0)r{rr|rrrr0r4rcsr/cs6dddddddddSrr0r0)rrrrrrrr0r4rc!s)rrTrfrr)rrrr{r|rr}r~rrrrrrrrrrrrrrrr0)r}rr~rrrrrrrrrr{rr|rrrrrrrr4rsJ    "rcCs|j}t|||||||| | | | | |||||||||t|||d|t|||d|t|||d|t|||d|t|||d|t|||d|t|||d|t|||d|dS)Nrr-rr/rrre)rrrrrr{r|rr}r~rrrrrrrrrrrrrrr0r0r4max_pool3d_backward_shape_check)s@rcCs|j}t||||||||| | | ddd| | |||||dt|||d|t|||d|t|||d|t|||d|dS)Nr/Trr-rr)rrrrr{r|rr}r~rrrrrrrrrrrr0r0r4rgs:rcCsBdd}|d|\}}tt|dvddt|dkr#||} } n|d|\} } |d |\} } |d |\} }|d }|d }|d }t|}|tjkr^t|dkddn|tjkrpt|dvddntdddt ||| | | |}t ||| | ||}t |||| | | | | ||||||||||fS)Ncrf)Nrgcr )Nz max_pool2d: rhr0r0rir0r4rcrzEmax_pool2d_checks_and_compute_shape..unpack..rr/rjrkr0rir4rnroz3max_pool2d_checks_and_compute_shape..unpackrrpcSri)NzOmax_pool2d: stride must either be omitted, a single int, or a tuple of two intsr0r0r0r0r4rcrkz5max_pool2d_checks_and_compute_shape..rrr^rrsrrrcSri)NzMnon-empty 4D (batch mode) tensor expected for input with channels_last layoutr0r0r0r0r4rcrkrcSri)Nz9non-empty 3D or 4D (batch mode) tensor expected for inputr0r0r0r0r4rcrkFcSri)NzAUnsupported memory format. Supports only ChannelsLast, Contiguousr0r0r0r0r4rcrk) rTrfrrrOrr rrrwrx)rrrr^rrTrnr{r|r}r~rrrrrUrrrrVrWr0r0r4rRsb              rRc st|||||\}tjjkfdd|jfdd} | | |t} tjjjj | dS)Ncr)NzExpected dtype z for `gradOutput` but got dtype rr0rr0r4rcrz7meta_max_pool2d_with_indices_backward..cs:t|dt|dt|ddS)Nr-rr/)re)r)rrrVrWr0r4_check_dim_sizesz>meta_max_pool2d_with_indices_backward.._check_dim_sizerW) rRrTrfr\rrOrrrr|) rrrrr^rrTrrUrrr0)rrrrVrWrr4%meta_max_pool2d_with_indices_backwards.    rc Cst||||||\}}}|dkr|dnd} t|} |dkr*|||g} n| |||g} tj| |j|j| dtj| tj |j| dfSrP) rRrrrOrrTrr\r|rrSr0r0r4meta_max_pool2d_with_indicess2     rc stjdvfddj}t|d|D]tdkfddqttdkddtt|dkd dd }d d |d kr[d}nd}tjjkddtjdkfddd}d}dt||kddt||kddtdkfddt|dddkfddt|dddkfddd kr|||d|dg} n ||d|dg} tj| jj dtj| tj j dfS)Nrcr)Nz:fractional_max_pool2d: Expected 3D or 4D tensor, but got: rr0rr0r4rc@rz,meta_fractional_max_pool2d..r-rcsdddS)Nz_fractional_max_pool2d: Expected input to have non-zero size for non-batch dimensions, but got rz emptyrr0)rrr0r4rcGs rcSri)NzNfractional_max_pool2d: kernel_size musteither be a single int or tuple of Intsr0r0r0r0r4rcNrkcSri)NzOfractional_max_pool2d: output_size must either be a single int or tuple of Intsr0r0r0r0r4rcSrkrsrrrr/cSri)Nz6Expect _random_samples to have the same dtype as inputr0r0r0r0r4rcarkcr)Nz1Expect _random samples to have 3 dimensions got, rr0)random_samplesr0r4rcercSri)Nz=Expect _random_samples.size(0) no less then input batch size.r0r0r0r0r4rcmrkcSri)Nz.rr/rcSriNzQmax_pool3d: stride must either be omitted, a single int, or a tuple of three intsr0r0r0r0r4rcrkcSriNzImax_pool3d: padding must either be a single int, or a tuple of three intsr0r0r0r0r4rcrkcSriNzJmax_pool3d: dilation must be either a single int, or a tuple of three intsr0r0r0r0r4rcrkrcSrirr0r0r0r0r4rcrkrrQrsrrzmax_pool3d_with_indices()rrr)rTrfrrrrwrrOrr#r5r@rrr)rrrr^rrTrr{r|rr}r~rrrrrrrsrrrrrrrr input_channels_last_checkrrrr0r0r4meta_max_pool3d_with_indicess            rcCs^tt|dvdd|d}t|dkr|n|d} t|dkr$|n|d} t| p2t|dvdd|s;|n|d} |sC| n t|dkrK| n|d} |sS| n t|dkr[| n|d} tt|dvdd|d}t|dkrw|n|d}t|dkr|n|d}tt|dvd d|d}t|dkr|n|d}t|dkr|n|d}t|jd vd d|d }|d }|d}|d}|d }|d}|d}t|||||| | | | | ||||||||||||d|jdkot|tjk}|jdkr| d}| o|j tjd}| |j }|r-|j tjd}|S)NrcSrirr0r0r0r0r4rcrkz7meta_max_pool3d_with_indices_backward..rr/rcSrirr0r0r0r0r4rcrkcSrirr0r0r0r0r4rcrkcSrirr0r0r0r0r4rc rkrcSrirr0r0r0r0r4rc(rkrQrsrrz"max_pool3d_with_indices_backward()rrr)rTrfrrrrrOrr#r5r@rrr)rrrrr^rrTrrr{r|rr}r~rrrrrrrrrrrrrr rrr0r0r4%meta_max_pool3d_with_indices_backwards               rgridcstjjkfddtjtjkojtjkfddtjdjdkfddtjdjdkfddtdjD]tjdkfd dqPdS) Ncr)NzNgrid_sampler(): expected input and grid to be on same device, but input is on z and grid is on r~r0rrr0r4rcdrz+check_grid_sampler_common..cr)NzTgrid_sampler(): expected input and grid to have torch.strided layout, but input has z and grid has )r{r0r r0r4rckrrcr)NzZgrid_sampler(): expected grid and input to have same batch size, but got input with sizes  and grid with sizes rr0r r0r4rcrrrrcsdjddjS)Nz+grid_sampler(): expected grid to have size rz, in last dimension, but got grid with sizes )rrr0r r0r4rcyscr)NzYgrid_sampler(): expected input to have non-empty spatial dimensions, but input has sizes rrrr0rr0r4rcr)rTrfr|r{rKrrr)rrr0)rrrr4check_grid_sampler_commonas,       r c@seZdZdZdZdZdS)GridSamplerInterpolationrr/rN)ru __module__ __qualname__BILINEARNEARESTBICUBICr0r0r0r4r sr interpolation_modecsPtjdko jjkfddtjdko |tjjk dddS)Nrcr)Nzdgrid_sampler(): expected 5D input and grid with same number of dimensions, but got input with sizes r rr0r r0r4rcs z'check_grid_sampler_3d..cSri)Nz.rrTr})rT sparse_coorfrr\r| is_sparsesparse_resize_and_clear_r sparse_dim dense_dimr _coalesced_r*rrfill_)rr\r{r|r}rr?r0r0r4rjs:   rjrzcCB|durt}|durt}|durtj}tj|||||dSrrTrget_default_devicerKrrr\r{r|r}r~r0r0r4 meta_ones  r.cCr*rr+r-r0r0r4 meta_zeros-r/r0cCrr=rOclone_preserve_strides)rrrrr0r0r4meta_select_scatterCrr3cCrr=r1)rrrrwrvstepr0r0r4meta_slice_scatterHrr5 dim_post_expr wrap_scalarcCsb|dkr |sJd}| }|d}||ks||kr'Jd|d|d|d|dkr/||7}|S)Nrr/zdim z out of bounds (rrrsr0)rr6r7rrr0r0r4rNs,rcCs|dkrdS|j|Srsrb)rrr0r0r4ensure_nonempty_sizeZsr8csttd}td}t||kddt|D]kr7tttkfddqdS)Nr/cSri)NzDIndex tensor must have the same number of dimensions as input tensorr0r0r0r0r4rcdrkz$gather_shape_check..cs$ddjdjdS)Nz!Size does not match at dimension z expected index  to be no larger than self  apart from dimension rr0rrrrr0r4rcjs)rrrTrfrr8)rrr self_dims index_dimsr0r;r4gather_shape_check_s r>csnddlm}t||}|dk}|s1tjtjkp$jtj kfddt ||| j S)Nrrcr)Nz8gather(): Expected dtype int32/int64 for index, but got rr0rr0r4rcxrzmeta_gather..) rrrrrrTrfr\rrr>rr)rrr sparse_gradr wrapped_dimis_index_emptyr0rr4 meta_gatheros    rBcCs|r*|dkrdS|dkrdS|dkrdS|dkrdS|d kr d Std d d dS|dkr0dS|dkr6dStd dd dS)Nr< REDUCE_ADDriREDUCE_MULTIPLYmean REDUCE_MEANamaxREDUCE_MAXIMUMaminREDUCE_MINIMUMFcSri)Nz=reduce argument must be either sum, prod, mean, amax or amin.r0r0r0r0r4rcrkz#get_operator_enum..addmultiplycSri)Nz/reduce argument must be either add or multiply.r0r0r0r0r4rcrkr)reduce_use_new_optionsr0r0r4get_operator_enums,rOcspddlm}||dkr"t|jtjkp|jtjkfdd|dur6t|j|jkfdddSdS)Nr)rc dS)Nz((): Expected dtype int32/int64 for indexr0r0 method_namer0r4rcrz,scatter_gather_dtype_check..crP)Nz0(): Expected self.dtype to be equal to src.dtyper0r0rQr0r4rcr)rrrrTrfr\rr)rRrrsrc_optrr0rQr4scatter_gather_dtype_checks    rTcCs t|dSr.)rrr0r0r4ensure_nonempty_dims rUc s0ddlm}|dkrdStttkddd}t}t|D]}t|}|kr:q.|t|krEd}nq.|scdurct|D]}t|}|t|krbd}nqPdurtttkddt| fdddSt| fd ddS) NrrcSriNzCIndex tensor must have the same number of dimensions as self tensorr0r0r0r0r4rcrkz%scatter_shape_check..FTcSrirVr0r0r0r0r4rcrkcs&djdjddjS)NExpected index r9r:z and to be no larger than src rr0rrrrSr0r4rcscsdjdjdS)NrWr9r:rr0)rrrr0r4rcs) rrrrTrfrUrrr8) rrrrSris_wrong_shaper<r index_d_sizer0rXr4scatter_shape_checksJ       r[cCsDt||}td|||t|||||dur t||dSdS)Nscatter)rrrTr[rO)rrrrrMrNr@r0r0r4scatter_meta_impls r]cCst||||d||jSNrKr]rrrrrrr0r0r4meta_scatter_adds racCst||||d|Sr^r]r`r0r0r4meta_scatter_add_rrccCs0t|tjr|nd}t|||||||jSr=)rlrTr r]rrrrr src_or_valuerGrr0r0r4 meta_scatters  rfcCs(t|tjr|nd}t||||||Sr=)rlrTr r]rdr0r0r4 meta_scatter_s rgqueryrr dropout_p is_causalreturn_debug_maskrc CsB|d}|d}|d} |d} |d} t|} tj||| ftj|jd} |rX| dkr3dnd}t| |}| dkrCd}n| dkrId}tj||| |f|j|jd}n tjd|j|jd}tj j rktj sqt |d krtjd tjd d}tjd tjd d}ntjdtjd d}tjd tjd d}| | dd| | |||f S) Nrr/rr-r@rr0ry)rrTrrrXr|rhceilr\versionhiprrMrrrv)rirrrjrkrlrr num_headsmax_seqlen_batch_qhead_dimmax_seqlen_batch_k attention logsumexp blocksize_c max_seqlen_k debug_maskseedoffsetr0r0r4(meta__scaled_dot_product_flash_attentionsN      r~ res_shape.cstjkrt}|StgdfddddfddD}fddttD}tj|jj d  |}|S) N)rr/rr-cs |Sr=rt)idx)rir0r4rcZrz,alloc_with_matching_layout..Trcsg|]}|qSr0r0)rJr)rr0r4rM\rz.alloc_with_matching_layout..csg|]}|qSr0rr) dim_orderr0r4rM]rNr) rerrTrsortedrrrr\r|r)rirr?permuted_shape final_permuter0)rrirr4alloc_with_matching_layoutRs   r attn_biascompute_log_sumexpc  Cs|d} |d} |d} |d} |d} | | | | f}t||}tj| | | dftj|jd}tjdtjdd}tjdtjdd}||dd| | ||df SNrr/rrrr0ryrrrTrrXr|r)rirrrrrjrkrlrrrlS_QS_KVD_Vrr? logsum_expr|r}r0r0r4(meta__scaled_dot_product_cudnn_attentiones0       rc Cs|d}|d} |d} |d} |d} || | | f} t|| }tj|| | ftj|jd}tjdtjdd}tjdtjdd}||dd| | ||df Srr)rirrrrjrkrlrrH_Qrrrrr?rr|r}r0r0r45meta__scaled_dot_product_fused_attention_overrideables0      rrrx cum_seq_q cum_seq_kmax_qmax_k philox_seed philox_offsetcCs(t|}t|}t|}|||fSr=r)rrirrrrxrrrrrjrkrrrgrad_qgrad_krr0r0r4'meta__scaled_dot_product_flash_backwards    r attn_maskc CsR|d}|d}|d} t|} tj|| |ftj|jddd} | | fS)Nrr/rr)rrTrrrXr|r) rirrrjrkrrrrsrtrwrxr0r0r40meta__scaled_dot_product_flash_attention_for_cpus"     rc CsXtj|d|j|jd} tj|d|j|jd} tj|d|j|jd} | | | fS)Nrrr/r-r)rTempty_permutedrr\r|) rrirrrrxrjrkrrrrrr0r0r49meta__scaled_dot_product_flash_attention_for_cpu_backwards& r dropout_maskcsdd}|\||\} } ||\} } j\| j\} } } fdd} fdd}dksF| krIdkrI|S| S) NcSs||dkr |ddfS|dkr:d}t|dD] }||j|9}q|||d|d|ddfS|d fS) Nr-rTrr/rsrrF)rr5rrviewr)r9rrr0r0r4 ensure_4d.s  &zBmeta__scaled_dot_product_attention_math_for_mps..ensure_4dcsj}r |}f}r<dkr'|d}||fStjdd|jdd}||}||fS)Nr-rrsr/r)rrview_asrsqueezerr)rattnr)rmax_seq_lengthnum_headq_q_sizeri unsqueezedr0r4sdpa_vector_fast_mps@s     zMmeta__scaled_dot_product_attention_math_for_mps..sdpa_vector_fast_mpscs,d}j}|f}||fS)Nrr*)blocksrr)r head_sizerrrr0r4sdpa_vector_2pass_mpsNs zNmeta__scaled_dot_product_attention_math_for_mps..sdpa_vector_2pass_mpsiir)rirrrrjrkrrrk_rRv_k_sizerrr0)rrrrrrrirr4/meta__scaled_dot_product_attention_math_for_mps#s  rcCs|dd}|dd}|dd}|d}|d} |d} |d} tj|| | | |j|jd} tjjrDtj rD |rA| nd} n |rOt | ddnd} tj|| | ftj |jd}| dd} tjdtj d d}tjdtj d d}| |||fS) Nr/rrrrrrr0ry)rrrTrr\r|rqrrrrMrhrprXr)rirrrrrjrkrrrIrsKvr? logsumexp_dimrr|r}r0r0r4,meta__scaled_dot_product_efficient_attentionZs*        rgrad_input_maskc Cs|d} |d}|d}|d}|d}|d}tj| |||fd|j|jd}tj| |||fd|j|jd}tj| |||fd|j|jd}d}|dur| dr|d}|ddkrb|n|d|d}t|}||d<tj||j|jd}|d d|f}||||fS) Nrr/rr-rrrr.)rrTrr\r|rr)rrirrrrrxrrrjrrkrrrsrru head_dim_vrrrr grad_biaslastDimlastDimAligned new_sizesr0r0r4+meta__scaled_dot_product_efficient_backwardsF              rcCs(t|}t|}t|}|||fSr=r)rrirrrrxrrrrrrrrjrkrrrrr0r0r4'meta__scaled_dot_product_cudnn_backwards    rwindow_size_leftwindow_size_right seqused_k alibi_slopescCs|dur |dn|d}|dur|dn|}|dur#|dn|}|d}|d}t|}|durFtj|||ftj|jd}n|d}tj||ftj|jd}| r|dkr_dnd}t||}|dkrod}n|dkrud}tj||||f|j |jd}n tjd|j |jd}d \}}tj j rtj rtjd tjd d}tjd tjd d}ntjd tjd d}tjd tjd d}|||||fS) Nrr/rrrrmrnroNNr0ryr)rrrTrrrXr|rhrpr\rqrrrrMrrv)rirrrrrrrjrkrlrrrrrrrtrvrsrurwrxtotal_qryrzr{r|r}r0r0r4meta__flash_attention_forwardsR     rcCs(t|}t|}t|}|||fSr=r)rrirrrrxrrrrrjrkrrrrr grad_querygrad_key grad_valuer0r0r4meta__flash_attention_backward0s    r cu_seqlens_q cu_seqlens_k max_seqlen_qrzcustom_mask_typecausal_diagonalseqlen_k window_sizecCs|d}|d}|d}|d}|d}tj|||||j|jd}|dur1|ddn|}|}|durA|dus?J|}|durG|n|}| rTt|ddnd}tj|||ftj|jd}tjdtjdd}tjdtjdd}||||||fS) Nrr/rrrrr0ry) rrTrr\r|rhrprXr)rirrr$rrrrzrjrrrrrrrrIrrsrr?logsumexp_batch_dimactual_max_seqlen_qactual_max_seqlen_krrr|r}r0r0r4!meta__efficient_attention_forwardOs,      rbias_requires_gradnum_splits_keyshared_storage_dqdkdvcCsL|rSt|jd|jdkddt|jd|jdkddtjg|jddd|jd|jdR|j|jd }|d d}|d d}|d d }nt|}t|}t|}|dur|d}|d dkrs|n|d |d }t |}||d<tj||j|jd }|d d|f}ntjd|jd}||||fS)Nr/cSri)Nz,seqlen must match for `shared_storage_dqdkdvr0r0r0r0r4rcrkz4meta__efficient_attention_backward..r-cSri)Nz3embedding dim must match for `shared_storage_dqdkdvr0r0r0r0r4rcrkrrrrrsrr.r0r~) rTrfrrr\r|rrrr)rrirrr$rrrrzrxrjrrrrrrrchunkrrrrrrrr0r0r4"meta__efficient_attention_backwards:*         rscale_ascale_b scale_resultuse_fast_accumcsdd}tdkodkfddt|jo$|jfddtdks9tdkrd d } d d } d d} t| pP| fddt| pb| fddtdddkfddtdddkodddkfddj\djtjkojtjkpjtj kojtj k}  dkrֈ dkrtjtj koψjtj kddn| r9jtj krd} dnd} d}t | }t |dd}|t |||t || kr+ kr+t ddt ddntdfddntjtj koGjtj kddtdkoZdkfdddkrddkrddkrdkrt o d dnadkrddkrt dkrnn dt dkrn6dkrddkrt dkrnn dkrntdfd!d|dur|nj}tjdd|jd"S)#NcS|tjtjtjtjtjfvSr=rTr? float8_e5m2float8_e4m3fnuzfloat8_e5m2fnuzfloat4_e2m1fn_x2rr0r0r4is_fp8_or_fp4_typez2_check_scaled_mm_sizes..is_fp8_or_fp4_typercddSNz%Inputs must be 2D but got self.dim()=z and mat2.dim()=rr0r0rr0r4rcr_z(_check_scaled_mm_sizes..crNz?Expected both inputs to be fp8 or fp4 types but got self.dtype=z and mat2.dtype=rr0rr0r4rcrrrcS|d|dko |ddkSrsr0rtr0r0r4 is_row_majorz,_check_scaled_mm_sizes..is_row_majorcS|ddko |ddkSrsr0rtr0r0r4 is_col_majorrz,_check_scaled_mm_sizes..is_col_majorcS|ddkp |ddkSrsr tensor_2dr0r0r4 has_zero_dimrz,_check_scaled_mm_sizes..has_zero_dimcrNz#self must be row_major, got stride rtr0rr0r4rcrcrNz#mat2 must be col_major, got stride rtr0r0r0r4rcrr/rrcddSNzBExpected self.size(1) to be divisible by 16, but got self.size(1)=r/rr0rr0r4rcrdcrNz?Expected both dimensions of mat2 to be divisible by 16 but got rr0rr0r4rcrcSri)NzNFor tensorwise scaling, both scale_a and scale_b must be float (fp32) tensors.r0r0r0r0r4rc rkrrnrcSri)Nzscale_a must be contiguousr0r0r0r0r4rc+rkcSri)Nzscale_b must be contiguousr0r0r0r0r4rc/rkFc s&ddddd S)NzTInvalid blockwise scaling configuration. For blockwise scaling, scale_a should have  elements, got z, scale_b should have rrrr0)expected_a_sizeexpected_b_sizerrr0r4rc4scSri)NzKFor rowwise scaling, both scale_a and scale_b must be float (fp32) tensors.r0r0r0r0r4rc=rkcsddS)NzLFor non-tensorwise scaling, scale tensors must be 2D, but got scale_a.dim()=z and scale_b.dim()=rr0rrr0r4rcBr_cSri)Nz@Both scale_a and scale_b must be contiguous for rowwise scaling.r0r0r0r0r4rcNrkcsddddtdd dtddtdddtdd dtddd d dd d d dd d SNz}Invalid scaling configuration. For tensorwise scaling, both scales should be scalar. For rowwise scaling, scale_a should be (z, 1), scale_b should be (1, z>). For (BlockWise1x128, BlockWise128x128), scale_a should be (rrrnz), zscale_b should be (z<). For (BlockWise1x128, BlockWise1x128), scale_a should be (z). Got scale_a.size()=(rr/z) and scale_b.size()=(rsr5rr0)_krBrrrr0r4rcbs2r)rTrfrr\rrrrfloat8_e8m0fnur?rr<r5r@rr|)rr0rrr$rr&rrrrris_blockwise_scaling block_size_k block_size_mn num_k_blockspadded_num_k_blocks _out_dtyper0) rrrrBr0rrrrr4_check_scaled_mm_sizess       "             .. rc Cst||||||||Sr=)r)rr0rrr$rr&rr0r0r4meta_scaled_mmss rscale_recipe_ascale_recipe_b swizzle_a swizzle_bc  sz dd} dd} t dkodk fddt| jo(| j fdd jd  jd jd | jrO| jrOd} | 9d d |D}d d |D} rgdd D ntjg rudd Dntjgt dkst dkrdd}dd}dd}t| p|  fddt|p|fddt d dd k fddt d dd koڈ d dd kfdddt t dt t fdd }dt t dt t fd!d"}dt t dt t fd#d$}dt t dt t fd%d&}dt t dt t fd'd(}dt t dt t fd)d*}dt t dt t fd+d,}|||rmt d  d kof d  d kof d jtj kof d jtj kd-dn<|||rt d jd ko d  ko d jtj ko d  ko d jtj k fd.dn|||rD d jtj koÈ d jtj k} d jd kojd d/kod d kod kpjd d kod d k} d  jd ko- jd d/ko- d d ko- d kp- jd d ko- d d k}t|o7|o7| fd0dne|||r d jtj koZ d jtj k}td/d1 d jd kojd d/kod d kod kpjd d kod d k} d  jd koˈ jd d/koˈ d d koˈ d kpˈ jd d koˈ d d k}t|o|o| fd2dn|||r d jtj ko d jtj k}td/d1 d jd ko5jd d/ko5d d ko5d kp5jd d ko5d d k} d  jd koj jd d/koj d d koj d kpj jd d koj d d k}t|ot|ot| fd3dn'|||rtjjrt jd d4 jd t jd d4 jd tjn)t jd d/tt jd d4d1tjd d/tt jd d4d1tjt d  ko d jtjko d  ko d jtjko d kod k fd5dn|||rtd/ttdd1td/ttdd1tjt d  ko d jtjko d  d ko d jtj ko d  ko d jtjko d  d ko d jtj ko d kod k fd6dntd7 fd8d|dur|n j}tj| jd9S):NcSrr=rrr0r0r4rrz5_check_scaled_mm_sizes_v2..is_fp8_or_fp4_typecSs |tjfvSr=)rTrrr0r0r4 is_fp4_typer z._check_scaled_mm_sizes_v2..is_fp4_typercrrrr0rr0r4rcr_z+_check_scaled_mm_sizes_v2..crrrr0rr0r4rcrrr/cSg|]}t|qSr0r%rJsir0r0r4rMrz-_check_scaled_mm_sizes_v2..cSr r0r rr0r0r4rMrcSr r0r&rr0r0r4rMrcSr r0rrr0r0r4rMrrrcSrrsr0rtr0r0r4rrz/_check_scaled_mm_sizes_v2..is_row_majorcSrrsr0rtr0r0r4rrz/_check_scaled_mm_sizes_v2..is_col_majorcSrrsrrr0r0r4rrz/_check_scaled_mm_sizes_v2..has_zero_dimcrrrtr0rr0r4rcrcrrrtr0rr0r4rcrrcrrrr0rr0r4rcrdcrrrr0rr0r4rcrrecipe_arecipe_bcS4t|dkot|dko|dtjko|dtjkSr)rr% TensorWiserrr0r0r4 is_tensorwise    z0_check_scaled_mm_sizes_v2..is_tensorwisecSrr)rr%RowWiserr0r0r4 is_rowwiserz-_check_scaled_mm_sizes_v2..is_rowwisecSrr)rr% BlockWise1x32rr0r0r4is_mxrz(_check_scaled_mm_sizes_v2..is_mxcSsPt|dko't|dko'|dtjko'|dtjko'|dtjko'|dtjkS)Nrrr/)rr% BlockWise1x16rrr0r0r4is_nvs      z(_check_scaled_mm_sizes_v2..is_nvcSrr)rr%BlockWise1x128rr0r0r4is_1x128_1x128rz1_check_scaled_mm_sizes_v2..is_1x128_1x128cS4t|dkot|dko|dtjko|dtjkSr)rr%rBlockWise128x128rr0r0r4is_1x128_128x128rz3_check_scaled_mm_sizes_v2..is_1x128_128x128cSr r)rr%r!rrr0r0r4is_128x128_1x128rz3_check_scaled_mm_sizes_v2..is_128x128_1x128cSri)Nz\For Tensorwise scaling, both scale_a and scale_b must be single element float (fp32) tensorsr0r0r0r0r4rc rkc s:djddddjdddd S)Nz'For Rowwise scaling, scale_a must have rz elements (got: z), and scale_b must have r/rs)rrr0)r0rrrr0r4rc)s  rncsRddddjddjddddjddjdS)Nz>For 1x128 x 1x128 blockwise scaling, scale a must have shape [rrrn] (got: ) and stride [1, )scale b must have shape [rsrrr0)rCrIrsasbr0r4rcF*rcsRddddjddjddddjddjdS)Nz]For 128x128 x 1x128 blockwise scaling, L4 = {round_up(K / 128, 4)}, scale a must have shape [rrrnr$r%r&rsr'r0rCL4rIrr(r)r0r4rcer*csRddddjddjddddjddjdS)Nz]For 1x128 x 128x128 blockwise scaling, L4 = {round_up(K / 128, 4)}, scale a must have shape [rrrnr$r%r&rsr'r0r+r0r4rcr*rcshdddddddtjddjddjdd dddd S) Nz!for MX scaling scale_a must have  (got: r) and scale_b must have z). Scales must have types z (for self: , mat_b: z) Must have swizzle type (got self: rs)rrTrr\r0expected_scale_a_elemsexpected_scale_b_elemsexpected_swizzlerrr r r0r4rcs" csHddddddddddddS) Nz!for NV scaling scale_a must have r-rr.z). Must have swizzle type r0r/rsrrr0r1r0r4rcs Fcsddddtdd dtddtdddtdd dtddd d d dd d d d d dd d d Srrr0)rCrIrrrr0r4rcs2    r)rTrfrr\rr& NO_SWIZZLErrrrr%rr<r;rqrrr5SWIZZLE_32_4_4rr?rr|)rr0rrrrr$r&r r rrr K_packed_multiplierrrrrrrrrr"r#types_ok scale_a_ok scale_b_okrr0)rCr,rIrr2r3r4r0r(r)rrrr r r4_check_scaled_mm_sizes_v2s          "      , ,   , ,   , ,       r;r9contraction_dimsc Cst|||||||| ||| d S)N)r$r&r r r)r;) rr0rrr rrr r$r9r<rr0r0r4meta_scaled_mm_v2sr=cCs t|||||dd||jSNT)rNr_rrrrrGrEr0r0r4meta_scatter_reduce_twos r@cCst|||||dd|Sr>rbr?r0r0r4meta_scatter_reduce__twosrAcshtdko dknfdddkr&tj|tjjdStjd|tjjdS)Nrrcr)NzAThe probability distributions dimensions must be 1 or 2, but got rr0rr0r4rcrz"meta_multinomial..r/r)rTrfrrrr|r)r num_samples replacementrr0rr4meta_multinomials  rDcCsd}|D]}||9}q|Sr.r0)vsr vr0r0r4multiply_integerss rGcsttkfdddttkfddttddddDo9tddDfdddd\}}||gRS)NcddtS)Nz%It is expected output_size equals to , but got size r]r0)num_spatial_dimsrr0r4rcrz'upsample_common_check..rcrH)Nz$It is expected input_size equals to rIr]r0)expected_input_dimsrr0r4rcrcs|]}|dkVqdSrNr0)rJr{r0r0r4ror1z(upsample_common_check..crp)NzDInput and output sizes should be greater than 0, but got input size z and output size r0r0)rrr0r4rcs )rTrfrr6)rrrJrschannelsr0)rKrrJrr4upsample_common_checks    * rOcsZtdkptddfddt|dd}|jt dS)Nrr/cr)Nz>Non-empty 3D data tensor expected but got a tensor with sizes rr0rr0r4rc*rz$upsample_nearest1d..rJr rTrfrrGrrOrrrOr)rrscalesfull_output_sizer0rr4upsample_nearest1d$   rTc stdkptddfddt|dd}|}t}j \}}}}j j dkr?|dkr?tj }|j |d }|S) Nrr/crNz>Non-empty 4D data tensor expected but got a tensor with sizes rr0rr0r4rc:rz$upsample_nearest2d..rrPrrr)rTrfrrGrrOrrOrrr|rtr contiguous) rrscales_hscales_wrSrrrR n_channelsr0rr4upsample_nearest2d4s     r[rrrXrYcstt||ddtjdkfddtdD]tkfddq|jt dS)NrrPrcrNzFExpected grad_output to be a tensor of dimension 4 but got: dimension rr0rr0r4rc`rz-upsample_nearest2d_backward..c &ddddS)NzCExpected grad_output to have the same shape as output; output.size() = z but got grad_output.size(rr0rSrrr0r4rcesr) rOrTrfrrrrrrOr)rrrrXrYr0r_r4upsample_nearest2d_backwardNs    r`csZtdkptddfddt|dd}|jt dS)Nrr/cr)Nz>Non-empty 5D data tensor expected but got a tensor with sizes rr0rr0r4rcwrz$upsample_nearest3d..r-rPrrQ)rrscales_drXrYrSr0rr4upsample_nearest3dqrUrbc Cst|tj|tjd}}|durQ|durQt|tsJt|ts$J|j}|} t||}t||}||| ||| t ||dt ||d||fS||fS)Nr)r2r3) rTrrrlrrrrrr ) rstabler descendingrrrFrr out_strider0r0r4 meta_sorts       rfcstjdkfddtjjkfddddurPtjdkfddtkfddtjjkfddtjdkfd dd tkfd dttfd d fDdddS)Nrc jdSNz != 2rr0 input_gatesr0r4rcrz%rnn_cell_checkSizes..cjdjSN != rr0) hidden_gatesrjr0r4rcrr/crg)Nz != 1rr0) input_biasr0r4rcrcsdSrlrrr0) gates_sizeror0r4rcrcrkrlrr0) hidden_biasror0r4rcrcrgrhrr0) prev_hiddenr0r4rcrrc s,dddddd S)Nrmrz * z // z (aka rs)rrr0)expected_prev_hidden_numelfactorrprjrrr0r4rcs,c3s|] }|jjkVqdSr=r~rIrir0r4ros  z&rnn_cell_checkSizes..cSri)Nz%expected all inputs to be same devicer0r0r0r0r4rcrk)rTrfrrrrr6)rjrnrorqrtrrr0)rsrtrprqrnrorjrrr4rnn_cell_checkSizess8          rucCsLt||||d|tj|tjd}tj|tjd}tj|tjd}|||fS)Nrr)rurTrr)rjrncxrorq workspacehycyr0r0r4_thnn_fused_lstm_cell_metas  rzcCs(t|dk}|rt|}|d}|jd}n| r|jdn|jd}| r)|jdn|jd}d}| r4dnd}|dkr<|n|}|rG|||g}n| rP||||gn||||g}||}| |||g}|durptjd|jd}n||}|| |||g}| rdnd}|j|tjd}|||||fS)Nrr/rrr~r)rrrrTrr|r1)rr"weight_stride0 weight_bufhxrvr hidden_size proj_size num_layers batch_firstdropouttrain bidirectional batch_sizes dropout_stateis_input_packed seq_length mini_batchbatch_sizes_sumnum_directionsout_sizerr cell_shaperyrx reserve_shapereserver0r0r4 _cudnn_rnns2       rcCs|r|jdn|jd}|r|jdn|jd}| }|r!|||gn|||g}||}|dur8tjd|jd}n||j}|durKtjd|jd}n||j}tjd|jtjd}||||fS)Nr/rr~r)rrrTrr|r1)rw0w1w2w3hx_cx_rrrr~r has_biasesrrrrroutput_chanelsrrrxryrwr0r0r4mkldnn_rnn_layer s      rcsT|jdkrtdkpdkfdddSt|dkfdddS)Nrrcr)Nz4: Expected reduction dim -1 or 0 for scalar but got r0r0rrr0r4rc7rz'zero_numel_check_dims..cr)Nz: Expected reduction dim z to have non-zero size.r0r0rr0r4rc<rd)rrTrr)rrrr0rr4zero_numel_check_dims3s    rcsF|durt||}t||dSt|dkfdddS)NrcrP)Nz@: Expected reduction dim to be specified for input.numel() == 0.r0r0rir0r4rcHrz%check_argmax_argmin..)rrrrTrfr)rrrr0rir4check_argmax_argminAs  rcCsDtd||t|j|dur|fnd}t|||}|j|tjdS)Nargmaxr)rrOrerrfrrTr)rrrhrrr0r0r4argmax_argmin_metaLs  rcCs$|tjkrtj}tjd||||dS)Nr0r)rTjaggedrKr)r{r\r{r|r}r0r0r4 scalar_tensorTs  rcCst||dd}|dkrdn||}t|dkt||kddt|j}t|dkr6|||<|||j|tj dfS)NTr7rr/cSri)Nzk not in range for dimensionr0r0r0r0r4rcfrkztopk_meta..r) rrrrTrfrrrrr)rrrlargestr sliceSizetopKSizer0r0r4 topk_meta`s  rc Cs@|dus |dus Jd|}|} tj|| j| j| jdS)Nz;segment_reduce(): Either lengths or offsets must be defined)r\r|r{)rWrTrr\r|r{) rrrVrGrQrRrSrU data_contig grad_contigr0r0r4meta__segment_reduce_backwardnsrcsddlm}t|dd|dkr|nd}t||dk||kfddt|jd|jdd}|rM|dkrM| d| ||j |tj dfS) Nr)rTrr/cr)Nz9kthvalue(): selected number k out of range for dimension r0r0rr0r4rcrzkthvalue_meta..r) rrrrrrTrfrrrrr)rrrrhrdimSizerr0rr4 kthvalue_metas  $ rcCs|dur|n|}t|dkdd|}|dur(t||kdd|dur8t||kddt||kddt||kddt|dkddt||d |d d kd ddS) NrcSriNr0r0r0r0r4rcrkz(checkLSTMBackwardSizes..cSrirr0r0r0r0r4rcrkcSrirr0r0r0r0r4rcrkcSrirr0r0r0r0r4rcrkcSrirr0r0r0r0r4rcrkcSrirr0r0r0r0r4rcrkrr/rcSrirr0r0r0r0r4rcrk)rTrfrrr)grad_hygrad_cyrvryrw defined_gradexp_sizer0r0r4checkLSTMBackwardSizess,rc Cs`|dur |dur dSt|||||tj|td}tj|td}|r)|jdddnd}|||fS)NNNNrrF)rh)rrTrlegacy_contiguous_memory_formatr<) rrrvryrwhas_bias grad_gatesgrad_cxrr0r0r4#_thnn_fused_lstm_cell_backward_impls rcCsfd}d}d}|dr||}|ds|dr.||d|df}||d}|||fS)Nrr/rrr)rrrrr grad_weightrr0r0r4linear_backwards rcstjdkrjd||dksJdjd|ddfdd }jd||}jd |}jd |}gjdd|||R}|}|j|d }|S) Nrrsrz'Invalid input shape for pixel_shuffle: z with upscale_factor = cSr r=r r r0r0r4rrz,meta_pixel_shuffle..is_channels_lastcsLrtdkr tjStjSjtjdrtjSjtjdr$tjSdS)Nrr)rrTrr r@preserve_formatr0rrr0r4rs z.meta_pixel_shuffle..pick_memory_formatrrr)rrrr)rupscale_factorrrHrWrrrr0rr4meta_pixel_shuffles&   rcCsZ||j}||j}||j}||j}||j}||j}|||||||fSr=r*)rweight0weight1weight2weight3rcx_tmprhy_cy_grad_output_r_opt grad_hy_r_opt grad_cy_r_optrrr~rrrrrrrwdiff_xdiff_hxdiff_cxdiff_w1diff_w2diff_br0r0r4mkldnn_rnn_layer_backwards      r) out_int32rcCstj||rtjntjtjdS)Nr\r)rTrr)rr)r boundariesrrr0r0r4meta_bucketizes rdcsdtdkrtfddtdkr#r#tdtttfddtdkfd dtttfd dtttfd dtkfd dtj j j d S)Nzhistc()rzcr)Nz%"histogram_cpu" not implemented for 'r~rr0rr0r4rcrzmeta_histc..rz%_histc_cuda with floating point inputcsdtS)Nz#: argument 'bins' must be int, not rr0binsrr0r4rcrrcr)Nz: bins must be > 0, but got r0r0rr0r4rcrcdtS)Nz%: argument 'min' must be Number, not rr0)rrr0r4rcrcr)Nz%: argument 'max' must be Number, not rr0)rrr0r4rc#rcrP)Nz: max must be larger than minr0r0)rr0r4rc%rr) rrTrfrrOrrlrrrr|r\)rrrrr0)rrrrrr4 meta_histc s.      rcsdt|dd}tdkptddddDfdd|jt d S) NrrPrcsrLrMr0)rJrr0r0r4ro7r1z,meta_upsample_bimode2d_aa..r/crrVrr0rr0r4rc8rz+meta_upsample_bimode2d_aa..r) rOrrTrfrr6rrrOr)rrrrXrYrSr0rr4meta_upsample_bimode2d_aa)s (  rcstt||ddtjdkfddtdD]tjkfddq|jt dS)NrrPrcrr\rr0rr0r4rcMrz4meta_upsample_bimode2d_aa_backward..c r])NzD Expected grad_output to have the same shape as output; output.size(r^z but got grad_output_size(rr0r_r0r4rcRsr) rOrTrfrrrrrrOr)rrrrrXrYr0r_r4"meta_upsample_bimode2d_aa_backward?s    rcCs\t|dkddt|dkddt|jjddt|jjdddS)Nr/cSri)Nz%found_inf must be a 1-element tensor.r0r0r0r0r4rc_rkz<_amp_foreach_non_finite_check_and_unscale_..cSri)Nz%inv_scale must be a 1-element tensor.r0r0r0r0r4rcbrkcSri)Nz!found_inf must be a float tensor.r0r0r0r0r4rcfrkcSri)Nz!inv_scale must be a float tensor.r0r0r0r0r4rcjrk)rTrfrr\r)rr inv_scaler0r0r4*_amp_foreach_non_finite_check_and_unscale_\srcCrr=r)rnanposinfneginfr0r0r4 nan_to_numorarcCs|jtjtjtjtjhvsJd|jd|j}t||}t||}||kr)|St| }t| }||||||<||<||||||<||<| |||S)Nz>torch.transpose_: in-place transposition is not supported for z layout) r{rTrL sparse_cscrM sparse_bscrrrrrr)rdim0rndimsrrr0r0r4rDus&      rDcCsz|j}|jr"|}|}|dkr|dks!Jd|d|dn|dks0Jd|dt|d|dkr:dSdS) NrrzEt_ expects a tensor with <= 2 sparse and 0 dense dimensions, but got z sparse and z dense dimensionsz6t_ expects a tensor with <= 2 dimensions, but self is rZr/)rr$r&r'rrD)rrr&r'r0r0r4t_s r)rrsidesortercsttjdkpjddjddkfddtdup)jjkfddt|dkp9| dd|rCtjntj}ttjrUtj|tj dStj d |j d S) Nr/rcsdtjdtjS)Nztorch.searchsorted(): boundaries tensor should be 1 dimension or the first N-1 dimensions of boundaries tensor and input value tensor must match, but we got boundaries tensor z and input value tensor rrr0)rsorted_sequencer0r4rcs z#meta_searchsorted..cs,dtjddurtjSgS)Nz[torch.searchsorted(): boundary and sorter must have the same size, but got boundary tensor z and got sorter tensor rr0)rrr0r4rcsrcSri)Nzetorch.searchsorted(): side and right can't be set to opposites, got side of left while right was Truer0r0r0r0r4rcrkrr0r) rTrfrrr)rrlr rrrr|)rrrrrrr\r0)rrrr4meta_searchsorteds&      rcs(ttjtjtjfvfdddS)Ncr)Nz/Unsupported input type encountered for isin(): r0r0rr0r4rcrz3_check_for_unsupported_isin_dtype..)rTrfr complex128 complex64rr0rr4!_check_for_unsupported_isin_dtypes rcCs|||df}|S)Nrr)rr num_weightsrvr~rr0r0r4meta_embedding_dense_backwardsrc Cs:| rt||||||||| | St||||||||| | Sr=)r*_embedding_bag_sparse_backward!meta_embedding_bag_dense_backward) rrrRrrmaximum_indicesrr~rrrqrvr0r0r4meta_embedding_bag_backwards2 rc sXtjtjtjtjtjfvfdd|tkr t|du| df} | S)Ncr)Nz$Unsupported input type encountered: rr0rr0r4rc$rz3meta_embedding_bag_dense_backward..r/) rTrfr\r=r>r<float64r|rr) rrrrrrr~rrqrvindex_grad_weightr0rr4rs  rc Cs|d}t|tkddt|dkt|dk|d}t|dkt|d|k||f} | S)Nr/cSri)NzHembedding_bag_backward: per_sample_weights only supported for mode='sum'r0r0r0r0r4rc9rkz@meta_embedding_bag_per_sample_weights_backward..rr)rrTrfr{rr) rr"rrRrrrvembedding_featuresrBrr0r0r4.meta_embedding_bag_per_sample_weights_backward,s   r) assume_uniqueinvertcCsxtt|tp t|tddt|tstj||jd}t|ts*tj||jd}t|jt|jtj|tj dS)NcSri)Nz.r~r) rTrfrlr rr|rr\rr)elements test_elementsrrr0r0r4 meta_isinDs    rrcCs4t|dkddt|tjd\}}tj||dS)NrcSri)Nz,polygamma(n, x) does not support negative n.r0r0r0r0r4rcYrkz meta_polygamma..rnr)rTrfrrror)rrrRrLr0r0r4meta_polygammaVs  rcCstd)Nz.Tensor.item() cannot be called on meta tensors)rrr0r0r4meta_local_scalar_denseasrcCrr=rrr0r0r4silufrarcCs t|tjd\}}tj||dSrm)rrrorTr)rrRrLr0r0r4sigmoidls  rc Csp|dk}|dk}|r:|r|d|d|dg}nVt|d|dkdd|d|dg}n;|rWt|d|dkdd|d|dg}nt|d|dkdd|d|d|dg}|py|j}tjjrd |j}|d|d||}||kr|d||dg} n|dg} tj|| ||j d } | Stj |||j d } | S) Nrrr/cSriNz matrix batch sizes have to matchr0r0r0r0r4rcrkz2_create_grouped_mm_output_tensor..rcSrirr0r0r0r0r4rcrkcSri)Nzbatched dimension has to matchr0r0r0r0r4rcrkrr) rrrTrfr\rqritemsizerr|r) r.r0offsr& mat1_is_2d mat2_is_2dr alignment size_paddedrerr0r0r4 _create_grouped_mm_output_tensorvs>     r mat_amat_brc stdudukddduodu} | r6tjjrtjntj} tj| ko-j| kfddntjtjkoCjtjkfddtdvoXdvfdddk} dk} | rp| st d d kd d| rd d } d d}t| fddt|fdddd}|d|ddurdurtjtj krȈjtj kpӈjtj koӈjtj kfddjtj kojtj kd#fdd }dur| r| rj dnd}|dd||dd|t|dudd| s| rLtdufdddurKtdkfddtjtj kfddn tdud dt|dud!dt|dupk|tjkd"dt|S)$NcSri)Nz,Either both scale factors are given, or noner0r0r0r0r4rcrkz)_meta_grouped_mm_common..cr/)Nz5Expected inputs of E4M3 FP8 type but got mat_a.dtype= and mat_b.dtype=rrr0r r r0r4rcrNcr/)Nz1Expected inputs of BF16 type but got mat_a.dtype=r rrr0rr0r4rcrN)rr-csddS)Nz3Multiplicands must be 2D or 3D but got mat_a.dim()=z and mat_b.dim()=rr0rr0r4rcr_rrrcSri)Nz3contraction dimension of mat_a and mat_b must matchr0r0r0r0r4rcrkcSs |}|ddko|ddkSNrr/rrtmat mat_strider0r0r4rz-_meta_grouped_mm_common..is_row_majorcSs |}|ddko|ddkSrrtrr0r0r4rrz-_meta_grouped_mm_common..is_col_majorcdddS)NzNExpected mat_a tensor to be row major in the last two dimensions, got strides rrtr0)r r0r4rcrNcr)NzQExpected mat_b tensor to be column major in the last two dimensions, got strides rrtr0)r r0r4rcrNcsdd}ddkr:tdjdkr:t|dkfdddSdkrddtdjkrdtd|dkfdddStdfdddS) Nr/rrcsddddS)Nr` stride along % dim to be multiple of 16 bytes, got rr0r0end_dimmat_namerr0r4rcrzF_meta_grouped_mm_common..check_valid_strides..cs$ddddddS)Nr`rr/rrr0r0rr0r4rc$FcsddjdS)NzInvalid strides/sizes, got z for strides and z for sizes.rr0rr0r4rcr)r element_sizerrrrTrf)rrrr0)rrrrr4check_valid_stridess*    z4_meta_grouped_mm_common..check_valid_stridesr r cr/)NzhFor FP8 scales must both be float32, or for MXFP8 both scales must be float8_e8m0fnu. Got scale_a.dtype=z and scale_b.dtype=rrr0rr0r4rcrNr/csdkrQtfddr(tkfdddStdkfddtjdjkfdddStd dkfd dtjdjdkfd drtjjdkfd dj\}}d }t||dt|dtjdkojdkfdddStdkfddtjdjdkfdddS)Nrcr )Nr`z to be contiguous.r0r0 scale_namer0r4rcrz>_meta_grouped_mm_common..check_scale..cddjdjS)NzKFor MXFP8, scale must have same number of dimensions as target tensor, but  has mat.ndim= and scale.ndim=rr0rrrr0r4rcr/cdddS)Nr`z to be 1D tensor, but got D tensor.rr0rrr0r4rcrNrcs(ddjdjddS)Nr`z to have rrz elements.rr0)rrscale_multiplierr scaled_dimr0r4rc!(rcr )Nr`z( to be contiguous in the last dimension.r0r0rr0r4rc&rcs$ddjddjddS)Nr`z batch dimension to be r, got rrr0r"r0r4rc*rcr)Nz0For MXFP8, 3d tensor should have 2d scales, but r r!rr0r"r0r4rc1r#rrrncs$djdddjS)NzFor MXFP8, expected mat.shape=z to have scale shape of (,z ), but got rr0)G blocked_K blocked_Nrrr0r4rc;rcr$)Nr`z to be 2D tensor, but got r%rr0r&r0r4rc@rNcs(ddjddjddS)Nr`z non-batch dimension to be r/r*rrr0)rrrr(r0r4rcDr))rrTrfr@rrrr;)rrrr(r'rCrrO)is_mxfp8)r,r-r.rrr'rr(r4 check_scale s^          z,_meta_grouped_mm_common..check_scalerrrcSri)Nz:Scale result tensor provided, but it is not supported yet.r0r0r0r0r4rcOrkcsdddS)Nz/Offsets tensor not provided, but is needed for zD/zD multiplicand layouts.rr0rr0r4rcUr#crY)Nz.Offsets tensor must be 1D, but got offs.dim()=rrr0rr0r4rcZrdcr)Nz7Offsets tensor must be integer (int32) tensor, but got rrr0r1r0r4rc^rcSri)NzJOffsets tensor provided, but is not needed for 3D/3D multiplicand layouts.r0r0r0r0r4rccrkcSri)Nz2Bias tensor provided, but it is not supported yet.r0r0r0r0r4rchrkcSri)Nz4If output dtype provided, it must be torch.bfloat16.r0r0r0r0r4rcmrkrO)rTrfrqrrrr?r\r>rrr<rrr)r )r r rrrr$rr&rscaled fp8_dtype mat_a_is_2d mat_b_is_2drrrr0r'r0)r/r r rrrr4_meta_grouped_mm_commons              =       r6c Cst||dd||d|dS)N)rrrr$rr&)r6)r r rr$r&r0r0r4meta_grouped_mmss r7c Cs$|ptj}t|||||||||d S)N)rrrr$rr&r)rTr>r6) r r rrrr$rr&rr0r0r4meta_scaled_grouped_mms r8r9 half_to_floatcCsR|r |jtjtjfvs Jtj|tjjd\}}|s|n|}tj||tj d}|S)Nrnr) r\rTrVr>rOrrrPrr)r9rr9computation_dtyperLr?r0r0r4softmaxs r;c sttddkfdd|jttd}|t|kfddtddDr||}tD]9dddkr_| |j}ddkrw|d|jd}q>|Std}t|D]1tddd}t|dkfd d| |qtj ||j |j |j t|d S) Nrrcr)Nz1Length of pad must be even but instead it equals r]r0rr0r4rcrz'_constant_pad_nd_meta..csdtddS)Nz`Length of pad should be no more than twice the number of dimensions of the input. Pad length is z while the input has z dimensions.r]r0)l_inprr0r4rcs css$|] }t|tjo |dkVqdSrM)rlrOIntWithoutSymInt)rJrr0r0r4ros"z(_constant_pad_nd_meta..r/c s6dddddd S)NzThe input size z, plus negative padding rr/zG resulted in a negative output size, which is invalid. Check dimension z of your input.r0r0)rrl_diffrpad_idxr0r4rcs )r\r|r~r)rTrfrrr6rnarrowrrrrr\r|r~r)rrrl_padc_input new_shapenew_dimr0)rrr>r<rr?r4_constant_pad_nd_metasP        rErvr~rc Csx|dks Jd|j}|j}|jdkr|df}n|jdkr)|d|df}n g||dR}|j}|j||dS)Nrz'weight' must be 2-Drr/r)rrrr\r) r"rrvr~r weight_shape indices_shaperr&r0r0r4 embeddings    rH max_lengths padding_valuecCs\t|dksJt|dksJ|djdd}|d}||g|jddR}||Sr)rrr)rrRrIrJrr rCr0r0r4$meta__jagged_to_padded_dense_forward s  rKcCt|tdd}|S)NcSrrrSrrorr0r0r4_f sz)_create_unary_float_meta_func.._frFr!funcrNr0r0r4_create_unary_float_meta_func  rRc Cs|js |js |jr td|dkr||j|dfS|d} |d}|d}|| ||}| rO| rC|| ||}||fS|| |||}||fS|d}||fS)NzP_native_multi_head_attention fake implementation does not support nested tensorsrr/) is_nestedr[rrrr)rirr embed_dimr qkv_weightqkv_bias proj_weight proj_biasr need_weightsaverage_attn_weights mask_typerT output_dimr attn_weightsr0r0r4 native_multi_head_attention_fake s$     r`cCrL)NcSrrrMr8r0r0r4rNR rz*_create_binary_float_meta_func.._frOrPr0r0r4_create_binary_float_meta_funcQ rSracs<tfdd}jd}||_ttt||}|S)Ncs(|g|Ri|}t|j|j|Sr=r)rrQrrr@r0r4_fnv sz#_register_inplace_meta.._fnrR)rrurFgetattrr*)rArb inplace_namer0r@r4_register_inplace_metau s  recsrtjjkfddg}ttr1jdkr,tjjkfdd|t|dtj iS)Ncr)Nrz for `end`, but got dtype rr0)rvrwr0r4rc rzlerp..rcrZ)Nrz for `weight`, but got dtype rr0)rwr"r0r4rc rrG) rTrfr\rlrrrrSrrP)rwrvr"rQr0)rvrwr"r4lerp s"       rf)rcCst|||tjdSrrrtensor1tensor2rr0r0r4addcmul s rjcCs8tt|jo t|j ddt|||tjdS)NcSri)N)zFInteger division with addcdiv is no longer supported, and in a future zErelease addcdiv will perform a true division of tensor1 and tensor2. z4The historic addcdiv behavior can be implemented as zA(input + value * torch.trunc(tensor1 / tensor2)).to(input.dtype) zfor integer inputs and as z6(input + value * tensor1 / tensor2) for float inputs. z?The future addcdiv behavior is just the latter implementation: z4(input + value * tensor1 / tensor2), for all dtypes.r0r0r0r0r4rc rkzaddcdiv..r)rTrfrOrr\rSrrPrgr0r0r4addcdiv s   rkcCs4i}dD]}t|}|D] }||vr||||<q q|D]y\}}t|tjjr*qt|ts1J|tjj j |tj | drR|tdvrQt |dq|jrVq| dvr]qd| vrjt||qd| vrwt||qd| vrt||qd | vrt||qt||qdS) N)ry post_autograd pre_autogradCompositeImplicitAutogradryz is a CompositeImplicitAutograd op, we shouldn't register meta function for it. Instead, we should let the decomposition run and write meta kernels for the base operators.> aten::clone aten::copy_ aten::rot90aten::_to_copyaten::empty_stridedaten::constant_pad_ndaten::as_strided_scatterzmkldnn::zmkl::zonednn::z quantized::)r itemsrlrT_opsHigherOrderOperatorrpy_impl_C DispatchKeyr,%_dispatch_has_kernel_for_dispatch_keyrris_view2_meta_lib_dont_use_me_use_register_meta_for_mkldnnimpl/_meta_lib_dont_use_me_use_register_meta_for_mkl2_meta_lib_dont_use_me_use_register_meta_for_onednn5_meta_lib_dont_use_me_use_register_meta_for_quantized'_meta_lib_dont_use_me_use_register_meta)activate_meta_tablertregistryopo op_overloadrAr0r0r4 activate_meta sN      rrrr=)NNNFrr/rr)Tr)r)rT)FF)TT)r)FTN)TFF)TF)r)rN)r<r)r0rrOF)r0rFTN)FrFNFr)NF)rF)rrFN)NNNNN)rNNr/)NNF)rhFFN)NrhFFN)rhFNN)NrhFNN)rhFN)FN)FNNNN)NNNF)NNNNF)NrFNN)NNNN)rTT)NNrN)rrr)r)rFF)rh)NTTN(rhcollections.abcrrenumr functoolsrtypingrtyping_extensionsrrTtorch._prims_commonr rOrr r torch._decompr r r r torch._opsr torch._primsrrrrrrrrrrrrrtorch._prims_common.wrappersrrrr r!rr"r#torch.fx.experimentalr$rtorch.nn.functionalr%r& torch.utilsr'rCr(r)opsr*libraryLibraryrrr{r}r|r5r;rFrSr^rhlinspacelogspacerKrtakerrrrrcummaxcumminrrrrrrr_fft_c2crrr_fft_r2crrandperm generator_outrrrrandintrrlow_outrrandr _fft_c2rr rrrr unsqueeze_r!_sparse_semi_structured_linearrr\r-_sparse_semi_structured_mmr2_sparse_semi_structured_addmmr5_cslt_sparse_mmrD index_reducerK index_reduce_rM index_selectrPsegment_reducer\r unary_outr`rrirrkrlrsrprt _assert_asyncrwmsgrz_printr|_make_dep_tokenrr_functional_sym_constrain_rangerr(_functional_sym_constrain_range_for_sizer_functional_assert_asyncrrrrrrrr _linalg_eighrr_linalg_eigvalslinalg_eigvalsr linalg_eigrrrrrrrr linalg_inv_exrlinalg_ldl_factor_exrerlinalg_ldl_solver linalg_lurlinalg_lu_factor_exrlinalg_lu_solver lu_unpackrr linalg_qrrrr _linalg_svdr!rrr.rAlinalg_solve_triangularrGrJrR _linalg_detrTr\rerureflection_pad1drzreplication_pad1drrreflection_pad1d_backwardrreplication_pad1d_backwardrrreflection_pad2drreplication_pad2dr_weight_norm_interface_backwardrreflection_pad2d_backwardrreplication_pad2d_backwardrrreflection_pad3drreplication_pad3drreflection_pad3d_backwardreplication_pad3d_backwardr_pdist_forwardrXr_pdist_backwardrbaddbmmr bernoullir bernoulli_rrrpoissonr_fused_moving_avg_obs_fq_helperrmmrrfrrrmiopen_batch_normr convolutionr!rz _has_mkldnnr~r"_convolution_pointwiser(_linear_pointwiser+has_mklrr, _mkl_linearr/rr0qconv2d_pointwiseqconv_pointwiserr;binary binary_tensorrDqlinear_pointwiserHrLlinear_dynamic_fp16linear_relu_dynamic_fp16rMrrN max_pool2drXint4mm_packed_weight_cpurare avg_pool2drravg_pool2d_backwardr avg_pool3dravg_pool3d_backwardr_adaptive_avg_pool2dr_adaptive_avg_pool3dr_adaptive_avg_pool2d_backwardr_adaptive_avg_pool3d_backwardrradaptive_max_pool2drrradaptive_max_pool3drrrrepeat_interleaverrmrrrr _unsafe_indexrconvolution_backwardraddbmmr randint_liker _fused_adam_ _fused_adamw_r _fused_adamr_int_mmr_convert_weight_to_int4packr#_convert_weight_to_int4pack_for_cpur_weight_int4pack_mmr!_weight_int4pack_mm_for_cpur#r'r(rH_dyn_quant_pack_4bit_weightrS_dyn_quant_matmul_4bitrW_weight_int8pack_mmrY_cdist_forwardrf_cdist_backwardro_embedding_bagr_embedding_bag_forward_onlyrrnansumrmedian nanmedianr dim_valuesrrr logical_not_rrepeatrzero_rmul_Scalardiv_ logical_and_ logical_or_ logical_xor_radd_sub_rrKsubrrounddecimalsrr __rshift__r __lshift__rzerorr)rfillrrelu_r _add_relurrrelu_with_noiserrrelu_with_noise_functionalrrrelu_with_noise_r index_put_unsafe_index_putr masked_fill_r _masked_scalermasked_scatter_rmasked_scatterrmasked_scatter_backwardr index_put_rrbmmrrrrrwrxrrrrR max_pool2d_with_indices_backwardrmax_pool2d_with_indicesrfractional_max_pool2drmax_pool3d_with_indicesr max_pool3d_with_indices_backwardrr r rgrid_sampler_2d_backwardrrr r"rjonesr.zerosr0select_scatterr3 slice_scatterr5rr8r>gatherrBrOrTrUr[r] scatter_addra scatter_add_rcr\rrrG value_reducerfscatter_rg#_scaled_dot_product_flash_attentionr~r#_scaled_dot_product_cudnn_attentionr0_scaled_dot_product_fused_attention_overrideabler,_scaled_dot_product_flash_attention_backwardr+_scaled_dot_product_flash_attention_for_cpur4_scaled_dot_product_flash_attention_for_cpu_backwardr*_scaled_dot_product_attention_math_for_mpsr'_scaled_dot_product_efficient_attentionr0_scaled_dot_product_efficient_attention_backwardr,_scaled_dot_product_cudnn_attention_backwardr_flash_attention_forwardr_flash_attention_backwardr_efficient_attention_forwardr_efficient_attention_backwardSymIntrr _scaled_mmrr; _scaled_mm_v2r=scatter_reducetwotwo_outr@scatter_reduce_rA multinomialrDrGrOrT_upsample_nearest_exact1dr[_upsample_nearest_exact2dr`"_upsample_nearest_exact2d_backwardrb_upsample_nearest_exact3drrc values_stablerfru_thnn_fused_lstm_cellrzrrrrrargminrrtopkr_segment_reduce_backwardrkthvaluerrrrrr pixel_shufflerr bucketize Tensor_outrhistcr_upsample_bilinear2d_aa_upsample_bicubic2d_aar _upsample_bilinear2d_aa_backwardrrrrDr searchsortedrrembedding_dense_backwardr_embedding_bag_backwardr_embedding_bag_dense_backwardr*_embedding_bag_per_sample_weights_backwardrisinr polygammar_local_scalar_denserrrr r6 _grouped_mmr7_scaled_grouped_mmr8_softmaxr;constant_pad_ndrErH_jagged_to_padded_dense_forwardrKrR_native_multi_head_attentionr`raspecial_airy_aispecial_bessel_y0special_bessel_y1special_modified_bessel_i0special_modified_bessel_i1special_modified_bessel_k0special_modified_bessel_k1!special_scaled_modified_bessel_k0!special_scaled_modified_bessel_k1special_chebyshev_polynomial_tspecial_chebyshev_polynomial_uspecial_chebyshev_polynomial_vspecial_chebyshev_polynomial_w&special_shifted_chebyshev_polynomial_t&special_shifted_chebyshev_polynomial_u&special_shifted_chebyshev_polynomial_v&special_shifted_chebyshev_polynomial_wspecial_hermite_polynomial_hspecial_hermite_polynomial_hespecial_laguerre_polynomial_lspecial_legendre_polynomial_prerfrjrklerp_addcmul_addcdiv_torch._refs.nn.functionaltorch._refs.specialrr0r0r0r4s      4   (   8      6     ;  " '  !  " .      #                      '   "  2  * *  "7  (*& %   :  /  Z &  5 ?'$ ,      "    u  % 2#  Q , H X N     .   *  "  $       n    #  c              -              +  T     \>     6 L + & T  ge(      ,  $  1           ?    +*         $  6,     7          K          /     <5   N             (  7'         "   0    *    "    /  T   (  6   0       E