
    Wj	                   k   U d dl mZ d dlZd dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
Z
d dlZd dlmZ d dlmZmZmZmZ d dlmZmZmZmZmZmZmZ d dlmZ d dlmZ d dlZd dl Z d dl!Z d dl"Z d dl#m$c m%Z& d dl'm(Z( d d	l)m*Z* d d
l+m,Z, d dl-m.Z. d dl/m0Z0 d dl1m2Z2m3Z3m4Z4m5Z5m6Z6m7Z7m8Z8m9Z9m:Z:m;Z;m<Z< d dl=m>Z>m?Z? d dl@mAZAmBZBmCZC d dlDmEZE d dlFmGZGmHZHmIZImJZJmKZK ddlLmMZM ddlNmOZOmPZPmQZQmRZR ddlSmTZTmUZU ddlQmVZVmWZWmXZXmYZYmZZZm[Z[m\Z\m]Z]m^Z^m_Z_m`Z`maZambZbmcZcmdZdmeZe ddl$mfZfmgZgmhZhmiZimjZjmkZkmlZlmmZmmnZnmoZompZpmqZq ddlrmsZsmtZt erddlumvZv  ed          Zw ed          Zx eEddg          Zy ejz        e{          Z|i Z}de~d <   i Zd!e~d"<    eEe j        j                             Ze js        j        Ze js        j        Ze js        j        Z eEe j        j                             Z eEe j        j                 e j        j        g          Z eEe j        j                             Zi Zd#e~d$<   e js        j        Zdd'Zdd+Zdd/Zdd3Zdd8Zdd;Zd d?Z eej        ej        ej        ej        ej        ej        ej        ej        ej        ej        ej        ej        ej        ej        g           e j        e j        e j        e j        e j        e j        e j        e j        e j        e j        e j        e j        e j        d@ZddDZddHZddJZdKdLddQZdR ZddVZddZZddbZddfZd	dhZdKe7j        dKe}fd
djZdk ZddlZ	 	 	 	 	 ddmZddnZdo ZdddqZ ee j        j        dr          ds             Z eej        dr          ddt            ZdKduddvZ eej        j        dr          ddw            ZdKdKdxddyZ eej        dr          dddz            Zdd{e7j        dKdddKdfd|Z end}e7j        d~            eej        d{d          dd            Zd Z eΦ             	 ddZ eej        dKd          d             Z eej        dKd          d             Z eej        ej        ej        ej        ej        g          d             Z eed          r  eej                  eצ            eej        dr          dd            Z eej        dr          dd            Z eej        g          dd            Z eej                  d             Z eej                  d             Z eej                  d             Z eej                  d             Z eej        j                  d             Z eej                  d             Z eej        dr          d             Z eej        dr          d             Z eej        dr          d             Z eej                  d             Z eej        dr           eej        dr           eej        dr          dd                                    Z eej        dr          d             Z eej        dr          dd            Z eej        dr          dd            Z eej        dr          dd            Z eej        dr          dd            ZddZ eej        dr          dd            Zd Z eej        j                  d             Z eej        j                  d             Z eej        dr          dddd            Z eej        j        dr          dd            Z eej        j        dr          dddd            Z eej        j        dr          dd            Z eej        j        dr          dddd            Z eej                  dd            Z eej         dr          ddd            Z  eej        dr          ddd            Z eej        dr          ddd            Z eej        dr          d             Z eej        dr          dd            Z eej        dr          dd            Z eej        dr          dd            Z eej        dr          d             Z eej        dr          d             Z eej	        dr          d             Z	ddÄZ
 eej                  ddń            ZddƄZej        dǄ             Zdd dʄZdd d˄Zdd!d΄Zd"dτZdЄ Z ee js        j        j        dr          dф             Z eej        dr          d҄             Z eej        dr          dӄ             Z eej        j        dr          dԄ             Z eej                  dՄ             Z ej        d          dք             Zdׄ Z eej        j                  Z  eej        j!                  Z" eej#        j                  Z$ eej#        j!                  Z% eej&                    ee js        j'        j(        j                    ee js        j'        j)        j                    eej                  d؄             Z eej#                  dل             Z# eePj*        dr          dڄ             Z+ eePj,        dr          d#dۄ            Z- eePj.        dr          d܄             Z/ eePj0        dr          d݄             Z1 eePj2        dr          d dޜd$d            Z3 eePj&        dr          d dޜd%d            Z4d&dZ5d'dZ6 eej7        j8        dr          dKdKdddd(d            Z7 eej9        j8        e7j:        r          dKdKdd)d            Z9d Z;d Z<d Z=d Z>d Z?d Z@d ZAd  ZB eejC                    eejD                    eejE                    eejF        dK            eejG        j        dK            eejH        e<            eejI        dK            eejJ        dK           e jK        L                                r eejM        dK           e jN        O                                r eejP        dK            eejQ                    eejR                    eejS                    eejT        j                    eejU        j                    eejV                    eejW        jX                    eejY        j                    eejZ        j                    eej[                    eej\        dK            eej]        e;            eej        eA            eej^        e;            eej_        e<            eej`        e;            eeja        e<            eejb        e<            eejc        e<            eejd                    eeje                    eeje                    eejf                    eejg                    eejh                    eeji        e;            eejj                    eejk                    eejl                    eejm                    eejn                    eejo                    eejp                    eejq                    eejr        e<            eejs                    eejt                    eeju        e<            eejv                    eejv        jw                    eejx                    eejy                    eejz                    eej{                    eej|                    eej}                    eej~                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej        j                    eej                    eej                    eej                    eej                    eej                    eej                    eej                    eej        j                    eej        j        dK            eej        e;            ee j        j        j                    ee j        j        j                    ee j        j        j                    eej                    eej                    eej        e<            eej                    eej                    eej                    eej                    eej                    eej        j                    eej        j        eBdK            eej        j        eBdK            eej        j        eBdK            eej        j        dK            eej        j        eBdK            eej        j        eBdK            eej        j        eBdK            eej        j        eBdK            eej        j        eBdK            eej        j        eBdK            eej        j        eBdK            eej        j        eB            eej        j                    eej        j        eB            eej        j        eB            eej        j        eB            eej                    eej        j8        d{            eej        j        e<            eej        dr          dd            Z eej                  ddd            Zd Z eed          r  eej                  e            eej                  d             Z eej        dr          d*d            Z eej        dr          d+d            ZĐd Z ee j        ej        g          ddddKdd            Z ee j                  dd            Z ee j                  d             Z eej                  d             Z eej                  d             Z eej                  	 d,dddd            Z̐d Z͐d Zΐd Z ee j        ej        g          dddddddd            ZАd Zѐd Z  eej                   eѐeЦ                    Z e ed                    Z e ed                     ZՐd Z eej                  dddddd            Z eej                  dddddd             Z eej                  dddddd!            Z eej        j                  d"             Z ee j        ej        g          d#             Z eej        dr          dd$            Z eej        dr          d-d%            Zݐd& Z	 dd'Zߐd( Zdd)Z eej        dr          d*             Z eej        dr          d+             Z eej        dr          dd,            Z eej                  dd-            Zd. Zd/ Z eej        dr          dd0            Z eePj        dr          dd1            Zdd2Z eej        j        dK3          Z eej        j        dK3          Z eej        dr          d4             Z eej        dr          d5             Zed6             Z eej        dr          dd7            Z eej        dr          d.d8            Zdd{d9d/d=Z eej        dr          dd>d0d?            Z eej        dr          d.d@            Z eej        dr          d.dA            Z eej        dr          d.dB            Z eej        dr          d{dCd1dD            Z	 	 d2d3dHZ eej        j                  dd4dJ            Z eej        j                  dd4dK            Z eej        j                  	 dd5dN            Z eej        j                  	 dd5dO            Z eej        j                  	 	 	 d,d6dQ            Z eej        j                  	 	 	 d,d6dR            ZdS Z eej        j                  dT             Zd7dXZ eej         dr          ddY            Z d8d]Zd9d^Zd:d_Z	 d;daZddbdcZdd ZddedfZdg Z eej	        dr          	 ddh            Z	d<doZ
 eej        dr          dp             Zdq Z eej        dr          	 	 	 	 d=dr            Z eej        dr          	 	 	 	 d=ds            Z eej        j        dK3          Z eej        dr          dt             Zd>dvZdw Zdx Zdy Z eej        j        dK3          Z eej                  dz             Z eej        j        dK3          Z eej                  d{             Zd| Z eej                  d}             Z eej                  d~             Zd Z eej        j                  	 d?d            Z eej        dr          	 	 	 	 	 d@d            Z eej        dr          	 	 	 	 	 d@d            Z eej        j        dK3           eej        j        dK3           eej        j        dK3          gZd Z eej        j        dK3          Z eej        dr          	 dd            Z eej         j        dK3          Z! eej         dr          	 dd            Z d Z"dddZ#ddAdZ$d Z% eej&                  dBddd            Z&d Z'd Z(d Z)d Z* eej+        ej+        g          dddKdd            Z, eej-                  dddKdd            Z-d Z.ed             Z/ eej0        j1        dK3          Z2 eej0        j3        dK3          Z4 eej0        j5        dK3          Z6 eej0        d{          d             Z0ddZ7 eej8                  d             Z8 eej9        dr          dd            Z9ed             Z:ed             Z; eej<        d{          dd            Z= eej>        gd{          d             Z>dCdZ? eej<        gd{          d             Z@ eejA        ej<        j8        gd{e7j                  d             Z< eejB        ejB        gd{          d             ZB eejC        ejC        g          dBddd            ZD eejE        j                  ZF eejG        j                  ZH eejI        j                  ZJ eejK        j                  ZL eejM        j                  ZN eejE                  dd            ZE eejG                  dd            ZG eejI                  d             ZI eejK        dr          dd            ZK eejM        dr          dd            ZM eejO                  dBddd            ZO eejP                  dBd            ZQ eejR        dr          dBd            ZS eejT        dr          dBd            ZU  eejV                   e$d                       eejW                   e$d                    ZX  eejY                   e$d                    ZZ  eej[                   e$de j                            Z\  eej]                   e$de j                            Z^ eej_        d{d          Z_ eejv        jw        dK3          Z` eejv        jw        dr          dddKdd            Za eejv        j        dr          dDd            ZvddZbdEdZc ebejd                  Zd eceje                  Ze ebejf                  Zf ebejg                  Zg eejh                  Zh eceji                  Zi ecejj                  Zj eejk                  Zk eejl        d{          Zl eejm        d{          ddd            ZmdFdÄZn eejo        j3        en            ecejp                    ecejq                    eejr                  Zr eejs                  Zs eejt                  Zt eeju        dĐŦ          Zu eejv                  Zv eejw                  Zw eejx                  Zx ebejy                    ebejz                  Zz  eej{        e7j        r          ez            ebej|                    ebej}                    ebej~                    ecej|                    eej        dd{e j        Ʀ          Z eej        dd{e j        Ʀ          Z eej        dd{e j        Ʀ          Z eej        dd{e j        Ʀ          Z eej                  Z eej                  Z  eej                  e             eej                  e            eej                  Z eejr                  Zr ebej                  Z eej                    eej        dǐŦ          Z eej                    eej        e j                     eej                  e            eej        e j                    eej        e j                    eej        e j                    eej        e j                  Z eej        e j                    eej        e j                    ebej                    ebej                    ebej                    ebej                    ebej                    ebej                    ebej                    ebej                    ebej                    ebej                    ebej                    ebej                    ebej                    ebej                    ebej                    ebej                   ddȐlmZmZ dɄ ZeD ][Z eee          D ]\  ZZZ eʐeeeeʦ            eee          D ]\  ZZZ eʐeeeeʦ           \ eej        j        e_d{          Z eej        j3        e_d{          Z eej        j8        e_d{            eej        j        e>          Z eej        j8        e>            eej        j3        e>          Z eej        j        el            eej        j3        el            eej        j        e            eej        j        er            eej        j3        e0            eej        j        e0            eej        j        e0            eej        j        e<          Z eej        j8        e<            eej        j3        e<          Z eej        ej            eej        ed            eej        j        e            eej        j3        e            eej        j        e            eej        j3        e            eej        j        e            eej        j3        e            eej        j        e            eej        j3        e            eej        e            eej        e            eej        e          ZÐd˄ Z eej        j        ej        j        e            eej        j3        ej        j3        e            eej        j        ej        j        e            eej        j3        ej        j3        e            eej        j        ej        j        e            eej        j3        ej        j3        e            eej        j        ej        j        eæ           d̄ Z eej        e_            eej        es            eej        et            eej        eu            eej        ev            eej        ew            eej        ex            eej        e>            eej        j8        e<            eej        j        e=            eej        e            eej        e            eej        e            eej        e            eej        el            eej        eh            eej        ei             eej                  es             eej                  et             eej                  ev             eej                  ew             eej                  ex            eej        ej                    eej        ej                    eej        ej                    eej        ej                    eej        ej                    eej                  dd̈́            Z eej        j                  d΄             Z eej        j                  dτ             Z eej                  dЄ             Z e>j                    D ]&\  ZZ  e e?e                    e           ' ee j                  dф             Z eej                  d҄             Z ee js        j        j                  dӄ             Z ee js        j        j                  dԄ             Z ee js        j        j        j                  dՄ             Z ee js        j        d֦          r, ee js        j        j9        j                  dׄ             Z ee js        j        j                  ddd؄            Zd dِlmZ  ee            ee,          dڄ             Z ee js        j        j        dr          dGd܄            Z ee js        j        j        dr          dd݄            Z  ee js        j        j         dr           ej        ed{ަ                      ee js        j        j        dr          dHd            ZdIdZd dlmZ  eedr          d             Z ee j        j        dr          dddJd            Z ee*dr          dKd            Z	 ee js        j        j
        j                  d             Z
 ee js        j        j        j                  d             Z ee js        j        j        dr          d             ZddlmZmZ  e              e              eePj        dr          d             Zd Z eePj        dr          d             ZddlNmZ  eMe           ddlNmZ  ej                      ej                     ddlNmZ  ej                     ddlNmZ  ej                     ej        dLd            ZdS (M      )annotationsN)defaultdict)Callable
CollectionIterableSequence)AnycastOptionalTYPE_CHECKING	TypeGuardTypeVarUnion)	ParamSpec)patch)counters)associative_scan_op)triton_kernel_wrapper_mutation)FakeScriptObject)get_layout_constraint_tag)canonicalize_dimcanonicalize_dimscheckdtype_to_typeelementwise_dtypesELEMENTWISE_TYPE_PROMOTION_KINDget_computation_dtypeis_boolean_dtypeis_float_dtypeis_integer_dtypeNumber)magic_methodsmethod_to_operator)free_unbacked_symbolshas_free_unbacked_symbolsresolve_unbacked_bindings)
OrderedSet)CeilDivFloorDivIdentityModModularIndexing   )import_submodule   )configinductor_primsirtest_operators)decompositionsget_decompositions)BaseView	DtypeView
ExpandViewIndexingConstantIRNode	is_triton
MutableBoxOnlineSoftmaxReductionops_wrapperPermuteView	Pointwise	ReductionSqueezeView	TensorBoxvalidate_irView)ceildivdecode_device
is_dynamicis_gpuis_pointwise_useis_view,needs_fallback_due_to_atomic_add_limitationspad_listlike#register_op_dtype_propagation_rules#register_op_requires_libdevice_fp64sympy_productuse_scatter_fallback)opsV)ReductionType_T_Pztorchvision::roi_alignzaten::index_add8dict[Union[Callable[..., Any], str], Callable[..., Any]]	loweringsz9dict[torch._ops.OpOverload, Optional[Callable[..., Any]]]_maybe_layout_constraintsz2dict[torch._ops.OpOverload, torch._ops.OpOverload]inplaceable_foreach_opsreturnboolc                     t           j        j        j        D ])} | j        D ]}|j        dk    r|j        t          v s  dS  *dS )Ncall_functionTF)rS   graphcurrent_nodeusersoptargetforeach_ops)nodeusers     ]/home/longshao/multi-rider-rag/.venv/lib/python3.11/site-packages/torch/_inductor/lowering.pycur_node_has_non_foreach_usersrh      s^    $*  J 	 	DG..DK;4N4Nttt 5O	 5    	arg_pairsIterable[Any]4defaultdict[tuple[Any, bool], list[tuple[int, Any]]]c                   t          t                    }d}t          |           D ]\  }}t          |t                    sd}|f}t          |  pt          j        }d }|D ]2}t          |t                    r|j	        
                                } n3|
J d            |r|\  }|||f                             ||f           |S )NFTz.foreach op should have at least one tensor arg)r   list	enumerate
isinstancer   rH   r0   #combo_kernel_foreach_dynamic_shapesrC   data
get_deviceappend)rj   outunpack_argsiargsuse_foreachdevicets           rg   group_foreach_argsr|      s     d

CKY'' 5 54$)) 	K7DD!!OV%O 	  	 	A!Y'' **,, !!#S!!! 	GTV[!"))1d)4444Jri   fnCallable[..., Any]Optional[Callable[..., Any]]c                    t          | t          j        j                  sdS t	          | d          x}rt          |          S | t          v rt          |          S dS )zHGet layout constraints. Returns None if there are no layout constraints.NF)with_default)rp   torch_ops
OpOverloadr   tag_to_layout_constraintrY   )r}   maybe_layout_tags     rg   maybe_layout_constraintsr      sd    b%*/00 t4ReLLLL :'(8999	&&&(,,4ri   tagtorch._C.Tag(Optional[Callable[..., tuple[Any, Any]]]c                $   | t           j        j        j        k    rt          S | t           j        j        j        k    rt          S | t           j        j        j        k    rt          S | t           j        j        j	        k    rd S t          d|            )NzUnknown layout constraint tag: )r   _CTagneeds_exact_stridesconstrain_to_fake_tensorsneeds_contiguous_stridesrequire_contiguous_stridesneeds_fixed_stride_orderconstrain_to_fx_stridesflexible_layoutAssertionError)r   s    rg   r   r      sz     ehl...((
ehl333))
ehl333&&
ehl***t
@3@@
A
AAri   condmsgstrNonec                .    | st          d|           d S )Nzinductor does not support NotImplementedErrorr   r   s     rg   
assert_nyir      s.     F!"Ds"D"DEEEF Fri   Union[Collection[Union[torch._ops.OpOverload, torch._ops.OpOverloadPacket]], torch._ops.OpOverload, torch._ops.OpOverloadPacket]Optional[list[Any]]c                    t           t          t          t          t          f          rd  D             S t           t
          j        j                  rt          	                                nWt           t
          j        j
                  r8t                               fd                                 D                        d S )Nc                ,    g | ]}t          |          S  )add_needs_realized_inputs.0xs     rg   
<listcomp>z-add_needs_realized_inputs.<locals>.<listcomp>   s!    999)!,,999ri   c              3  8   K   | ]}t          |          V  d S N)getattr)r   overloadr}   s     rg   	<genexpr>z,add_needs_realized_inputs.<locals>.<genexpr>   s>       %
 %
&.GB!!%
 %
 %
 %
 %
 %
ri   )rp   rn   settupler'   r   r   r   needs_realized_inputsaddOpOverloadPacketupdate	overloads)r}   s   `rg   r   r      s     "tS%455 :99b9999"ej+,, 
!!"%%%%	B
3	4	4 
$$ %
 %
 %
 %
24,,..%
 %
 %
 	
 	
 	
 4ri   9Union[torch._ops.OpOverloadPacket, torch._ops.OpOverload]
constraintCallable[..., tuple[Any, Any]]c                    t          | t          j        j                  r1|                                 D ]}|t
          t          | |          <   d S |t
          | <   d S r   )rp   r   r   r   r   rY   r   )r}   r   r   s      rg   add_layout_constraintr      sk     "ej122 3 	J 	JH?I%gb(&;&;<<	J 	J )3!"%%%ri   )r   r/   r-                     	   
         dtypeUnion[int, torch.dtype]torch.dtypec                |    t          | t                    s| S | t          v sJ d|  d            t          |          } | S )Nzid z missing from DTYPE_ID_LOOKUP)rp   intDTYPE_ID_LOOKUPr   s    rg   decode_dtyper     sM    eS!! O###%O5%O%O%O###E"ELri   r   r	   ,TypeGuard[Union[TensorBox, sympy.Expr, int]]c                    t          | t                    rBt          |                                           p t	          |                                           S t          | t
          j                  r	| j        du S t          | t                    S NT)	rp   rC   r    	get_dtyper   sympyExpr
is_integerr   r   s    rg   is_integer_typer     sq    !Y "..Q2B1;;==2Q2QQ	Auz	"	" "|t##!S!!!ri   !TypeGuard[Union[TensorBox, bool]]c                    t          | t                    r!t          |                                           S t          | t                    S r   )rp   rC   r   r   r\   r   s    rg   is_boolean_typer   '  s:    !Y #...!T"""ri   F)return_compute_dtyperx   type_promotion_kindr   r   c                R    ddfd|D             }t          |d| i\  }}|r|n|S )Ninpr	   r[   c                    t          | t          t          j        f          r| S t	          |                                           }t          j        dg|z  |                                           S )Nr/   r   )	rp   r!   r   Basiclenget_sizer   zerosr   )r   dims     rg   construct_inputz+get_promoted_dtype.<locals>.construct_input3  sY    cFEK011 	AJcllnn%%C;sSy@@@@ri   c                &    g | ]} |          S r   r   )r   argr   s     rg   r   z&get_promoted_dtype.<locals>.<listcomp>;  s#    111SOOC  111ri   r   )r   r	   r[   r	   )r   )r   r   rx   inpscompute_dtyperesult_dtyper   s         @rg   get_promoted_dtyper   .  sf    
A A A A 2111D111D"4	##6# #M< 1B==lBri   c                P   t          | t          t          f          s| g} nt          |           } t          |           D ]f}t          |t          j        j                  rE|                                D ]0}t          ||          }|t          vr| 	                    |           1g| S r   )
rp   rn   r   r   r   r   r   r   rX   rt   )aten_fnr}   r   other_fns       rg   get_overloadsr   B  s    ge}--  )w--7mm - -b%*566 	-LLNN - -"2x009,,NN8,,,Nri   rb   >Union[Any, torch._ops.OpOverloadPacket, torch._ops.OpOverload]	namespacec                    t          | t          j        j                  r	|| j        v S t          | t          j        j                  r||                                 v S dS NF)rp   r   r   r   _qualified_op_namer   name)rb   r   s     rg   in_namespacer   R  sX     "ej122 &B111	B
-	.	. &BGGII%%5ri   rC   rz   torch.devicec           	        t          | j        t          j                  r!t	          |                                           r| S d |                                 D             }|                                 }||j        dk    rx||k    rrt          |          dk    st          |          dk    rL|d         dk    r@t          t          j
        t          j                            | |d                              S | S )zB
    Copy cpu scalar if doesn't not match with given `device`
    c                V    g | ]&}t           j        j                            |          'S r   rS   r_   sizevarssize_hint_or_throwr   ss     rg   r   z)maybe_copy_cpu_scalar.<locals>.<listcomp>d  s+    IIIqAG//22IIIri   Ncpur   r/   F)rp   rr   r2   ReinterpretViewr%   r   rs   typer   rC   
StorageBox
DeviceCopycreate)r   rz   size
cur_devices       rg   maybe_copy_cpu_scalarr  \  s     afb011 5N	

6 6  IIAJJLLIIIDJOu$$&  YY!^^D		Q47a<<r}';';Az5'Q'QRRSSSHri   	list[Any]kwargsdict[str, Any]	broadcast)Optional[ELEMENTWISE_TYPE_PROMOTION_KIND]convert_input_to_bool tuple[list[Any], dict[str, Any]]c           	     b    d t                     D             }d                                 D             }|s|s fS |s|r|rt          j        nId  D             }|                    d                                 D                        t          |d|i|r |d                  n|d                                                  |D ]}t           |                    |<   |D ]}	t          |	                   |	<   dfd
fd D              fd                                D             |rht          t          t          j         fd|D             fd|D                                  }
t          |
d                                                   }t          ||
dt          |                             D ]
\  }}| |<   t          ||
t          |          d                   D ]
\  }	}||	<   t!          t                               D ]@}t#           |         t$          j                  rt)          j         |         |           |<   AD ]@}	t#          |	         t$          j                  rt)          j        |	         |          |	<   A fS )zB
    Transforms arguments for broadcasting and type promotion
    c                B    g | ]\  }}t          |t                    |S r   rp   rC   r   rw   r   s      rg   r   z"transform_args.<locals>.<listcomp>{  s+    NNN$!QZ95M5MNANNNri   c                B    g | ]\  }}t          |t                    |S r   r  )r   kvs      rg   r   z"transform_args.<locals>.<listcomp>|  s+    OOODAqjI6N6NOaOOOri   c                t    g | ]5}t          |t          t          j        f          st	          |d           3|6S r   )rp   r!   r   r   hasattrr   as     rg   r   z"transform_args.<locals>.<listcomp>  sT       a&%+!677 <C1g;N;N  ri   c              3  :   K   | ]}t          |d           |V  dS )r   N)r  r  s     rg   r   z!transform_args.<locals>.<genexpr>  s1      !T!T7@S@S!T!!T!T!T!T!T!Tri   r   r   r   r	   r[   c                    t          | t                    rt          |           S t          | t          j                  rt          j        | j                  S | S )Nvaluer   rz   )rp   rC   to_dtyper2   Constantr  )r   rz   r   s    rg   promoteztransform_args.<locals>.promote  sW    #y)) U+++C-- {%OOOO
ri   c                &    g | ]} |          S r   r   )r   r  r  s     rg   r   z"transform_args.<locals>.<listcomp>  s!    )))q

)))ri   c                .    i | ]\  }}| |          S r   r   )r   r  r  r  s      rg   
<dictcomp>z"transform_args.<locals>.<dictcomp>  s'    ;;;DAq!WWQZZ;;;ri   c              3  (   K   | ]}|         V  d S r   r   r   rw   rx   s     rg   r   z!transform_args.<locals>.<genexpr>  s'      33T!W333333ri   c              3  (   K   | ]}|         V  d S r   r   )r   r  r  s     rg   r   z!transform_args.<locals>.<genexpr>  s'      771VAY777777ri   N)r   r	   r[   r	   )ro   itemsr   r\   extendvaluesr   rs   r  broadcast_tensorsrn   	itertoolschainr   zipr   rangerp   r2   r  r8   r  )rx   r  r  r   r
  args_indiceskwargs_indicespromoting_argsrw   r  broadcastedr  r   rz   r   r  s   ``           @@@rg   transform_argsr0  p  sM    ON)D//NNNLOOFLLNNOOON  V| %<3 %<  	JEE   N !!!T!TV]]__!T!T!TTTT&$7 E &2PDa!!vnQ>O7P
*,, 	  	= 	=A+DGV<<DGG 	A 	AA-fQi@@F1II	 	 	 	 	 	 	 *)))D)));;;;FLLNN;;; ?'3333l3337777777  
 KN++--..k2EC4E4E2E&FGG 	 	DAqDGGC4E4E4G4G(HII 	 	DAqF1IIs4yy!! 	; 	;A$q'2;// ;$+DGT::Q 	? 	?A&)R[11 ?&-fQi>>q	<ri   r   torch._ops.OpOverload	decomp_fnc                    t          j                  dfd            }t          |           }t                              |           t
                              t                              ||                     |S )a  
    Add a foreach lowering to lowerings dict.

    Arguments:
        aten_fn: torch.ops.aten.* fn we are lowering
        decomp_fn: alternate implementation on our IR
        broadcast: True to apply broadcasting to tensor inputs
        type_promotion_kind: kind of type promotion applied to tensor inputs, `None` means no type promotion
        convert_input_to_bool: some logical ops require inputs are converted to bool
    rx   r	   r  r[   c                 6     | i |}t          |           |S r   )rD   )rx   r  ru   r2  s      rg   wrappedz+_register_foreach_lowering.<locals>.wrapped  s*    i(((C
ri   )rx   r	   r  r	   r[   r	   )	functoolswrapsr   rd   r   rX   dictfromkeys)r   r2  r5  aten_fnss    `  rg   _register_foreach_loweringr;    s     _Y      
 W%%Hx   T]]8W55666Nri   lowering_dictc                     t          j                   fd            }t                      |                    t                               |                     |S )a  
    Add a lowering to lowerings dict

    Arguments:
        aten_fn: torch.ops.aten.* fn we are lowering
        decomp_fn: alternate implementation on our IR
        broadcast: True to apply broadcasting to tensor inputs
        type_promotion_kind: kind of type promotion applied to tensor inputs, `None` means no type promotion
        convert_input_to_bool: some logical ops require inputs are converted to bool
    c                    t          |           } t          |          }d}t          |           dk    r9t          | d         t           t          f          rd}t          | d                   } t          d D                       s#t          d |D                       r
J d            t          | |          \  } }|r| g}  | i |}t          |           |S )NFr/   r   Tc              3  H   K   | ]}|t           v pt          |d           V  dS )_c10d_functionalN)	fallbacksr   )r   r}   s     rg   r   z6_register_lowering.<locals>.wrapped.<locals>.<genexpr>  sG       
 
JLR9_DR1C D D
 
 
 
 
 
ri   c              3  "   K   | ]
}|d k    V  dS )ru   Nr   r   s     rg   r   z6_register_lowering.<locals>.wrapped.<locals>.<genexpr>  s&      66!1:666666ri   zout= ops aren't yet supported)	rn   r8  r   rp   r   allanyr0  rD   )	rx   r  unpackedru   r   r  r
  r2  r   s	       rg   r5  z#_register_lowering.<locals>.wrapped  s   t**!%ft99>>ja4-@@>HQ==D 
 
PW
 
 
 
 
 	X 66v66666WW8WWW6%&)%8:O
 
f  	6Di(((C
ri   )r6  r7  r   r   r8  r9  )r   r2  r  r   r
  r<  r5  s   `````  rg   _register_loweringrF    s    & _Y         6 G$$Gw88999Nri   .Callable[[Callable[_P, _T]], Callable[_P, _T]]c                @    t          j        t          | ||||          S )z+
    Shim to support decorator syntax.
    )r  r   r
  r<  )r6  partialrF  )r   r  r   r
  r<  s        rg   register_loweringrJ    s0     /3#   ri   c                &   t          |          }| r| |k    r|S g }t          j        t          |           t          |          t          j        j                  D ]\  }}t          j        j	        
                    |          r|                    |           @t          j        j	        
                    |          r|                    |           zt          j        j	                            ||           t          t	          j        |          j                  t          t	          j        |          j                  k     r|                    |           |                    |           t          t          |                    S )z
    Broadcasting logic based on symbolic shapes.

    We give the shapes 0 and 1 concrete values, while all other shapes
    are symbolic sympy formulas.
    )	fillvalue)r   r(  zip_longestreversedr   SOnerS   r_   r   is_size_one_or_falsert   check_equalsr   expandfree_symbols)r  boutputr   ys        rg   broadcast_symbolic_shapesrX  '  sJ    	aA QF%hqkk8A;;%'+VVV 
! 
!170033 		!MM!W22155 	!MM!G))!Q///5<??/003u|A7S3T3TTTa    a    &!!"""ri   c                   ||
J d            ||t           j        }t          d | D                       s| S t          d | D                       r!|pt	          | d|ifdfd| D             S t          d | D                       }g }| D ]L}t          |t          t          f          r|	                    t          j        t          j        ||                                |                                          t!          |                                                               t          |t$          j                  r}|	                    t          j        t)          ||                                |                                	          t!          |                                                               7|	                    |           N|S )
NzEonly one of override_return_dtype or type_promotion_kind may be givenc              3  d   K   | ]+}t          |t          j        t          t          f          V  ,d S r   )rp   r   r   r   floatr   s     rg   r   z$promote_constants.<locals>.<genexpr>I  s5      HHAz!ek3677HHHHHHri   c              3  d   K   | ]+}t          |t          t          t          j        f          V  ,d S r   )rp   r   r[  r   r   r   s     rg   r   z$promote_constants.<locals>.<genexpr>K  s5      
D
D:a#uek233
D
D
D
D
D
Dri   r   c                    t          | t          j                  r$t          j        | t          d                     S t          j        | t          d                     S )Nindexr   rz   r  )rp   r   r   r2   r9   rG   r  r   r   s    rg   
const_funcz%promote_constants.<locals>.const_funcR  sb    !U[)) U*5t1D1D    {%d@S@STTTTri   c                &    g | ]} |          S r   r   )r   r   ra  s     rg   r   z%promote_constants.<locals>.<listcomp>Z  s!    ...!

1...ri   c              3  h   K   | ]-}t          |t          t          t          j        f          )|V  .d S r   )rp   rC   r8   r2   r  r   s     rg   r   z$promote_constants.<locals>.<genexpr>[  s:      WWA:a)Z1U#V#VWaWWWWWWri   r  r^  )r   DEFAULTrD  rC  r   nextrp   r   r[  rt   r8   r  r2   r  r   get_device_or_errorrn   r   r   r   r9   )inputsoverride_return_dtyper   exru   r   ra  r   s         @@rg   promote_constantsrj  A  s*    (,?,G,GO -H,GG $)<)D=EHHHHHHH 

D
DV
D
D
DDD /% 
);*
 !4*
 *
	U 	U 	U 	U 	U /...v....	WWWWW	W	WB
C  a#u&& 	JJ!Kr||~~b>T>T>V>V   ''	     5;'' 
	JJ!$r||~~b>T>T>V>V   ''	     JJqMMMMJri   c                .     d dd fd}|S )Nalpharg  rC   c           	       	
 "t          d D                       r	rJ   S t                    r1| .| dk    r(t                    t          d         |           d<   n| J d D             	d                                         pd                                         dd          D ]s}t          |t          j                  sWt                    t          |                                          k    s%J d d d|                                             tt          j        t          j        f
t          j        d uo^t          t          j        dd           d uoAt          j        j        j        d uo)t          j        j        j                            d	d
          o
v 	
f	d}s_d }D ]>}t'          |                                j                  r|                                } n?|sd                                         }p|}t-          j        ||          S )Nc              3  ^   K   | ](}t          |t                    ot          |          V  )d S r   rp   r:   r;   r   r   s     rg   r   z0make_pointwise.<locals>.inner.<locals>.<genexpr>  sH       /
 /
;>JsF##6	#/
 /
 /
 /
 /
 /
ri   r/   c                6    g | ]}|                                 S r   make_loaderr   s     rg   r   z1make_pointwise.<locals>.inner.<locals>.<listcomp>  s     333q1==??333ri   r   zndim mismatch  r`   low_precision_pointwise_barrierFc                  	  t                     t                    k    sJ d  d             t          j        k    r  fdD              S g }t                    D ]q\  }} |           }|                                         }	r0|v r,t          j        ||d          }t          j        ||          }|                    |           r 
| }r,t          j        |d          }t          j        |          S |S )Nzwrong ndim rv  c                &    g | ]} |          S r   r   )r   loadr_  s     rg   r   zCmake_pointwise.<locals>.inner.<locals>.inner_fn.<locals>.<listcomp>  s!    4U4U4UTTT%[[4U4U4Uri   F)use_compute_types)r   r   r\   ro   r   rR   r  rt   )r_  inputs_loaded	inp_indexrz  ru   	inp_dtypedowncastr   emulate_output_castemulate_precision_castsr}   rg  loaders	low_pr_fpoverride_fn_when_input_boolrangess   `      rg   inner_fnz/make_pointwise.<locals>.inner.<locals>.inner_fn  sH   u::V,,,.LE.L.LF.L.L,,,
""'B'N224U4U4U4UW4U4U4UVV "'0'9'9 . .OIt$u++C &y 1 ; ; = =I. @9	3I3I#&<YRW#X#X#X!l8Y??!((----b-(& 9  #|C%PPPH<%888
ri   rz   r   r  r  )rD  rj  rn   mulr   r   rp   r2   BaseConstantr   r   bfloat16float16rS   r_   r   r`   metagetrI   rs   r   r@   r  )rm  rg  otherr  rz   rw   r   r  r  r  r  r  allow_alphar}   override_devicer  rh  triton_fallbacks    `    @@@@@@rg   innerzmake_pointwise.<locals>.inner~  s   &3 /
 /
BH/
 /
 /
 ,
 ,
& #""?"?F++"6+@AA 	! UaZZf U33r
===33F333##%%%>)<)<)>)>ABBZ 	A 	AEeR_55 AV  I I : : :@@@V@@enn.>.>@@: :  ^U]3	G4 X66dBX$)5X $)--.OQVWW	 	  6L%9:L	 	 	 	 	 	 	 	 	 	 	 	 	,  	0F  !,,..-.. \\^^FE  0--// !*F	
 
 
 	
ri   )rg  rC   r   )r}   rh  r  r  r  r  r  s   `````` rg   make_pointwiser  v  sc     )- L
 L
 L
 L
 L
 L
 L
 L
 L
 L
 L
 L
 L
\ Lri   c                     ddd fd}|S )Nr/   rl  rg  list[list[TensorBox]]c                F    t          t          j        j        j                  dk    p*t          j        j        j        t          v pt                      }d }|D ]"}t          |t          t          f          r|} n#|
J d            g }|D ]Z}t          |t          t          f          s'|                    |gt          |          z             E|                    |           [t          t          |           } 	fd}t          |t          |          ||          S )Nr   z1at least one input must be a list to a foreach opc                "    r | diS  |  S )Nrm  r   )rx   r  rm  pw_fns    rg   apply_fnz7make_foreach_pointwise.<locals>.inner.<locals>.apply_fn  s,     $ud0%000ud|#ri   )r   rS   r_   r`   ra   rc   inplace_foreach_opsrh   rp   rn   r   rt   r|   r*  foreach_group_loop)
rm  rg  realize_outputsa_list_inputinputbroadcast_inputsgroupsr  r  r  s
   `       rg   r  z%make_foreach_pointwise.<locals>.inner  sY   $*++q0 0w#*.AA0-// 	  	 	E%$// $ ''? (''
  	/ 	/EedE]33 / ''#l2C2C(CDDDD !''....#C)9$:;;	$ 	$ 	$ 	$ 	$ 	$ 	$ "&#l*;*;XWWWri   )rg  r  r   )r  r  r  s   `` rg   make_foreach_pointwiser    sE    45 !X !X !X !X !X !X !X !X !XF Lri   c                   dg|z  }|                                  D ]\  \  }}}g }|D ]~\  }	}
 ||
          }|||	<   t          j                            |t          j                  r?|r=|r;|                                 |                    |                                           |rt          j        	                    |           t          d |D                       sJ |S )aa  
    Common loop over grouped foreach arguments.

    Args:
        groups: Result of group_foreach_args - dict mapping (device, use_foreach) to groups
        num_outputs: Number of outputs to produce
        apply_fn: Function to apply to each set of args, returns the output
        realize_outputs: Whether to realize outputs for foreach fusion
    Nc              3     K   | ]}|d uV  	d S r   r   r   s     rg   r   z%foreach_group_loop.<locals>.<genexpr>  &      ..q}......ri   )r$  rS   r_   has_featureBackendFeatureFOREACHrealizert   get_operation_nameregister_operation_listrC  )r  num_outputsr  r  outputsrz   ry   groupoperation_list
output_indrx   rV  s               rg   r  r    s    f{"G(. < <$u$& % 
	C 
	CJXd^^F"(GJ ##FN,BCCCC $C
    %%f&?&?&A&ABBB 	<G++N;;;..g........Nri   copyc                    |                                  k    r|rt          |           n| S fd} t          |          |           S )Nc                2    t          j        |           S )N)	src_dtype)rR   r  )r   r   r  s    rg   	_to_dtypezto_dtype.<locals>._to_dtype  s    |Au	::::ri   rh  )r   cloner  )r   r   r  r  r  s    `  @rg   r  r    so    IE&uQxxxQ&; ; ; ; ; ; B>)5AAA!DDDri   r   c                   ddl m} |}| j        } ||t          j                  }t          j        |          5   |j        |  ddd           n# 1 swxY w Y   |j        }|sJ t          |          }dgt          |          z  }	|
                                D ]\  \  }
}}g }|D ]q\  }}||	|<   t          j                            |
t          j                  r=|r;|                                 |                    |                                           r|rt          j                            |           t%          d |	D                       sJ |	S )aI  
    This lowers an invocation of foreach_map
    The way this works is that an arbitrary N-arg func is provided by the user, looped over by the
    polyfill with the same semantics as a foreach op (a loop applying an n-ary function to n args)
    and then traced into a subgraph by dynamo.
    This code allows us to inline the subgraph into the main graph lowering using the PontwiseSubgraphLowering.
    The graph outputs represent the vertically fused sequence of ops, and then register_operation_list
    below registers the buffers as horizontally fuseable in the scheduler.
    r/   )PointwiseSubgraphLowering)root_graph_loweringNc              3     K   | ]}|d uV  	d S r   r   r   s     rg   r   z_foreach_map.<locals>.<genexpr>J  r  ri   )subgraph_loweringr  graph_modulerS   r_   set_graph_handlerrungraph_outputsr|   r   r$  r  r  r  r  rt   r  r  rC  )subgraphrx   r  r  rg  gmpw_subgraphsub_outputsr  r  rz   ry   r  r  r  rV  s                   rg   _foreach_mapr  !  s    =<<<<<F		B++BAGLLLK	
	[	)	) ! !  ! ! ! ! ! ! ! ! ! ! ! ! ! ! ! +K;,,Ffs;'''G(. < <$u$& 	C 	C 
"(GJw""6>+ABB C{ C   %%f&?&?&A&ABBB 	<G++N;;;..g........Ns   AAAc                H   |j         s|                                 j         rq|                                 r3t          | |          }t          j                            ||            |S  t          t          j	        j
        d          | |          S t          | |d          S )Nr   Fadd_to_fallback_setTr  )
is_complexr   r   
empty_liker2   InplaceCopyFallbackr  fallback_handlerprimsconvert_element_typedefaultr  )r   r   dsts      rg   _convert_element_typer  N  s     
1;;==3 
::<< 		 Qe,,,C"))#q111J#*2     Au4((((ri   r  c               .   |                                  }||k    r|rt          |           n| S d } ||          } ||          }||k    r( t          t          j        j                  | |          S t          t          j        | |                    S )Nc                t    | j         rt          j        |           j        S t          j        |           j        S r   )is_floating_pointr   finfobitsiinfor   s    rg   _get_primitive_bitwidthz1to_dtype_bitcast.<locals>._get_primitive_bitwidthc  s3    " 	+;u%%**;u%%**ri   )	r   r  r  atenviewr   rC   r7   r  )r   r   r  x_dtyper  src_bitsdst_bitss          rg   to_dtype_bitcastr  ^  s    kkmmG%&uQxxxQ&+ + + '&w//H&&u--H80	00E:::)!U33444ri   c                    |j         s|                                 j         rLt          j        t          j                            t          j        j        j	        j
        | |                    S t          | |          S r   )r  r   rC   r  r2   ComplexViewr   rR   r  r  r   r  r`  s     rg   _view_dtyper  r  sg     
1;;==3 
N!!%)."5";QFF
 
 	
 Au%%%ri   r  non_blockingc                   t          |          }|                                 |k    r|rt          |           n| S t          j        t
          j                            | ||                    S r   )rG   rs   r  rC   r  r2   r   )r   rz   r  r  s       rg   	to_devicer  {  s\    6""F||~~&uQxxxQ&BM00FLIIJJJri   c                (    t          | |d|          S )NTr  )r  )r   rz   r  s      rg   _device_putr    s    QTEEEEri   Tc	                `   |p| j         }t          |          }	t          |||           |t          |          }t          |	||||          }	 t	          | |||          |	          }	t          t          |          r. t	          t          t          |          d|          |	           |	S )z3A pointwise function that maps ops.{name} to inputsN)rh  r  r  r  )r  r   r
  )r   r
  )__name__r>   rN   r  rJ  r  r  r   )
r   r   r  r   r
  rh  r  r  r  r}   s
             rg   register_pointwiser    s     #7#D	T		B'!#8   #.&12M&N&N#	
3$?'
 
 
B
	/3	
 
 

 	

 

B ud 	
E4   $"7	
 	
 	
 		 	 	
 Iri   ldexp)r   rh  )r  r   nc                p  	 t          d          |                                 }|                                }|j        }|j         o|t          j        k    }|r |rfd} t          |          | |          S t          |           rt          j        n|		fd} t          |	          | |          S )Nr  c                     | |          S r   r   )r   r  ldexp_fns     rg   compute_ldexpz%ldexp_lowering.<locals>.compute_ldexp  s    8Aq>>!ri   c                    t          j        |          }t          j        d          }t          j        ||          }t          j        | |          S )Ng       @)rR   r  constantpowr  )r   r  
n_out_typetwo
pow_result	out_dtypes        rg   compute_fallbackz(ldexp_lowering.<locals>.compute_fallback  sH    a33J,sI..Cj11J71j)))ri   r  )r>   r   r  r   r\   r  r   float32)
r   r  r  n_dtype
x_is_floatn_is_intr  r  r  r  s
           @@rg   ldexp_loweringr    s    7##HkkmmGkkmmG*J,,FEJ1FH h 	" 	" 	" 	" 	" -~m,,Q222 &5Q%7%7DEMMW		* 	* 	* 	* 	*
~"+
 
 
 Q  	ri   c                 f   d} t          d          fd}fd}t          |          t          |t          j                  gfd} t	          t
          j                  |          }t          t          |           r- t	          t          t          |           d          |           |S )z2A pointwise function that maps ops.frexp to inputsfrexpc                       | i |d         S Nr   r   rx   r  r   s     rg   frexp0zregister_frexp.<locals>.frexp0      ud%f%%a((ri   c                       | i |d         S Nr/   r   r  s     rg   frexp1zregister_frexp.<locals>.frexp1  r  ri   r  c                 <     d         | i | d         | i |fS Nr   r/   r   )rx   r  pw_fnss     rg   r}   zregister_frexp.<locals>.fn  s4    vay$)&))96!9d+Ef+E+EEEri   Nr  )
r>   r  r   int32rJ  r  r   r  r  r   )r   r  r  r}   r   r  s       @@rg   register_frexpr    s   D  E) ) ) ) )) ) ) ) ) 	vvU[AAAF
F F F F F
	

 


 

B ud 	
E4   $	
 	
 	
 	 	 	 Iri   c                H    t          ||          }t          | |          }|S )Nr  )r  r;  )r   pointwise_lowering_fnr  r}   s       rg   register_foreach_pointwiser    s+    
 
  5;	O	O	OB	#GR	0	0BIri   c           	     f   d }t          |t          t          f          r t          |          |          }t          |t          t          f          r t          |          |          }| ||gt	          d         d         t
          j                  }d t                    D             }t          |t          fd|D                        D ]
\  }}||<   t          t                              D ]k}t          |         t          j                  rIt          j        |         t!          |d                                                                      |<   l t%          ||          d         t'          d         |          t'          d         |                    S )	Nc                     t          j        |  S r   )rR   whererx   s    rg   r}   zwhere.<locals>.fn	  s    y$ri   r/   r-   r  c                B    g | ]\  }}t          |t                    |S r   r  r  s      rg   r   zwhere.<locals>.<listcomp>  s+    IIITQ
1i0H0HIqIIIri   c                     g | ]
}|         S r   r   r"  s     rg   r   zwhere.<locals>.<listcomp>  s    0J0J0JQa0J0J0Jri   r   r  )rp   r[  r   constant_liker   r   rd  ro   r*  r'  r+  r   r2   r  r8   r  rn   r   r  r  )	r   r  rU  r}   r   indicesrw   r   rx   s	           @rg   r  r    s         !eS\""  M!Q!eS\""  M!Q!Q<DQa.M.U  E JIYt__IIIGG.0J0J0J0J'0J0J0JKLL  1Q3t99 T Td1gr{++ 	T 'Qd71:6F6O6O6Q6Q1R1RSSDG:>"E:::Q$q'5))8DGU+C+C  ri   c                    t          |           dk    r3t          | d         t          t          f          rt	          | d          S | S t          j        t          d | D             d          }g }| D ]}t          |                                          x}|k    rnWt          |          t          |          k    s't          d t          ||          D                       rt          ||          }|                    |           |S )Nr/   r   c              3  >   K   | ]}|                                 V  d S r   )r   r   s     rg   r   z$broadcast_tensors.<locals>.<genexpr>'  s*      #A#AQAJJLL#A#A#A#A#A#Ari   r   c              3     K   | ]Q\  }}t           j        j                            |          t           j        j                            |          k    V  Rd S r   )rS   r_   r   rQ  r   r  rU  s      rg   r   z$broadcast_tensors.<locals>.<genexpr>.  sj       .
 .
 1 G11!44w44Q778.
 .
 .
 .
 .
 .
ri   )r   rp   rn   r   r'  r6  reducerX  r   rD  r*  rS  rt   )rg  rc   r  r   sizess        rg   r'  r'     s   
6{{afQi$// 	1$fQi00(/!#A#A&#A#A#A2   F G 
 
1::<<(((EV33ZZ3v;;&&# .
 .
 E6**.
 .
 .
 +
 +
&
 q&!!AqNri   c                    | S r   r   r   s    rg   nopr!  8  s    Hri   
lift_freshc                .   t          | t                    sJ |&t          t          j        | j                            S t          |t
          t          j        f          r$t          j	        j
                            |          nt          d |D                       }t          t          |                                           |          }t!          t          |t                    s|fn|          }g }t#          |                                           D ]U\  }}||v r7t          j	        j
                            t          j        |d                    s|                    |           V||                                 k    rt+          | |          n| S )Nc              3  ^   K   | ](}t           j        j                            |          V  )d S r   rS   r_   r   	guard_intr   ds     rg   r   zsqueeze.<locals>.<genexpr>J  s5      >>Q17#--a00>>>>>>ri   r/   )rp   rC   rB   r  rr   r   r   r   rS   r_   r   r&  r   r   r   r   r'   ro   guard_or_falseEqrt   r  )r   r   dims	new_shaper(  r   s         rg   squeezer-  A  s]   a#####
{+AF33444 cC,--	?""3'''>>#>>>>> 
 C

--s
3
3CJsE$:$:CsffDDDI!**,,''    1T		ag.==ehq!nnMM	Q "+ajjll!:!:49Ari   c                <    t          t          | |                    S r   )r  r-  )r   r   s     rg   squeeze_copyr/  X  s    C!!!ri   c                    t          | |          }t          | t                    sJ t          |t                    sJ |j        | _        | S r   )r-  rp   rC   rr   r   r   vals      rg   squeeze_r3  ]  sI    
!S//Ca#####c9%%%%%XAFHri   c                    t          |           rt          | dt          j                  S t	          d          } t          |t          j                  |           S )NFr   isinfr  r   	full_liker   r\   r>   r  r   r}   s     rg   r5  r5  f  U    q 5E4444	W		B?>"EJ???BBBri   c                    t          |           rt          | dt          j                  S t	          d          } t          |t          j                  |           S )NFr   isnanr  r6  r8  s     rg   r;  r;  n  r9  ri   c                    t          |           rt          |           S t          d          } t          |          |           S )Nceilr   r  r>   r  r8  s     rg   r=  r=  v  sB    q Qxx	V		B>"a   ri   c                    t          |           rt          |           S t          d          } t          |          |           S )Nfloorr>  r8  s     rg   r@  r@  ~  B    q Qxx	W		B>"a   ri   c                    t          |           rt          |           S t          d          } t          |          |           S )Nroundr>  r8  s     rg   rC  rC    sB    q %Qxx!!!~b!!!$$$ri   c                    t          |           rt          |           S t          d          } t          |          |           S )Ntruncr>  r8  s     rg   rE  rE    rA  ri   c                X   t          | g          \  } t          | t          j                  r"t	          j        | t          |                    S t          | t                    sJ t          |t          t          f          sJ t          | 	                                          t          |          k    r| S t          | 	                                          st          j        j                            t          | 	                                                    }|dk    rVt          |          sG|                     t          j        j                            t          |                    |z             t          t	          j        | j        t          |                              S r  )rj  rp   r2   r  r8   r  r   rC   rn   r   r$   rS   r_   r   r   rP   
mark_reuserr   )r   r  x_size_products      rg   rS  rS    sW   aS!!DQ!R_%% 2 E%LL111a#####edE]+++++QZZ\\eEll** .. )<<!**,,''
 
 A&;E&B&BLL 33M%4H4HII!"   Z&qvuU||<<===ri   c                    t          |          }|D ]}d||<   | }t          |          D ]\  }}|dk    rt          ||          }t          ||          S Nrr  )rn   ro   	unsqueezerS  )r  shapebroadcast_dimensionsr   broadcast_dimensionr  idxr   s           rg   broadcast_in_dimrP    sr    UA3 $ $!#
	AA,, " "Q77!S!!A!Uri   c                F    t          | |                                          S r   )rS  r   )r   rW  s     rg   	expand_asrR    s    !QZZ\\"""ri   c                V   t          |                                           t                    t                    k    rRt          j        j        gt                    t                    z
  z  z   t          | t                              } t                    t          |                                           k    sJ t          |                                           }d}t          t                              D ]$}|         dk    rd}||         |         z  ||<   %|r6t          || 	                                | 
                                          S t          d t                    D                       rt          t          | |                    S fd}t                    st          |          s~t           j        j                            t)                              }|dk    rG|                     t           j        j                            t)          |                    |z             |                                 t/          j        | 
                                | 	                                |t          |                    S )NFr   Tr   rz   c              3  4   K   | ]\  }}|d k    p|d k    V  dS r/   Nr   r  s      rg   r   zrepeat.<locals>.<genexpr>  s3      
A
A$!QAFa1f
A
A
A
A
A
Ari   c                T   t          |           t                    k    sJ t          |           } t          t                              D ]O}|         dk    rA|         dk    rt          j        j        | |<   /t          | |         d|                   | |<   P |           S r  )r   rn   r+  r   rO  Zeror,   )r_  rw   old_sizerepeatsx_loaders     rg   r  zrepeat.<locals>.inner_fn  s    5zzS\\))))Us7||$$ 	I 	IAqzQA;!##$w|E!HH.uQxHQKHHE!Hxri   r  )rn   r   r   r   rO  rP  r  r+  emptyr   rs   rC  r*  r  rS  r$   rS   r_   r   r   rP   rG  ru  r@   r  )	r   rZ  new_sizezero_tensorrw   r  old_size_productrY  r[  s	    `     @@rg   repeatr`    s]   AJJLL!!H
7||c(mm##GK=CLL3x==$@AHLDNN##w<<3qzz||,,,,,,AJJLL!!HK3w<<   / /1:??KqkGAJ. KXQ[[]]1<<>>JJJJ

A
A#gx*@*@
A
A
AAA *VAx(()))	 	 	 	 	 	 	 !** 3H3R3R 7+>>}X?V?VWWa LL 33M(4K4KLL#$  
 }}H||~~kkmmH~~	   ri   r  Sequence[sympy.Expr]c                P    t          t          j        | j        |                    S r   )rC   rE   r  rr   )r   r  s     rg   r  r    s      T[//000ri   c                    t          | t                    sJ t          |t          t          f          sJ t          t	          j        | j        t          |                              S r   )rp   rC   rn   r   r?   r  rr   )r   r+  s     rg   permuterd    sU    a#####dT5M*****['d<<===ri               c           
        ddl m}m} t          | t                    sJ t          | |d          }|                                 |         }t          j        |          }t          |t          j	                  s|dk    s
J |            	 |dk    r-t          j        j                            ||          r|dk    r| S n# t          $ r Y nw xY wdd}	d\  }
}|}|r$ |	||d          }
 |	|||          }|
||
|}}d}|s7t	          t          j                            | j        |||||                    S |sJ  |t          j        j        j        t          j        j        j        d	                   }|J t-          |          d
k    s
J |            d\  }}|                                D ]=\  }}| |d          t1          j        |          fk    r|}+| |d          fk    r|}>|
|J t          j        |||||                                 |                   }t          j                            |          |_        t          j                            |           |}|                                 |                                  |
<|J |                                  j!        |
| "                                |         z  z   }nt          j#        |||                                  j!        | "                                |         |                                 |         d          }t          j                            |          |_        t          j                            |           |}tI          |                                           }tI          | "                                          }|||<   ||xx         |z  cc<   tK          | |||          S )a	  
    Lowers a slice call, creating ExternKernels for the output size & storage offset symbols,
    if the indices are unbacked and appropriate semantics aren't known.
    If they are known (indices are static/backed/unbacked with info), a SliceView is created.
    r   )CallMethodKeyr&   r/   Nc                   | |S d }t          j        |           } t          j        |          } |t          j        | d                    r  |t          j        | |                    r| S  |t          j        | d                    r$ |t          j        | |                     r| |z   S  |t          j        | |                    r|S  |t          j        | |                     rdS d S )Nc                J    t           j        j                            |           S r   )rS   r_   r   r)  r   s    rg   <lambda>z5slice_.<locals>.compute_slice_index.<locals>.<lambda>+  s    qw'66q99 ri   r   )r   rS  GeLeLtGt)r_  r  r  r}   s       rg   compute_slice_indexz#slice_.<locals>.compute_slice_index'  s   =N99U##|D!!2ehua  !! 	bb%)>)>&?&? 	LR""## 	58ED5+A+A(B(B 	4<R%%&& 	KR&&'' 	1tri   NNFclampunbacked_bindingsr-   r  storage_offsetTr   )&%torch.fx.experimental.symbolic_shapesrg  r&   rp   rC   _validate_dimr   r   rS  r   rS   r_   r   statically_known_leq	TypeErrorr2   	SliceViewr  rr   	shape_envr`   r  r   r$  pytreeSequenceKeyDynamicSliceSizeregister_bufferr   register_operationmaybe_get_layoutr  
get_layoutoffset
get_strideDynamicSelectStorageOffsetrn   
as_strided)r   r   startendsteprr  rg  r&   r  ro  start_index	end_indexambiguous_slicers  sym_sizesym_storagesymkeypathb_sizer]  new_storage_offset	b_storage	new_sizesnew_stridess                           rg   slice_r    s4          
 a#####
3
"
"C::<<D<DdEJ''94!888T883QJJ 55dC@@ 		H      " (KO $))%q99''T488	"y'<$i3E#O  
LUCUKK
 
 	
 LL511	"AG$8$=>Q$R  ((( !!Q&&&(9&&&&Hk)//11  W}}V,,f.@.E.EFFFHH'788:::K)"3"33 	

S F '))&11FKGv&&&H#			"""\\^^2[1<<>>RUCV5VV1LLNN!LLNN3JJLL
 
 
	 00;;		""9---(QZZ\\""Iq||~~&&KIcNaK1CDDDs   
2B> >
C
Cc           	     h   d }d }t          | t                    rSt          | j        t          j                  r4|                                 }| j        }| j                                        } |                                  t          j	        |           st          d|  d          t          j        |           \  }}t          j        |r|n|j        |r|n|j        d |D             d |D             t          j        |pd                    }t          t          j        ||                    S )Nzunrealized as_strided(z, ...)c                6    g | ]}t          j        |          S r   r   rS  r   s     rg   r   zas_strided.<locals>.<listcomp>  s     '''Qa'''ri   c                6    g | ]}t          j        |          S r   r  r   s     rg   r   zas_strided.<locals>.<listcomp>       )))Qa)))ri   r   rr   layout)rp   rC   rr   r2   r6   rs   r   unwrap_viewr  is_storage_and_layoutr   as_storage_and_layoutFixedLayoutrz   r   rS  r   )	r   r  stridert  
new_device	new_dtypestorage
old_layout
new_layouts	            rg   r  r    s/   JI!Y 
!Jqvr{$C$C 
! \\^^
G	F  IIKKK#A&& F!"D1"D"D"DEEE2155GZ 7

j&74		J$4''$'''))&)))^(q)) J R'WZHHHIIIri   c                l    t          | t                    sJ t          | |||          j        | _        | S r   )rp   rC   r  rr   )r   r  r  rt  s       rg   as_strided_r    s5    a#####488=AFHri   c                D    t          | |||          }t          |          S r   )r  r  )r   r  r  rt  results        rg   as_strided_copyr    s!    488F==ri   c                    g d} D ]B}                     |||                                         z   f           d         d         }Cd  D              fd}t           d                                                   }d         d         |<   t          j         d                                          d                                         ||          S )Nr   rr  c                6    g | ]}|                                 S r   rt  rq  s     rg   r   z!pointwise_cat.<locals>.<listcomp>  s"    :::Ccoo'':::ri   c           	       
 t          j        |          t          j                  }g }g }t	          t                              D ]X

dk    rt          j        dt          j                  n*t          j        
         d         t          j                  }t          j        
         d         t          j                  }t          j        ||          }t          j        ||          }
dk    r|}n.
t                    dz
  k    r|}nt          j	        ||          }|
                    |           t          |           t                   
         d         z
            <   |
                    t          j        |
fdd                     Z|d         }	t	          t                    dz
  dd          D ]$
t          j        |
         |
         |	          }	%|	S )Nr   r/   c                 &                         S r   r   )rw   idx_loadinputs_loaderss   rg   rj  z1pointwise_cat.<locals>.inner_fn.<locals>.<lambda>  s    -N1-h77 ri           rr  r-   )rR   
index_exprr   int64r+  r   r  geltand_rt   rn   r*   maskedr  )rO  idx_dimmasksmasked_loadsr  r  
start_condend_condmasknext_valrw   r  r   rg  r  inputs_rangess             @@rg   r  zpointwise_cat.<locals>.inner_fn  s   .S5;77s6{{##  	  	A 66 Q,,,^M!$4Q$7EE 
 .q!1!!4ekBBC//Jvgs++HAvvc&kkAo%%!x
H55LLCyyH %Xc]]15Ea5H%HIIHSM
777777      #Fq("b11 	 	AyaQ HH
 ri   r  )rt   r   rn   r@   r  rs   r   )rg  r   prev_endr   r  r]  r  r  s   ``    @@rg   pointwise_catr    s   9;MH ) )h3<<>>#3F(FGHHH $R(::6:::N. . . . . . . .` F1I&&(())H!"%b)HSMay##%%Qi!!##	   ri   r  scaleszero_pointsaxisr   	quant_min	quant_maxc           	     v  	
 t                                                    dk    s
J d            t                                                    dk    s
J d            |                                 t          j        k    rt          | t          j                  } |                                 t          j        k    sJ d|                                              t          |                                           k     s,J dt          |                                                        |                                                                 	                                
	
f	d}t          j	        | 
                                ||                                           S )Nr/   expect scales 1 dimexpect zero_points 1 dim<Expecting input to have dtype torch.float32, but got dtype: Expecting axis to be < c                <  	 | 
         f} |           } |          } |          }t          t          j                  \  }}j        t          j        k    rt	          j        |t          j                  }j        t          j        k    rt	          j        |t          j                  }t	          j        |          }t	          j        ||z            |z   }t	          j	        |t	          j
        ||                    }	t	          j        |	          S Nr   )_create_constantsr   r  r   rR   r  r  
reciprocalrC  maximumminimum)rO  channel_idxr  scale
zero_pointqminqmax	inv_scaler2  clampedr  r   input_loaderr  r  r  scales_loaderr  zero_points_loaders             rg   r  z;quantized_decomposed_quantize_per_channel.<locals>.inner_fn
  s    4ylS!!k**''44
&y)5=QQQ
d<5=((L66E++j%+>>JN5))	i	)**Z7+dCKc$:$:;;|GU+++ri   r  )r   r   r   r   r  r  r  ru  r@   r  rs   )r  r  r  r  r  r  r   r  r  r  r  s    `````` @@@rg   )quantized_decomposed_quantize_per_channelr    s    v  !!Q&&&(=&&&{##%%&&!+++-G+++EN**..??---ZuGXGXZZ .-- #enn&&'''''9#enn&6&6"7"799 ('' $$&&L&&((M$0022, , , , , , , , , , , , ," !!~~	   ri   c           	     P                                       t           t          j                    fd}t	          j                                                                          |t           	                                                    }|                                  |S )Nc                    t           j                                        5  t          j                                         |                     cd d d            S # 1 swxY w Y   d S r   )r2   ComputedBufferforce_realizerR   device_assert_asyncru  )r_  r   r   s    rg   r  z_assert_async.<locals>.inner_fn'  s    ,,.. 	K 	K*+=4+;+;+=+=e+D+DcJJ	K 	K 	K 	K 	K 	K 	K 	K 	K 	K 	K 	K 	K 	K 	K 	K 	K 	Ks   0AA!$A!r  )
r  r  r   r\   r@   r  rs   r   rn   r   )r   r   r  assertion_ops   ``  rg   _assert_asyncr  #  s    LLNNND%*%%DK K K K K K #  nnDMMOO$$	  L ri   c                "    t          | |          S r   r  r   s     rg   lower_assert_asyncr  5      s###ri   c                "    t          | |          S r   r  r   s     rg   lower_assert_functional_asyncr  :  r  ri   )r  r  Optional[torch.dtype]c                 	
 t                                                    dk    s
J d            t                                                    dk    s
J d            |                                 |k    s"J d| d|                                              t          |                                           k     s,J dt          |                                                        t          j        |                                 	                                
                                	
fd}t          j        |                                 ||                                           S )	Nr/   r  r  Expecting input to have dtype , but got dtype: r  c                   |          f} |           } 
|          } |          }	j         t          j        k    rt          j        |t          j                  }j         t          j        k    rt          j        |t          j                  }t          j        t          j        |t          j                  |          |z  }t          j        |          }|S r   )r   r   r  rR   r  sub)rO  r  r  r  r  r2  r  r  r  r  r  r  r  s         rg   r  z=quantized_decomposed_dequantize_per_channel.<locals>.inner_fn]  s    4ylS!!k**''44
<5=((L66E--j%-@@Jgcl5%-88*EEMl3	**
ri   r  	r   r   r   r   r  ru  r@   r  rs   )r  r  r  r  r  r  r   r  r  r  r  r  s    ```   ` @@@rg   +quantized_decomposed_dequantize_per_channelr  ?  s    v  !!Q&&&(=&&&{##%%&&!+++-G+++??%%%TTTARARTT &%% #enn&&'''''9#enn&6&6"7"799 ('' M	$$&&L&&((M$0022           !!~~	   ri   r  r[  r  c                   |                                  t          j        k    rt          | t          j                  } |                                  t          j        k    sJ d|                                               |                                 fd}t          j        |                                 t          j
        |t          |          t          |                    |                                           S )Nr  c                J    
|           }t          d|z  |t          j                  \  }}t          j        ||z            |z   }t          t          j                  \  }}t          j        t          j        ||          |          }t          j        |	          S )N      ?r   )r  r   r  rR   rC  r  r  r  )rO  r  r  r  r  r2  r  r  r  r   r  r  r  s            rg   r  zBquantized_decomposed_quantize_per_tensor_default.<locals>.inner_fn  s    S!! 1%K5=!
 !
 !
	: i	)**Z7&y)5=QQQ
d+ck#t44d;;|GU+++ri   r  r  r  )r   r   r  r  r  ru  r@   r  rs   r6  rI  r[  r   r   )r  r  r  r  r  r   r  r  s      ``` @rg   0quantized_decomposed_quantize_per_tensor_defaultr  t  s
    EN**..??---ZuGXGXZZ .-- $$&&L, , , , , , , , !!"E%LLS__
 
 
 ~~   ri   c                  |                                  |k    s"J d| d|                                               t          j        |                                 fd}t	          j        |                                 t          j        |t          |          t          |                    |                                           S )Nr  r  c                     |           }t          ||t          j                  \  }}t          j        t          j        |t          j                  |          |z  }t          j        |          }|S r  )r  r   r  rR   r  r  )rO  r  r  r  r2  r  r  s        rg   r  zDquantized_decomposed_dequantize_per_tensor_default.<locals>.inner_fn  sg    S!!-eZu}UUUzgcl5%-88*EEMl3	**
ri   r  r  )r   r   r  ru  r@   r  rs   r6  rI  r[  r   r   )	r  r  r  r  r  r   r  r  r  s	         ` @rg   2quantized_decomposed_dequantize_per_tensor_defaultr    s     ??%%%TTTARARTT &%% M	$$&&L      !!"E%LLS__
 
 
 ~~   ri   c                  	 |                                  t          j        k    rt          | t          j                  } |                                  t          j        k    sJ d|                                               t                                                    dk    sMt                                                    dk    r                                d         dk    s
J d            t                                                    dk    sMt                                                    dk    r                                d         dk    s
J d            |                                                                                                 		fd}t          j	        | 
                                ||                                           S )Nr  r   r/   expect scale as scalar tensor"expect zero_point as scalar tensorc                    	|           } t                                                    dk    rdnd          } t                                                    dk    rdnd          }j        t          j        k    rt          j        |t          j                  }j        t          j        k    rt          j        |t          j                  }t          j        |t          j        |          z            |z   }t          
t          j                  \  }}t          j
        t          j        ||          |          }t          j        |          S )Nr/   r   r   r   )r   r   r   r   r  rR   r  rC  r  r  r  r  )rO  r  _scale_zero_pointr2  r  r  r  r   r  r  r  r  scale_loaderr  zero_point_loaders           rg   r  zAquantized_decomposed_quantize_per_tensor_tensor.<locals>.inner_fn  s#   S!!c%..*:*:&;&;q&@&@ddbII''ENN4D4D0E0E0J0JPRSS;%-''\&%-88Fu},,,{EMBBKiv 6 6677+E&y)5=QQQ
d+ck#t44d;;|GU+++ri   r  )r   r   r  r  r  r   r   ru  r@   r  rs   )
r  r  r  r  r  r   r  r  r  r  s
    ````` @@@rg   /quantized_decomposed_quantize_per_tensor_tensorr    s    EN**..??---ZuGXGXZZ .-- u~~  A%%ENN""u~~'7'7':a'?'?'?& (@'? z""$$%%**J!!""a''J,?,?,A,A!,D,I,I,I+ -J,I $$&&L$$&&L"..00, , , , , , , , , , , , !!~~	   ri   c               j  	
 t                                                    dk    sMt                                                    dk    r                                d         dk    s
J d            t                                                    dk    sMt                                                    dk    r                                d         dk    s
J d            |                                 |k    s"J d| d|                                              t          j        |                                                                 	                                
	
fd}t          j        |                                 ||                                           S )	Nr   r/   r  r  r  r  c                J    |           } t                                                    dk    rdnd          } 
t                                                    dk    rdnd          }j        t          j        k    rt          j        |t          j                  }	j        t          j        k    rt          j        |t          j                  }t          j        t          j        |t          j                  |          |z  }t          j        |          }|S )Nr/   r  r   )r   r   r   r   r  rR   r  r  )rO  r  r   r  r2  r  r  r  r  r  r  s        rg   r  zCquantized_decomposed_dequantize_per_tensor_tensor.<locals>.inner_fn  s    S!!c%..*:*:&;&;q&@&@ddbII''ENN4D4D0E0E0J0JPRSS;%-''\&%-88Fu},,,{EMBBKgcl5%-88+FFOl3	**
ri   r  r  )r  r  r  r  r  r   r  r  r  r  r  s    ``   ` @@@rg   1quantized_decomposed_dequantize_per_tensor_tensorr    s    u~~  A%%ENN""u~~'7'7':a'?'?'?& (@'? z""$$%%**J!!""a''J,?,?,A,A!,D,I,I,I+ -J,I ??%%%TTTARARTT &%% M	$$&&L$$&&L"..00
 
 
 
 
 
 
 
 
 
 !!~~	   ri   c                h  
 | d                                          j        dk    }|rt          d | D                       rt| D ]}|                                 t          d | D                       rt	          t
          j        g| R  \  } } t          t
          j        j                  | |          S t          |           dk    rt          | d                   S t          | d         |d          }t          | dt          j        ifd| D             } ddd fdt          fd| D                       }dfdt           j        rt%          | |          S |r-t'          t(          j                            | |                    S fdd}d
ddt          |           |k    s9t          |           t           j        k    rt          
fd| D                       rt          fdt0          j        j        D                       }t          fd| D                       o|}t          fd| D                       ot          fd| D                        }	|s|	r|st%          | |          S t'          t(          j                            | |                    S ) Nr   r   c              3  n   K   | ]0}|                                 t          j        t          j        fv V  1d S r   )r   r   int8uint8r   r  s     rg   r   zcat.<locals>.<genexpr>'  sK        ;@ej%+66     ri   c              3  `   K   | ])}t          |                                          d k    V  *dS )r   N)r   r   r  s     rg   r   zcat.<locals>.<genexpr>.  s9      >>es5>>##$$)>>>>>>ri   r/   r   c                0    g | ]}t          |          S r   r  )r   r   r   s     rg   r   zcat.<locals>.<listcomp>9  s#    555shsE""555ri   r   Union[TensorBox, ir.StorageBox]r[   	ir.IRNodec                    t          | t                    r?t          | j        t          j                  r| j                                        S | j        S t          | t          j                  r| j        S | S r   )rp   rC   rr   r2   r6   r  r   r   s    rg   unwrap_tensorzcat.<locals>.unwrap_tensor;  se    a## 	!&"+.. v))+++va'' 	6Mri   c                t    t          | t          j                  ot          | j        t          j                  S r   )rp   r2   r  rr   rA   )r{   s    rg   is_reductionzcat.<locals>.is_reductionG  s(    !R.//TJqvr|4T4TTri   c                   t          | t          t          j        f          r  |                     S  |           pFt          | t          j                  o,t          fd|                                 D                       S )Nc              3  h   K   | ],} t           j                            |                    V  -d S r   )rS   r_   
get_buffer)r   readcan_fuse_reductions     rg   r   z2cat.<locals>.can_fuse_reduction.<locals>.<genexpr>P  sR         #"17#5#5d#;#;<<     ri   )rp   rC   r2   r   r@   rD  get_read_names)r{   r  r  r  s    rg   r  zcat.<locals>.can_fuse_reductionJ  s    a)R]344 	8%%mmA&6&6777LOO !R\**     ,,..    	
ri   c              3  .   K   | ]} |          V  d S r   r   r   r{   r  s     rg   r   zcat.<locals>.<genexpr>W  s/      BBa..q11BBBBBBri   r\   c                D   t          j        |           r9t          j        | d          \  }}t           j                            |           S t          | t          t           j        f          r  |                     S t          | t           j                  rdS dS )NF)freezeT)	r2   r  r  ConcatKernelcan_realize_into_without_copyrp   rC   r   r@   )r   r  _should_lower_cat_inputr  s      rg   r#  z#cat.<locals>.should_lower_cat_inputY  s     #A&& 	N1!EBBBJGQDDWMMMMa)R]344 	<))--*:*:;;;a&& 	4uri   c                `   t          | t          t          j        f          r  |                     S t          | t          j                  sdS |                                 j        }|                                 D ]-}| t          j	        
                    |                    z  }.|S r  )rp   rC   r2   r   r@   inner_fn_opcountnum_opsr  rS   r_   r  )r   countr  op_countr  s      rg   r(  zcat.<locals>.op_countp  s    a)R]344 	.8MM!,,--- !R\** 	1""$$,$$&& 	8 	8DXXag0066777EEri   r   r-   rb   r1  c                J    | t           j        j        t           j        j        fv S r   )r  catr  constant_pad_ndrb   s    rg   additional_pointwise_opsz%cat.<locals>.additional_pointwise_ops  s    dh&(<(DEEEri   c              3  6   K   | ]} |          k    V  d S r   r   )r   r{   MAX_SIMPLE_OP_COUNTr(  s     rg   r   zcat.<locals>.<genexpr>  s1      CCq22CCCCCCri   c              3  8   K   | ]}t          |          V  d S r   )rJ   )r   user-  s     rg   r   zcat.<locals>.<genexpr>  sB       
 
 S":;;
 
 
 
 
 
ri   c              3  .   K   | ]} |          V  d S r   r   r   r   r#  s     rg   r   zcat.<locals>.<genexpr>  s/      >>&&s++>>>>>>ri   c              3  .   K   | ]} |          V  d S r   r   r3  s     rg   r   zcat.<locals>.<genexpr>  s@       "
 "
,/""3''"
 "
 "
 "
 "
 "
ri   c              3  .   K   | ]} |          V  d S r   r   r  s     rg   r   zcat.<locals>.<genexpr>  s/      <<((++<<<<<<ri   )r   r  r[   r  r[   r\   rb   r1  )rs   r   rC  r  require_channels_lastr  r*  r  r  r   r  rv  r   r   rd  rD  r0   force_pointwise_catr  rC   r2   r   r  max_pointwise_cat_inputsrS   r`   ra   )rg  r   
cpu_devicer  r"  fusable_reductionMAX_COMPLEX_POINTWISE_CATpointwise_usesfuse_pointwise_usehorizontal_fuse_catr/  r-  r  r   r  r(  r#  r  s             @@@@@@@@rg   r*  r*  $  s   %%'',5J 	?c  DJ     	?
  	 	EMMOOOO>>v>>>>> 	A-dh@@@@IFA1 011&#>>>
6{{aVAY
q	3
*
*C	%D%L E 6555f555F
 
 
 
U U U

 

 

 

 

 

 

 BBBB6BBBBB       ! *VS)))  >//<<===     $ !"F F F F 6{{///	V7	7	7CCCCCFCCCCC 
8  
 
 
 
~+
 
 
 
 
 >>>>v>>>>>Q> 	 " "
 "
 "
 "
39"
 "
 "
 
 
 =<<<<V<<<<<< 	  	."5 	.>O 	. ---R_++FC88999ri   r  dim1dim2c                  	
 |                                  
t          
          	t          	          t          	          t          k    fd           t          j        j                            t          j	        |d                    }|rXt          j        j        
                    t          j        j                            
         |z   
                   d          }nWt          j        j        
                    t          j        j                            
         
         |z
            d          }d|r| dfnd|ffdt          
          D             }|                    |           	
fd}t          t          j                            | ||                    S )N)rO  rankc                     d  d S )Nz(diagonal dimensions cannot be identical z, r   )rA  rB  s   rg   rj  zdiagonal.<locals>.<lambda>  s    WWWQUWW ri   r   )r   r   c                &    g | ]\  }}|fv|S r   r   )r   rw   r   rA  rB  s      rg   r   zdiagonal.<locals>.<listcomp>  s+    NNN41a$8M8MQ8M8M8Mri   c                
   | d         }dgt          	          z  }d}t                    D ]<}|k    r|d         z   ||<   |k    r|d         z   ||<   ,| |         ||<   |dz  }=|t          	          dz
  k    sJ |S )Nrr  r   r/   r-   )r   r+  )
rO  diag_idxoriginal_idxcur_dimr(  base_idxrA  rB  num_dimsoriginal_shapes
        rg   	reindexerzdiagonal.<locals>.reindexer  s    r7sS000x 	 	ADyy"*Xa["8Qd"*Xa["8Q"%g,Q1#n--11111ri   )r   r   r   r   rS   r_   r   evaluate_exprr   rm  evaluate_maxevaluate_minro   rt   rC   r2   GenericViewr  )r  r  rA  rB  offset_negative	diag_sizer  rN  rK  rL  rM  s     ``    @@@rg   diagonalrU    s   ^^%%N>""H8444D8444D	WWWWW   g&44UXfa5H5HIIO 
G$11G))t$v-~d/C  	
 
		 G$11G))t$nT&:V&C  	
 
	 H GQ<v;NNNNN9^44NNNE	LL          R^**5%CCDDDri   c                @    t          t          | |||                    S r   )r  rU  )r  r  rA  rB  s       rg   diagonal_copyrW    s    %t44555ri   c                h    t          |           }t          ||||          }t          ||           |S r   )r  rU  	mutate_to)r  srcr  rA  rB  rV  rc   s          rg   diagonal_scatterr[    s5    5\\FffdD11FfcMri   c                   t          j        |          }t          j        |                                 |                   }d }t          j        j                            t          j        |d                    r||z   }n9t          j        j                            t          j        |d                    r|}|t          |          ry| 
                                 |                                 }|                                 }|                                 j        ||         |z  z   }||= ||= t          | |||          S t          | |||dz   d          }t!          ||          S t#          t          j        j        j        t          j        j        j        d                   }	|	J t+          |	          dk    s
J |	            t-          t/          |	                                                    \  }
}| 
                                 |                                 }|                                 }|
}t3          j        |
||                                 j        ||         |                                 |         d          }t          j                            |          |_        t          j                            |           ||= ||= t          | |||          S )Nr   r/   Frq  rs  )r   rS  r   rS   r_   r   r)  rm  rk  r%   r  r  r  r  r  r  r-  r&   rz  r`   r  r   re  iterr$  r2   r  r~  r   r  )r   r   rO  r  actual_indexr]  
new_strider  slice_resultrs  unbacked_offset_symr"  buffers                rg   selectrc    s   
,s

C<

S)**DLw&&uxQ'7'788 Tz	
		(	(#q)9)9	:	: $S)) 	. IIKKKzz||HJ!"!6C<9W!W3a:7IJJJ "!S,q8HPUVVVL<--- 2	"AG$8$=>Q$R  ((( !!Q&&&(9&&&!$'8'>'>'@'@"A"ABB IIKKKzz||HJ,*	3	

S  F '))&11FKGv&&&3a:/ABBBri   c           
        t          | |d          }|}t          |t          t          f          sf|                                 |         }t
          j        j                            t          ||z   dz
  |                    }|g|z  }||dz
  |z  z
  |d<   g }d}|D ]0}||z   }	|
                    t          | |||	d                     |	}1|S )Nr   r/   rr  Frq  )rv  rp   rn   r   r   rS   r_   r   r&  r)   rt   r  )
r   r  r   sizes_x_sizechunksr  r  r  r  s
             rg   splitrh  &  s    
3
"
"CF edE]++ 3c"!++HVe^a5G,O,OPP6!vzU22r
 FE  dl 	fQUCu===>>>Mri   c                $    t          | ||          S r   )rh  )r   r  r   s      rg   split_with_sizesrj  A  s    E3ri   c                     t           d          t          j        j                                                                                } fdt          |          D             }|S )Nr   c                2    g | ]}t          |          S r   )rc  )r   rw   r   r   s     rg   r   zunbind.<locals>.<listcomp>J  s%    777AfQQ777ri   )rv  rS   r_   r   r&  r   r+  )r   r   rf  r  s   ``  rg   unbindrm  F  sa    
3
"
"CW''

S(9::F77777v777FMri   c                   |                                  }t          |          }t          ||          |dk    r t          t	          | d          |d          S |         }t
          j        j        }|                    ||           |	                    d           t          ||z
            dz   }|                    |          dk    r9|                     |                    t          ||z  |                               g |d          ||dz   d          |}	fd}
t          t          j                            | |	|
                    S )Nr   F)r  rr  r/   c                d    | d         |          z  z   }g | d          || dz   d         R S )Nrr  r/   r   )rO  dim_idxr   r  s     rg   rN  zunfold.<locals>.reindexerd  sF    b'CHtO+8TcT8G8c#'B,&7888ri   )r   r   r   r  rK  rS   r_   r   	check_leqcheck_ltr)   r   rG  r(   rC   r2   rR  r  )r   	dimensionr  r  r  ndimdim_sizer   new_dim_sizeout_sizerN  r   s      `       @rg   unfoldrx  N  sj   JJLLEu::D
4
+
+Cqyyi1oo4u====SzHwHtX&&&aHtOT22Q6L""8,,q00	''t0CX(N(NOO	
 	
 	
 EttDlDU3799-=DtDH9 9 9 9 9 9 R^**1h	BBCCCri   c                    t          | |d          }t          |                                           }|                    |t          j        j                   t          | |          S r  )rv  rn   r   insertr   rO  rP  r  )r   r   r,  s      rg   rK  rK  k  sR    
3
"
"CQZZ\\""IS%'+&&&9ri   c                    t          | |          }t          | t                    sJ t          |t                    sJ |j        | _        | S r   )rK  rp   rC   rr   r1  s      rg   
unsqueeze_r|  s  sK    
As

Ca#####c9%%%%%XAFHri   c                   t           j        j        j                            t          j        |                    }t          |                                           }|dk     r|||z   z  }d|cxk    r	||z   k     sn J |S r  )	rS   r_   r   rz  rO  r   sympifyr   r   )r   r   r  rt  s       rg   rv  rv  |  s    
'

$
2
25=3E3E
F
FCqzz||D
Qwwtf}####dVm######Jri   rr  c                2   t          | |d          }t          j        j                            |                                 |                   dz  }t          | |d|d          }t          | |||dz  d          }t          |t          |                    S )Nr   r-   Frq  )	rv  rS   r_   r   r&  r   r  r  sigmoid)r   r   new_lenr  rU  s        rg   glur    s    
3
"
"Cg((c):;;q@Gq#q'///Aq#w!5999Aq'!**ri   c                X     |rt                                            fd}d|_        |S )Nc                 d    d }t          j        |t          j        j        g| R i |          S )Nc                b    t          | t          j                  rt          j        |           n| S r   rp   r2   r:   rC   r  r   s    rg   wrap_tensorsz7fallback_handler.<locals>.handler.<locals>.wrap_tensors  )    *4Q	*B*BI9#A&&&Iri   )r{  tree_mapr2   FallbackKernelr  )rx   r  r  kernels      rg   handlerz!fallback_handler.<locals>.handler  sN    	J 	J 	J "+26KDKKKFKK
 
 	
ri   T)rA  r   _is_fallback_handler)r  r  r  s   `  rg   r  r    sE     f
 
 
 
 
 $(G Nri   c                 .    t          j        d           d S )NzjTorchinductor does not support code generation for complex operators. Performance may be worse than eager.)warningswarnr   ri   rg   _warn_complex_not_supportedr    s#    Mt    ri   r{   torch.Tensorc                   |                                  rt                       dS | j        rdS | j        rdS | j        t
          j        k    r|sdS t          |j        t
          j	        j
                  rI|j        t          j        j        t          j        j        t          j        j        t          j        j        fv p7t          |j        t
          j	        j
                  ot#          |j                   S dS )z0Do not support reading or writing to this tensorTF)r  r  is_meta	is_sparser   r   float8_e8m0fnurp   rc   r   r   r  r  r*  r  r  
_scaled_mmrK   )r{   re   s     rg   unsupported_input_tensorr    s    ||~~ #%%%ty t{ tw%&&& 	4
 t{EJ$9:: 	 
"'	Y 4;
(=>>W74;CWCW

 
	
 5ri   c                    t           j        j        t          j        j        j        j        f}||j        |v r| 	                                rdS t          | |          rdS | j        ot          j        S )z2Do not support writing tensor but can read from itNFT)r  r  r   r   rR   r  r  r  rc   r  r  is_cpur0   disable_cpp_codegen)r{   re   supported_complex_viewss      rg   unsupported_output_tensorr    sn     			,4 DK+BBBq||~~Bu4(( t8222ri   re   torch.fx.Nodec                     j         t          j        j        u rdS  j        dk    rdS  j         t          j        j        u rdS  fd}t          j         j        i  j	        D ]} ||d          r dS  | d          S )NFplaceholderc                6   t          | t          j        j                  sdS d| j        vrdS t          j        | j        d                   D ]K}t          |t          j        j                  s"|rt          |          r dS 8t          |          r dS LdS )NFr2  T)rp   r   fxNoder  r{  tree_leaves_subclasses
FakeTensorr  r  )inp_out_node	is_outputr  re   s      rg   check_skip_conditionzCfallback_node_due_to_unsupported_type.<locals>.check_skip_condition  s    ,66 	5)))5&|'8'?@@ 		  		 DdE$5$@AA   ,T488  44  ,D$77  44  uri   )r  T)
rc   r  view_as_complexr  rb   lift_fresh_copyr{  arg_tree_leavesrx   r  )re   allow_cpu_inputsr  r   s   `   rg   %fallback_node_due_to_unsupported_typer    s    {d*222uw-u {d*222u    * %ty@DK@@  u555 	44	  5555ri   c                   t           j        o.| t          j        t          j        j        t          j        j        hv }| t          vs|s|sJ d|              |rt          t          j
        d                    rt          | g          rt           j        r| t          j        j        j        v s[|sYt          j        j         j        r0dt          j        j         _        t&                              d           t+          d|  d          fd}t-          | t          j        j                  r4|                                 D ]}t5          | |          } ||           d S t-          | t          j        j        t          j        j        f          r ||            d S t;          d|  d	t=          |                      )
Nz*both a fallback and a decomp for same op: CIFznA make_fallback error occurred in suppress_errors config, and suppress_errors is being disabled to surface it.zmake_fallback(a.  ): a decomposition exists, we should switch to it. To fix this error, either add a decomposition to core_aten_decompositions (preferred) or inductor_decompositions, and delete the corresponding `make_fallback` line. Get help from the inductor team if unsure, don't pick arbitrarily to unblock yourself.c                    t          |            t          |             t          | d           t          |                     S Nr  )r   r   rJ  r  )op_overloadlayout_constraints    rg   register_fallbackz(make_fallback.<locals>.register_fallback*	  sU    !+...(!+/@AAAG $GGG[))
 
 	
ri   zUnsupported fallback z with type )r0   r  r  addcmul_foreach_addcmulScalar_foreach_addcdivr4   r\   osgetenvr5   fallback_randomr   _decompdecompositions_for_rngextra_random_decomps_dynamosuppress_errorslogwarningr   rp   r   r   r   r   r   HigherOrderOperatorRuntimeErrorr   )rb   r  r  override_decompskip_decomp_for_precisionr  olr  s    `      rg   make_fallbackr  	  s    !' > !2$$J D
 ^###:S##9R99 $#S 	
4!!
 t$$
 "
 em:OOO P =/ 	38EM 0KKH   fR f f f
 
 	

 
 
 
 
 "ej122 N,,.. 	+ 	+B!"b//Kk****	+ 	+ 
B.
0NO	P	P N"L2LL$r((LLMMMri   c                P    d}| D ]}||z  }t          |t          j                  S )z
    TorchInductor offset calculation differs from PyTorch eager offset
    calculation for random ops (tl.rand vs torch.rand). In future, we should
    strive for same impl for tl.rand and torch.rand.
    r/   r   tensorr   r  )rL  numelr   s      rg   philox_rand_offsetr  <	  s8     E  	%u{++++ri   c                f  	
 t          j        || t           j                            |                                                     
|                                |                                		
fd}t          j        ||t          |                     }t          |           }||fS )Nc                n   t          j         g           t          j                  }t          j         g           t          j                  }t          j        t          j         |           t          j                  |          }t          j        ||          }t          j        |          S r   )rR   r  r   r  r   r  rand)	r_  seed_index_exproffset_index_exprrand_index_exprr  r   offset_loader
random_posseed_loaders	        rg   r  zphilox_rand.<locals>.inner_fnU	  s     ,{{2DDLr):):EKHH'N::e,,ek::<M
 
 
 
 |FE***ri   r  )
r2   r  FlexibleLayoutcontiguous_stridesmake_indexerru  r@   r  rn   r  )r  seedr  r  rz   r   r  random_values_nodeoffset_noder  r  r  s        `   @@@rg   philox_randr  H	  s     
,,T22	 
 lnn  ""$$K&&((M+ + + + + + + + #)Dzz	   %T**K{**ri   c           	         t           j        rNt          j        t          j        t          j                            t          j	        j
        | ||                    S t          d          )Nz&should be handled in replace_random.py)r0   r  r{  r  rC   r  r2   r  r  native_dropoutr  r   )r   ptrains      rg   r  r  o	  sY     G$$T%8%@!QNN
 
 	

 EFFFri   c                t   t           j        s4|                                 t          j        d          k    s
J d            |                                  t          |          dk    st          |d         t                    rt          j
        j        nt          j
        j        }t          j        || g|R   | S )Nr   Tthis should be handled in decomps unless config.fallback_random or the device is CPUr   )r0   r  rs   r   rz   r  r   rp   r[  r  
bernoulli_Tensorr2   InplaceBernoulliFallback)r   rx   r  s      rg   r  r  z	  s    ! Q\\^^u|E7J7J%J%J%J^ &K%JJ IIKKK t99>>ZQ77> 	_# 
 Q66666Hri   c                    t           j        s4|                                 t          j        d          k    s
J d            t          t          |           g|R  S )Nr   r  )r0   r  rs   r   rz   r  r  )r   rx   s     rg   bernoulli_pr  	  sY    ! Q\\^^u|E7J7J%J%J%J^ &K%JJ eAhh&&&&&ri   c                    t           r   r   )r"  s    rg   _foobarr  	  s    
ri   c                :    t                               d           d S )Nz1using triton random, expect difference from eager)r  info)salts    rg   _warn_triton_randomr  	  s    HH@AAAAAri   c                 B    t          t          j        j                   d S r   )r  rS   r_   creation_timer   ri   rg   warn_triton_randomr  	  s    -.....ri   c                     |                     d          t          | i |S t          j        r"|                    dd            t          | i |S t          d          N	generatorz-should have been handled in replace_random.py)r  fallback_rand_generatorr0   r  popfallback_rand_defaultr   rx   r  s     rg   r  r  	  sh    zz+*&7777		 6

;%%%$d5f555
H
I
IIri   c                     |                     d          t          | i |S t          j        r"|                    dd            t          | i |S t          d          r  )r  fallback_randn_generatorr0   r  r  fallback_randn_defaultr   r  s     rg   randnr  	  sh    zz+*'8888		 7

;%%%%t6v666
H
I
IIri   c                j    t          j        |          }t           j                            | |          S r   )r2   get_stride_orderExternKernelrequire_stride_order)input_tensorr  stride_orders      rg   inductor_force_stride_orderr  	  s*    &v..L?//lKKKri   c                     t          d          )Nz.should be handled in fuse_seed_creation_pass()r  )rz   s    rg   inductor_seedr  	  s    
I
J
JJri   c                    t                       t          j        t          j        | t          |                              S r   )r  rC   r  r2   RandomSeedsrG   )r'  rz   s     rg   inductor_seedsr
  	  s3    BN5-2G2GHHIIIri   c                      fd}t          j                                                                          |g           S )Nc                R    t          j                                                  S r   )rR   	load_seedget_name)r"  r_  seedss    rg   r  z&inductor_lookup_seed.<locals>.inner_fn	  s    }U^^--u555ri   r  )r@   r  rs   r   )r  r_  r  s   `` rg   inductor_lookup_seedr  	  s`    6 6 6 6 6 6 !!oo	   ri   r  r  	list[int]r  modec                 	 t           j        rJ dv sJ g | } t          j        }|                                }t          j        ||| t
          j                            |           |          	                                |
                                		fd}t          j        |||g |           }|                                 |S )N)r  r  r  c                     t          t                     g           t          j         |           t          j                            S r   )r   rR   r  r   r  )r_  r  r  r  s    rg   r  z!inductor_random.<locals>.inner_fn	  sG    !wsD!!KOON::e,,ek::
 
 	
ri   r  )r0   r  r   r  rf  r2   r  r  r  r  ru  r@   r  r  )
r  r  r  r  r   rz   r  r  r  r  s
     `     @@rg   inductor_randomr  	  s    %%%%$$$$$T7DME%%''FtR.AA$GGPV  lnn  ""$$K
 
 
 
 
 
 
 w	  F NNMri   lowhighc               j   	 t           j        rJ g |}t          j        }|                                }t          j        |||t
          j                            |          |          	                                |
                                	 	fd}t          j        |||g |          S )Nr  c           	        t          j         g           t          j         |           t          j                  t          j        t          j                  t          j        t          j                            S r   )rR   	randint64r  r   r  r  )r_  r  r  r  r  s    rg   r  z"inductor_randint.<locals>.inner_fn
  s_    }KOON::e,,ek::N3,,N4--	
 
 	
ri   r  )r0   r  r   r  rf  r2   r  r  r  r  ru  r@   r  )
r  r  r  r  r  r   rz   r  r  r  s
   ``      @@rg   inductor_randintr  	  s     %%%%T7DKE%%''FtR.AA$GGPV  lnn  ""$$K
 
 
 
 
 
 
 
 w	   ri   tb.tuple[str, sympy.Expr, sympy.Expr, sympy.Expr]c                    |                                  }|                                 }t          d t          ||          D                       |d         z   }|                                 |d         ||d         fS )Nc              3  ,   K   | ]\  }}|d z
  |z  V  dS rV  r   )r   r   sts      rg   r   z%_boundaries_helper.<locals>.<genexpr>
  s.      AAeaa!er\AAAAAAri   rr  )r   r  sumr*  r  )r  r  r  
max_offsets       rg   _boundaries_helperr$  
  ss    
 ;;==D]]__FAAs4/@/@AAAAAF2JNJ
Rr
	 ri   tuple[str, sympy.Expr]c                ^    |                                  |                                 d         fS rJ  )r  r  r  s    rg   _sorter_helperr(  '
  s"    ;;=="--//"---ri   	out_int32rightsidesortersorted_sequenceselfr*  r+  r,  Optional[str]r-  Optional[TensorBox]c               x   
 d } |           r ||          r: |          s/ t          t          j        j        d           |||          S ||dk    rd|rt          j        nt          j        
|                                                                                                   t           
                                          dk    r

 fd}n	
 fd	}|                                }t          j        |
||j        
          }	|	                                 |	S )Nc                V    t           j                            | t          j                  S r   )rS   r_   r  r  	BUCKETIZEr'  s    rg   rj  zsearchsorted.<locals>.<lambda>5
  s     AG$7$7
N$% % ri   Fr  r)  r+  Tr/   c           
          |           }t          j        |t                    dd nt                    d nd          S )Nr   r-  sorter_indicesrR   	bucketizer$  r(  )rO  r2  index_dtyper+  r.  r-  values_loaders     rg   r  zsearchsorted.<locals>.inner_fnZ
  s`    -$$C="?33%~tt>&3I3I'-~tt1   ri   c                                 }d fd}t          j        |t                     |          d nt                    d n
 |                    S )Nr  rC   c                    |                                  }t          j        t          j        t
          j        d t          |d d         d d                   D                                 S )Nc              3  &   K   | ]\  }}||z  V  d S r   r   )r   r   rw   s      rg   r   zNsearchsorted.<locals>.inner_fn.<locals>.get_flattened_index.<locals>.<genexpr>q
  s*      &U&UAq1u&U&U&U&U&U&Uri   rr  )r  rR   r  r6  r  operatorr   r*  )r  stridesrO  r:  s     rg   get_flattened_indexz;searchsorted.<locals>.inner_fn.<locals>.get_flattened_indexm
  sj    --//~$ &U&UWSbS\3sPRs89T9T&U&U&U   	  ri   r6  )r  rC   r8  )rO  r2  rA  r:  r+  r.  r-  r;  s   `  rg   r  zsearchsorted.<locals>.inner_fnh
  s    -$$C       ="?33##O44%~tt>&3I3I'-~tt;N;Nv;V;V   ri   r  )r  r  searchsortedr  r   r  r  ru  r  r   r   rs   r@   r  rL  )r.  r/  r*  r+  r,  r-  validate_bucketizer  rz   r  r:  r;  s   `  ` `    @@rg   rB  rB  +
  s     //
!!$''
 '9'9&'A'AT 1 8eTTT
 
 
 	
 DGOO!*;%++K$$&&M 
?##%%&&!++
	 
	 
	 
	 
	 
	 
	 
	 
	 
		 	 	 	 	 	 	 	 	0 __Fz	  F NNMri   r*  r+  
boundariesc                  t                                                    dk    sJ t          j                            | t
          j                  r*t          j                            t
          j                  s- t          t          j	        j
        d          | |          S                                  |                                 }|                                 |rt          j        nt          j        fd}t#          j        |||                                           }|                                 |S )Nr/   Fr  rD  c                h     |           }t          j        |t                    d          }|S r  )rR   r9  r$  )r_  r2  r  rE  r:  r  r+  s      rg   r  zbucketize.<locals>.inner_fn
  sA    l5!!-z**
 
 ri   r  )r   r   rS   r_   r  r  r4  r  r  r9  r  r  rs   ru  r   r  r  r@   r  )	r  rE  r*  r+  rz   r  r  r:  r  s	    ` `   @@rg   r9  r9  
  sW    z""$$%%**** 	
E>#;<<
G
N,DEE
 R 55QQQ:%
 
 
 	
 F$$&&L!*;%++K
 
 
 
 
 
 
 
 ~~	  F NNMri   c                r    t          j        t          j        t          j        j        ||f          \  }}||fS r   )r{  tree_map_onlyr2   r:   r  require_stride1r"  rx   r  s      rg   require_denserL  
  s6    '
	2?2T6N LD& <ri   c                r    t          j        t          j        t          j        j        ||f          \  }}||fS r   )r{  rI  r2   r:   r  require_contiguousrK  s      rg   rN  rN  
  s6    '
	2?5f~ LD& <ri   c                r    t          j        t          j        t          j        j        ||f          \  }}||fS r   )r{  rI  r2   r:   r  r   rK  s      rg   r   r   
  s8     '
	2?=f~ LD& <ri   c                r    t          j        t          j        t          j        j        ||f          \  }}||fS r   )r{  rI  r2   r:   r  r8  rK  s      rg   r8  r8  
  s6    '
	2?84. LD& <ri   c                     S t          t                    r S t           t          j                  r2t          j                                                                       S t           t                    r fd D             S t           t          t          f          r0 t                     d t                     D                       S  S )Nc                J    i | ]}|t          |         |                    S r   constrain_to_fake_tensor)r   keyr   fake_args     rg   r   z,constrain_to_fake_tensor.<locals>.<dictcomp>
  s.    VVV3-c#hFFVVVri   c              3  <   K   | ]\  }}t          ||          V  d S r   rS  )r   r  f_as      rg   r   z+constrain_to_fake_tensor.<locals>.<genexpr>
  sB       
 
19!S$Q,,
 
 
 
 
 
ri   )rp   r   r2   r:   r  require_exact_stridesr  r8  r   rn   r   r*  )r   rV  s   ``rg   rT  rT  
  s    
(,-- 
#ry!! M44S(//:K:KLLL#t 
VVVVVRUVVVV	C%	'	' 
tCyy 
 
=@h=O=O
 
 
 
 
 	
 Jri   c                    t          d t          | |          D                       } fd|                                D             }| |fS )Nc              3  <   K   | ]\  }}t          ||          V  d S r   rS  )r   r   rV  s      rg   r   z,constrain_to_fake_tensors.<locals>.<genexpr>
  sD        C 	!h//     ri   c                D    i | ]\  }}|t          ||                   S r   rS  )r   r  r  fake_kwargss      rg   r   z-constrain_to_fake_tensors.<locals>.<dictcomp>
  s.    XXXAa)![^<<XXXri   )r   r*  r$  )rx   r  	fake_argsr]  s      `rg   r   r   
  sf       y11    D YXXXXXXF<ri   c                     fdt          fdt          | j                  D                       } fd|                                D             }||fS )Nc                V    t           t          j                  rft          j        j        d                                         t          j        j        j	                  }t          j
                             |          S t           t                    r fd D             S  S )Nr2  c                B    i | ]}| |         |                   S r   r   )r   rU  apply_constraintr   fx_args     rg   r   zEconstrain_to_fx_strides.<locals>.apply_constraint.<locals>.<dictcomp>	  s1    PPPSC))#c(F3K@@PPPri   )rp   r2   r:   r   r  r  rS   r_   r   rz  r  r  r8  )r   rc  r  rb  s   `` rg   rb  z1constrain_to_fx_strides.<locals>.apply_constraint  s    c29%% 	K.E"))++QW-=-G L ?77\JJJc4   	QPPPPPPCPPPP
ri   c              3  6   K   | ]\  }} ||          V  d S r   r   )r   r   rc  rb  s      rg   r   z*constrain_to_fx_strides.<locals>.<genexpr>  sF        *5#vf%%     ri   c                F    i | ]\  }}| |j         |                   S r   r  r   r  r  rb  fx_nodes      rg   r   z+constrain_to_fx_strides.<locals>.<dictcomp>  s4    SSSDAqa!!!W^A%677SSSri   )r   r*  rx   r$  rh  rx   r  rb  s   `  @rg   r   r     s             9<T7<9P9P    D TSSSSFLLNNSSSF<ri   c           	           fdt          fdt          t          | j                            D                       } fd|                                D             }||fS )Nc           
     
   t          |t          j                  s|S |j        d         }d |                                D             }t
          j        j        j        }t          j	        |                                |          }|rT|d         dk    rHt          t          t          t          |                                                                        }j        t           j        j        k    r| dv rt          |          dk    sJ d}|j        s t          j                            ||          S dj        t,          j        j        j        j        k    o| d	k    }t          |t2                    sJ t          |                                          d
vr|S t          j        |          }|r2t          j        t          j                            |          |          S t          |t                    rH|                                4|r2t          j        t          j                            |          |          S |rt          |                                          }	g }
|                                }t          t          |                                          dz
            D ]o}t
          j        j                            ||         d          s-|@t
          j        j                            ||         d          r|
                    |           pdgt          |	          z  }d|d<   d}t          t          |	          dz
  dd          D ]}||dz            dk    r||	|dz            z  }||
v r7t
          j        j                            ||dz            z  d          rd||<   Zt
          j        j                            |z  d          stA          |          z  }|||<   t          j        !                    ||          S |r2t          j        t          j                            |          |          S t          |t                    rH|                                4|r2t          j        t          j                            |          |          S fd}t          |j"        t          j#                  rZ ||          sO ||$                                          r2t          j        t          j                            |          |          S t          j                            ||          S )Nr2  c                ^    g | ]*}t          |t          j                  r|j        j        n|+S r   )rp   r   SymIntre   exprr   s     rg   r   z=sdpa_constraint.<locals>.apply_constraint.<locals>.<listcomp>  s>     
 
 
BC:a66=AFKKA
 
 
ri   rr  r   )r   r   r   )r   r/   r-   r   r   r   r   r   r/   r-   c                    t           j        j                            t	          j        t          |                                 d                   d                    S )Nrr  r   )rS   r_   r   r)  r   r*  r+   r   )r   	ALIGNMENTs    rg   
is_alignedz=sdpa_constraint.<locals>.apply_constraint.<locals>.is_aligned  sF    7#22QZZ\\"-y991==  ri   )%rp   r2   r:   r  r  rS   r_   r   rz  r   rn   rN  r+  r   r   rc   r  0_scaled_dot_product_efficient_attention_backwardr  is_cudar  r  r   rR   '_scaled_dot_product_efficient_attentionrC   is_aligned_realized_tensortry_match_insignificant_stridesrealize_inputmaybe_get_stridestatically_known_equalsrt   rF   rY  rr   r6   r  )rO  r   rc  meta_valmeta_stride_exprrz  r  effn_attn_fwd_biasis_aligned_tensorrw  expanded_dimsmaybe_striderw   out_stridesr  rr  rq  rh  s                   @rg   rb  z)sdpa_constraint.<locals>.apply_constraint  sA   #ry)) 	J;u%
 
GOGXGX
 
 
 G$.	*8??+<+<iHH 	FL,11s3<<>>/B/B)C)C D DEEL NDLM Mv|$$))))
 (L 	K?77\JJJ 	 Ny~EMN q 	 #y)))))s||~~f,,J9#yII 	5--c224D  
 sF##	$$&&2! 3 5--c224D    $	KCLLNN++HM //11L3s||~~..233 , ,7#;;<LQ<OQRSS , ,(@@aRSTT - "((+++ $X.KKOF3x==1,b"55 ( (q1u%**#hq1uo5F %%w'??#AE*Y6  ! *+A w'??@RTUVV D$VY77)CF!'A?88kJJJ 	5--c224D  
 sF##	$$&&2! 3 5--c224D  	 	 	 	 	
 ch,, 	:c?? :coo//00 =55c::<L   33CFFFri   c              3  >   K   | ]\  }\  }} |||          V  d S r   r   )r   rO  r   rc  rb  s       rg   r   z"sdpa_constraint.<locals>.<genexpr>  sN        C#v 	c6**     ri   c           	     H    i | ]\  }}| d |j         |                   S rr  rf  rg  s      rg   r   z#sdpa_constraint.<locals>.<dictcomp>  s6    WWW1a!!"a):;;WWWri   )r   ro   r*  rx   r$  ri  s   `  @rg   sdpa_constraintr    s    |G |G |G |G |G|     "+Cgl,C,C"D"D    D XWWWWWWWF<ri   )r  )r  c                   t          |t          j                  s6t          ||                                 |                                           }|}|                                 |                                k    r"t          ||                                           }|                                 |                                k    r"t          ||                                           }|                                 |                                k    r1t          ||                                           }t          |          S t          |          S NrT  )rp   r2   r:   r  r   rs   r  r  r   rS  r  )r/  rZ  r  r   ru   s        rg   r  r  ~  s    c29%% LS 0 09J9JKKKACNN,,,,a**++~~3==??**Q(())}}#,,..((Q((Szz88Ori   )memory_formatc          	         t          j        |                                 |                                 |                                 t          |                                                     S Nr  )r@   r  rs   r   ru  rn   r   )r   r  s     rg   r  r    sR     ||~~kkmmAJJLL!!	   ri   c                   g }t          | t                    rt          | j        t          j                  rx| j        } t          | t          j                  rH|                    |                                            | j        } t          | t          j                  Ht          |           } t          |           } |r:| j        } |d d d         D ]}t          j        | |          } t          |           } | S )Nrr  r  )rp   rC   rr   r2   r   rt   r  r  )r   reinterpret_view_layoutsr  s      rg   clone_preserve_reinterpret_viewr    s    !!Y Jqvr7I$J$J FB.// 	$++ALLNN;;;A B.// 	 aLLaA F.ttt4 	: 	:F"&999AAaLLHri   r  c               b    fd}t          j        t          |          || g          S )Nc                H    t          j        | d         z  z             S )Nr   r   rR   r  )r_  r   r  r  s    rg   r}   ziota.<locals>.fn  s%    ~dU1Xo5UCCCCri   r  )r@   r  rG   )lengthr  r  r   rz   requires_gradr}   s    ```   rg   iotar    sa    D D D D D D D V$$x	   ri   r   r_  c           	        t          ||                                           }|                                 t          | d          t          j        j                            t          j	        d                    r| 
                                         z   nbt          j        j                            t          j        d                    rn* t          t          j        j                  | |          S t          j        j                            d           t          j        j                            | 
                                                    t%          t'          |          | 
                                          }|                                fd}t)          j        |                                 |                                 |t/          | 
                                                    S )Nr   c           	         t          j        t          j        t          j        |          t          j                  t          j        t          j                             |            |                     S r   )rR   r  eqr  r   r  )rO  r   r_  
src_loaderr[  s    rg   r  z select_scatter.<locals>.inner_fn  sd    yFs3x55uek22  JsOOHSMM
 
 	
ri   r  )r  r   ru  rv  rS   r_   r   r)  r   rm  r   rk  r  r  select_scatterr  rq  rr  rS  rK  r@   r  rs   rn   )r   rZ  r   r_  r  r  r[  s     `` @@rg   r  r    s   
3
&
&C}}H
3
"
"Cw&&uxq'9'9:: Q

S))	
		(	(%););	<	< Q = 3 ;<<QS%PPPGq%(((GeQZZ\\#%6777
3$$ajjll
3
3C""J
 
 
 
 
 
 
 
 ||~~kkmmAJJLL!!	   ri   c           	        	
 t          |                                           }                                 
t           d                                                    t
          j                                       \  t                                                     }t          z
  dz
  z             |<   t          ||          }|                                		 
fd}t          j                                                                          |t                                                               S )Nr   r/   c                2   	dk    rk    r
dk    r |           S t          j        |          t          j                  }t	          |           t          |          	z
  
          <   g }	dk    rW|                    t          j        |t          j        t          j	        	          t          j                                       k    rW|                    t          j
        |t          j        t          j	                  t          j                                       
dk    rz|                    t          j        t          j        t          |          	z
  d
          t          j                  t          j        dt          j                                       |sJ t          j        t           j        |          }t          j        |fdt%                    rdnd          }t          j        || |                     S )Nr   r/   c                                 S r   r   )src_idxr  s   rg   rj  z1slice_scatter.<locals>.inner_fn.<locals>.<lambda>  s    JJw'' ri   r  )rR   r  r   r  rn   r)   rt   r  r   rS  r  r  r,   r  r6  r  r  r  r   r  )rO  r  r  src_valr  r   ru  r  r  r  r  r   r[  s       @rg   r  zslice_scatter.<locals>.inner_fn  s   A::#//daii:c??".S5;77s))C5 0$77A::KKN5<#6#6DD    (??KKN5<#4#4ekBB    199KKN'C5(8!TBBEK  LEK00	    t$//*''''' ##,AA
 

 yHSMM
 
 	
ri   r  )r  r   ru  rv  r   r2   ry  normalize_start_endrn   r)   rS  r@   r  rs   )r   rZ  r   r  r  r  src_sizer  ru  r  r[  s   ` ````  @@@rg   slice_scatterr    sJ   
3
&
&C}}H
3
"
"Czz||C H 11!S%EEJE3AJJLL!!HS5[D1H5t<<HSM
h

C""J,
 ,
 ,
 ,
 ,
 ,
 ,
 ,
 ,
 ,
 ,
 ,
\ ||~~kkmmAJJLL!!	   ri   c                    t          | t          t          f          r(t          |           dk    rt	          | d                   S | S r  )rp   rn   r   r   _unwrapr   s    rg   r  r  0  s;    !dE]## A

qt}}Hri   r   rz   r  
pin_memoryc               :    t          |d t          j        fv d|            t          | d           t          t	                     t
                    rpt          j        npt          j                    g }t           t          j	                  r fd}nt           t          t
          f          r fd}nt                     dk    s5t           d         t          t
          f          rNt                     dk    r;|                    t          j        t                                           fd}n4t          j                            t          j         |                    S t%          j        t)          |          ||	          S )
Nlayout=r  c                .    t          j                  S r   r  r_  rr   r   s    rg   r  ztensor.<locals>.inner_fnC  s    >$...ri   c                .    t          j                  S r   rR   r  r  s    rg   r  ztensor.<locals>.inner_fnH  s    <e,,,ri   r   r   c                      fdt                    dk    rt          j        d          S  dt                              S )Nc           	     l   | |k     sJ || z
  dk    rt          j        |                    S || z
  dz  | z   }t          j        t          j        t          j        d         t
          j                  t          j        |t
          j                             | |           ||                    S )Nr/   r-   r   )rR   r  r  r  r  r   r  )r  r  midbinary_searchrr   r   r_  s      rg   r  z/tensor.<locals>.inner_fn.<locals>.binary_searchP  s    s{{{{;!##<UU;;;U{q(50yFuQx==S%+66  "M%--!M#s++  ri   r   )r   rR   r  )r_  r  rr   r   s   `@rg   r  ztensor.<locals>.inner_fnO  sh            4yyA~~|Au--- =CII...ri   rT  r  )r   r   stridedrp   r  r   r  get_default_dtyper   r   r[  r   rt   IntegerrS   r_   add_tensor_constantr  r@   r  rG   )rr   r   rz   r  r  r  r  s   ``     rg   r  r  6  s   v$..0B&0B0BCCC:~|,,,'$--%% 3$2022!F$$$ $
	/ 	/ 	/ 	/ 	/ 	/ 	/ 
D5#,	'	' 
	- 	- 	- 	- 	- 	- 	- 
Ta:d1gs|<<TaemCII..///	/ 	/ 	/ 	/ 	/ 	/ 	/( w**LU6:::
 
 	
 V$$	   ri   c                    t          | t                    r&|t          | |          } |t          | |          } | S t	          | ||          S r  )rp   rC   r  r  r  )rr   r   rz   s      rg   	as_tensorr  o  sX    $	"" D%((DT6**D$eF3333ri   c                8    t          | t          j                  S r  r  )rr   s    rg   long_tensorr  z  s    $ek****ri   c                   t          t          j        j        j        t          j        j        j        d                   }|J t          |          dk    s
J |            t          t          |
                                                    \  }}t          j        |||           }t          j                            |          |_        t          j                            |           t          j        j        j        d         }t!          |t"          j        t"          j        t"          j        f          r|j        j        S t/          j        |          S )Nrs  r/   r2  )r&   rS   r_   r   rz  r`   r  r   re  r]  r$  r2   DynamicScalarr~  r   r  rp   r   rm  SymFloatSymBoolre   rn  r   r~  )rr   rs  binding_symr  rb  r2  s         rg   _local_scalar_denser    s    2	"AG$8$=>Q$R  ((( !!Q&&&(9&&&   %6%<%<%>%> ? ?@@Kk7D99F'))&11FKGv&&& '

#E
*C#enemDEE "x}}S!!!ri   c                    d S r   r   )rr   r   s     rg   _assert_scalarr    s	     4ri   )rz   r  c                   d S r   r   )r  r  r  r   rz   r  s         rg   _assert_tensor_metadatar    s	     4ri   c                   | t          | t          t          f          st          d          rj        t          t          t          f          rfd}nat          t
          j                  rfd}n@t                                                    dk    sJ 	                                fd}t          j        ||t          |                    S )Nr  c                .    t          j                  S r   r  r_  r   r  s    rg   r  z_full.<locals>.inner_fn  s    <u---ri   c                .    t          j                  S r   r  r  s    rg   r  z_full.<locals>.inner_fn  s    >%///ri   r   c                     g           S r   r   )r_  value_loaders    rg   r  z_full.<locals>.inner_fn  s    <###ri   r  )rp   r   r[  r  r  r   r   r   r   ru  r@   r  rn   )
fill_valuerz   r   r  r  r  r  s     `  @@rg   _fullr    s'   Ej3,// GE74K4K %#u&& $	. 	. 	. 	. 	. 	. 	. 
E5;	'	' 
$	0 	0 	0 	0 	0 	0 	0 5>>##$$))))((**	$ 	$ 	$ 	$ 	$ Dzz	   ri   c                H     t          t          |                    | fi |S r   create_tensor_liketensor_constructor)r   r  r  s      rg   r7  r7    s,    =0<<==aJJ6JJJri   c                "     d d d d dd d fd
}|S )NF)namesr   rz   r  r  r  c                   t          | d u d           t          |d t          j        fv d|            t          | d           t          |          }|pt          j                    }t          |          dk    rBt          |d         t          t          t          j	        f          rt          |d                   }|D ]}t          |t          j
                  rJ d |D             }t          |||          S )Nnamed tensorsr  r  r/   r   c                6    g | ]}t          j        |          S r   r  r   s     rg   r   z5tensor_constructor.<locals>.inner.<locals>.<listcomp>  s     ...AQ...ri   )r   r   r  rG   r  r   rp   rn   r   Sizerm  r  )	r  r   rz   r  r  r  r  r   r  s	           rg   r  z!tensor_constructor.<locals>.inner  s     	5D=/2226dEM224Ff4F4FGGGz><000v&&2022t99>>ja4
2KLL>a>>D  	3 	3A!!U\222222.....Z555ri   r   )r  r  s   ` rg   r  r    sA     6 6 6 6 6 6 6. Lri   )r  r   r  rz   r  r  c                   t          | d u d           t          |          }t          |          dk    rBt          |d         t          t
          t          j        f          rt          |d                   }t          |d ||||          S )Nr  r/   r   r   r  rz   r  )	r   rG   r   rp   rn   r   r   r  empty_strided)r  r   r  rz   r  r  r  s          rg   r\  r\    s     u}o...6""F
4yyA~~*T!WtUEJ.GHH~T!W~~d%v*   ri   c                      dddddd fd
}|S )zZ
    Shim to convert X_like(...) into X(...).  For example zeros_like() into zeros().
    NF)r   rz   r  r  r  c               B   t          | d           t          |d t          j        fv d|            ||                                 }nt	          |          }|p|                                 }t          |                                           } |||||          S )Nr  r  r  )r   r   r  r   r   rs   rn   r   )r   r   rz   r  r  r  r  creation_fns          rg   _constant_likez*create_tensor_like.<locals>._constant_like  s     	z><0006dEM224Ff4F4FGGG=KKMMEE ''E)1<<>>AJJLL!!{fV

 
 
 	
ri   r   )r  r  s   ` rg   r  r    s<     duTX
 
 
 
 
 
 
 ri   c                :    t          t          |                     S r   r  )r  s    rg   r  r  %  s    0<<===ri   c                     d d d d d fd
}|S )Nr  c               l   t          |t          t          f          sJ t          | d           t          |d t          j        fv d|            t          |          p|                                 }|p|                                 }d |D             }t          t          |          ||          S )Nr  r  c                6    g | ]}t          j        |          S r   )r   r  r   s     rg   r   z7new_constant.<locals>._new_constant.<locals>.<listcomp>8  s"    ///Qa  ///ri   )rp   rn   r   r   r   r  r   r   rs   r  rG   )r   r  r   r  rz   r  r  s         rg   _new_constantz#new_constant.<locals>._new_constant/  s     $u.....z><0006dEM224Ff4F4FGGGU##4q{{}})1<<>>//$///Zv!6!6tDDDri   r   )r  r  s   ` rg   new_constantr  .  s>    tDT
E 
E 
E 
E 
E 
E 
E ri   r  c          	         ||                                  }||                                 }t          |d ||t          |          |          S Nr  r   rs   r  rG   )r   r  r   r  rz   r  s         rg   	new_emptyr  >  sW    }~V$$   ri   c               &   t          | t          t          f          sJ t          |t          t          t          d           f          sJ t	          |d t
          j        fv d|            t          |          pt          j                    }|pt          j	        d          j
        }t          |          }t          d|||           }|                                 |j        j        }t          j        |j        dgt#          |           z            |_        t          |t$          j                  sJ d | D             } |rd |D             nt$          j                            |           }t%          j        ||| ||pd	          |_        |S )
Nr  r  r   )r  rz   r   r  )r  c                6    g | ]}t          j        |          S r   r  r   s     rg   r   z!empty_strided.<locals>.<listcomp>_  s     ***ELOO***ri   c                6    g | ]}t          j        |          S r   r  r   s     rg   r   z!empty_strided.<locals>.<listcomp>a  r  ri   F)rz   r   r  r  	is_pinned)rp   rn   r   r   r   r   r  r   r  r  rz   rG   r  r  rr   dataclassesreplacer   r2   r  r  r  r  r  )r  r  r   r  rz   r  	pointwiserb  s           rg   r  r  N  s    dT5M*****ftUDJJ788888v$..0B&0B0BCCC<5#:#<#<E/u|C((/F6""F6TJJJI^ F%fk1#D		/JJJFKfb/00000**T***D 	8))&))))11$77 
 N%  FM ri   c          	         ||                                  }||                                 }t          ||||t          |          |          S r  r  )r   r  r  r   r  rz   r  s          rg   new_empty_stridedr  o  sY     }~V$$   ri   c                    d |D             }t          t          t          |                    |j                  }t          j                            | |          S )Nc                V    g | ]&}t           j        j                            |          'S r   r   r   s     rg   r   z copy_strided.<locals>.<listcomp>  s+    EEEag11!44EEEri   )rU  )sortedr+  r   __getitem__r2   r  r  )r   r  r  s      rg   copy_stridedr    sP    EEfEEEF%F,,&2DEEEL?//<@@@ri   c                l    |                     d          
J d             t          |          | fi |S )Nr   z(dtype should be handled by decomposition)r  r  )r  r  r  s      rg   fullr    sD    ::g**,V***)j))$99&999ri   c                D   t          | t                    sJ |                                dk    r"t          | |                                          S |                                 t                    dk    }t          | |          |rt          | dg          } dg|                                 |                                fd}t          j
        |                                 |                                 ||                                          S )Nr   r/   c                    t          |           } t          j         |                              }t          |           dk    r|g} n|| <    |           S r  )rn   rR   indirect_indexingr   )rO  
gather_idxr   index_loaderr  r[  s     rg   r}   zgather.<locals>.fn  s]    3ii*<<+<+<d3iHH
s88q==,CC!CHx}}ri   r  )rp   rC   	get_numelr  r   r   rv  rS  ru  r@   r  rs   r   )	r   r   r_  sparse_gradr  r}   r  r  r[  s	    `    @@@rg   gatherr    s     a#####AENN,,---::<<DYY!^F
3
'
'C 1qcNNs}}H$$&&L        ||~~kkmm~~	   ri   c                  	
 |r+ t          t          j        j                  | ||||          S |rJ t	          | t
                    sJ t	          |t
                    sJ dt          |                                          v sJ |                                 	|                                t          |
                                          | 
                                
g |
                                
dd          	
fd}t          j        |                                 |                                 |          S )Nr   r/   c                    t          |           t                    k    sJ |  d              | d                    }t          j        |d                   gg | d          z   } |          S )Nz != r   )r   rR   r  )rO  	var_index
weight_idxindices_loaderindices_ndimr]  weight_loaderweight_sizes      rg   r}   zembedding.<locals>.fn  s    3xx3x==(((S*@*@h*@*@((("N3}}#566	+I{1~FFG K
K
 

 }Z(((ri   r  )r  r  	embeddingr  rp   rC   r   r   ru  r   r   r@   r  rs   )weightr  padding_idxscale_grad_by_freqsparser}   r   r  r]  r  r  s         @@@@@rg   r  r    su    
7 677G[*<f
 
 	
 :fi(((((gy)))))C))++,,,,,,&&((M((**Nw''))**L//##K6!!##6k!""o6H) ) ) ) ) ) ) ) )   ""  	   ri   c                    t          d  D                       sJ dd  D                          t          d  D                       rt          d          d t                     D             }t	          |          dk    s
J d            d gt	                     z  }t          |t           fd	|D                        D ]1\  }}|                                |k    rt          d
          |||<   2||fS )Nc              3     K   | ]H}||                                 t          j        t          j        t          j        t          j        fv V  Id S r   )r   r   r  r  r\   r  r   rw   s     rg   r   z.check_and_broadcast_indices.<locals>.<genexpr>  sN        = 	
%+u{EJLL=== ri   z)indices must be int64, byte or bool. Got c                :    g | ]}||                                 S r   r   r  s     rg   r   z/check_and_broadcast_indices.<locals>.<listcomp>  s&    4e4e4eqWXWdQ[[]]WdWdWdri   c              3  r   K   | ]2}||                                 t          j        t          j        fv V  3d S r   )r   r   r\   r  r  s     rg   r   z.check_and_broadcast_indices.<locals>.<genexpr>  sA        78q}%*ek22}}}} ri   zFallback for bool indicesc                B    g | ]\  }}t          |t                    |S r   r  r  s      rg   r   z/check_and_broadcast_indices.<locals>.<listcomp>  s+    OOO1jI6N6NO!OOOri   r   z"requires at least 1 non-None indexc                     g | ]
}|         S r   r   )r   rw   r  s     rg   r   z/check_and_broadcast_indices.<locals>.<listcomp>  s    3S3S3S1GAJ3S3S3Sri   z.Fallback when indices is on a different device)rC  rD  r   ro   r   r*  r'  rs   )r  rz   
valid_idxsnew_indicesrw   r   s   `     rg   check_and_broadcast_indicesr    sb           
 	h4e4eG4e4e4egg     <C     ? ""=>>>OO	' 2 2OOOJz??Q D&3w<<'KJ 13S3S3S3S
3S3S3S TUU  1 <<>>V##%&VWWWA
""ri   c	           
     f    dt          j                  D ]\  }	}
|
|	z
  dk    rd fdt                    D             g  t                    t                    z   d          d         }rz   nd |         z   |d          z   f
d}|fS )NFr/   Tc                *    g | ]\  }}||         S r   r   )r   rw   r2  rf  s      rg   r   z2index_output_size_and_inner_fn.<locals>.<listcomp>  s"    MMMC6!9ri   r   c                v  
 t          |           t                    k    sJ t                    t          
          k    sJ t                    }g }d         }rdn|}d}t          d         dz             D ]}||k    r||z  }|         6|t          |           k     sJ |                    | |                    |dz  }K|         }|J 
|         }|                    t          j         || |||z                      |	                     g || |d          }|n
 |          S )Nr   rr  r/   r   wrap_neg)r   r+  rt   rR   r  )rO  rD  	new_indexfirst_tensor_indexstart_offsetnext_idxrw   loaderr  r   indexed_sizer  indices_loadersnon_consecutive_tensorsoutput_sizetensor_indicestensor_sizer  r[  s            rg   r}   z*index_output_size_and_inner_fn.<locals>.fn  s   3xx3{++++++?##s<'8'88888;	+A.3Kqq9K~b)A-.. 	 	AL  D qz!#c((****  X///A(+)))#A  )s<,2E#EFGG#!)	     

^
	 %,yy((92E2EEri   )r(  pairwisero   r   )rf  r  r"  r#  r  r  r[  r   r  previouscurrentr  r}   r   r!  s   `````````    @@rg   index_output_size_and_inner_fnr'    sM   $ $&/?? + +'X""&*#MMMM9W+=+=MMMKSKS&[)9)9C<O<O)O)Q)Q"RSK'* 
!K/ +++,,--./ 	 F  F  F  F  F  F  F  F  F  F  F  F  F  FD ?ri   c                    t          | ||          \  }}}t          j        |                                 |                                 ||          S r  )index_impl_helperr@   r  rs   r   )r   r  r   r!  r  r"  s         rg   
index_implr*  8  sP    0GUCCK1||~~kkmm	   ri   c                  
 t          t          t          f          sJ |                                 t	          |                                           \  }t          |          dk    s
J d            d D             }t          |d                                                            }|                                 fdt          t                              D             }|rd|v rd|vrt          d          fdt          t                              D             }t          ||||d ||	  	        \  }

fd}	||	
fS )	Nr   z Must have at least one valid idxc                >    g | ]}||                                 nd S r   rt  r  s     rg   r   z%index_impl_helper.<locals>.<listcomp>I  (    SSSa!-q}}TSSSri   c                0    g | ]}|         
|         S r   r   )r   rw   r  rf  s     rg   r   z%index_impl_helper.<locals>.<listcomp>Q  s&    UUU!gaj>TF1I>T>T>Tri   z0index is out of bounds for dimension with size 0c                     g | ]
}|         S r   r   r   rw   rf  s     rg   r   z%index_impl_helper.<locals>.<listcomp>U      ;;;!F1I;;;ri   r  c                ,      |                     S r   r   )rO  index_inner_fnr[  s    rg   r  z#index_impl_helper.<locals>.inner_fnb  s    xs++,,,ri   )rp   rn   r   ru  r  rs   r   r   r+  
IndexErrorr'  )r   r  r   r  r"  r  r#  r  r!  r  r3  r[  rf  s    `        @@@rg   r)  r)  C  s   ge}-----}}H9'1<<>>RRG^~"""$F"""SS7SSSO w~a01::<<==KZZ\\FUUUUUuS\\':':UUUL Ml""q';';KLLL;;;;uS\\':':;;;L"@
# 
# 
#K- - - - - - .00ri   c                    	 t          | |d          S # t          $ rA |                                   t          t          j        j        d          | |          cY S w xY w)NTr   Fr  )r*  r   r  r  r  r_  r  r   r  s     rg   r_  r_  h  sz    
!WD1111 
 
 
			M
 1uMMMw
 
 	
 	
 	

s    AAAc                &    t          | |d          S )NFr6  )r*  r7  s     rg   _unsafe_indexr9  t  s    a....ri   c                F    t          t          |           |||dd          S )NTFr   may_realizeindex_put_impl_r  r   r  r&  
accumulates       rg   	index_putrA    s,    a'6:Tu   ri   c                F    t          t          |           |||dd          S )NFr;  r=  r?  s       rg   _unsafe_index_putrC    s,    a'6:U   ri   c                   |                                 |                                  k    r"t          ||                                            }|rt          | |          }t          | t	          |d         ||                     S r  )rs   r  r   rY  r  )r/  r  r  r@  s       rg   index_put_as_masked_fillrE    ss    T__....%!2!233 !D%  T5UD99:::ri   c                j   ddl m} t          t          j        t
          j        j        j        j	                  }t
          j        j        }t          j        sD|B ||          r7d}|j                            dd           x}r| d| }|t
          j        _        t          j        || |||           | S )Nr/   ),_fx_node_is_input_dependent_cudagraph_unsafezLindex_put_ fallback with boolean indexing is not compatible with CUDA graphsstack_trace Found from : 
 )utilsrG  r   r  
index_put_rS   r_   r`   rc   _overloadnamer0   graph_partitionr  r  disable_cudagraphs_reasonr2   IndexPutFallback)	r/  r  r&  r@  rG  r  rh  r   rH  s	            rg   index_put_fallbackrP    s    CCCCCC$/17+?+F+TUUK
 g"G"088AA   ]!,**=$???; 	988;88C,/)T7FJGGGKri   c                ,    t          | |||dd          S )NTr;  r>  r/  r  r&  r@  s       rg   rK  rK    s&    gvz4   ri   c                ,    t          | |||dd          S )NFTr;  rR  rS  s       rg   _unsafe_index_put_rU    s&    gvzD   ri   c           
        |rZd t          j        |           |                                v r/t          fd|D                       s|                                 |                                dk    rt          |          dk    r|d                                         t          j	        t          j
        fv r||d         }t          t          |                                          t          |                                                     D ]}t          |d          }t          | |g||          S t          j                    rt!          | |||          S |D ]D}|@|                                t          j	        t          j
        fv rt!          | |||          c S E|                                 t                    }	|rbt#          |                                           rA|	dk    rt%          | dg          } t!          | |||          } |	dk    rt%          | g           } | S t'          ||                                           }	 t)          ||                                           \  }}
n"# t,          $ r t!          | |||          cY S w xY wd |D             }t/          | t0                    sJ |                                  |	dk    rt%          | dg          } t3          ||
d                                                            }fdt          t          |                    D             }t5          ||
|||d |          \  }}t7          ||          }|                                 }|J t          j        ||                                 |                                |||rd	nd 
          }t          j        d t          j        |           |          }t@          j!        "                    |          |_#        t@          j!        $                    |           |	dk    rt%          | g           } | S )Nc                   t          | t                    rt          | j        t          j                  r| j                                        } t          | t          j                  o`t          | j        t          j                  oAt          | j        dd           o+| j        j	        j
        t          j        j        j        j        u S dS )Nrh  F)rp   rC   rr   r2   r6   r  r   r  r   rh  rc   r   rR   r  randpermr  )indices    rg   indice_slice_from_randpermz3index_put_impl_.<locals>.indice_slice_from_randperm  s     &),, FK1U1U 0022vr}55 V"6;@@VY==V +2ein6M6UU	 5ri   c              3  .   K   | ]} |          V  d S r   r   )r   rY  rZ  s     rg   r   z"index_put_impl_.<locals>.<genexpr>  sH       H
 H
39&&v..H
 H
 H
 H
 H
 H
ri   r/   r   rr  c                >    g | ]}||                                 nd S r   rt  r  s     rg   r   z#index_put_impl_.<locals>.<listcomp>  r-  ri   c                     g | ]
}|         S r   r   r0  s     rg   r   z#index_put_impl_.<locals>.<listcomp>  r1  ri   r6  
atomic_addrz   r   r  r  output_indexerscatter_moder   r  rr   )%r2   try_get_namer  rC  r  r  r   r   r   r\   r  r+  r   rK  rE  $are_deterministic_algorithms_enabledrP  rL   r  r  r  rs   r   rp   rC   rn   r'  rS  Scatterru  r  MutationLayoutSHOULDREMOVErS   r_   r~  r   r  )r/  r  r&  r@  r   r<  r  r"  r_  x_ndimr"  r  r#  r  expected_vals_sizer  rz   scatterrb  rZ  rf  s                      @@rg   r>  r>    s{    	 	 	 ?4  F$9$9$;$;;;C H
 H
 H
 H
=DH
 H
 H
 E
 E
; NN 	aLLAAJ  ""uz5;&???qzs4==??++S-A-ABB 	' 	'AT2&&DD'tffjIII 133 E!$DDD  I I!2!2uz5;6O!O!O%dGVZHHHHH]]__F[[F B4>>CSCSTT Q;;qc??D!$DDQ;;b>>Dfdnn..//FE"=T__&&#
 #
  E E E!$DDDDDE TS7SSSOdI&&&&&LLNNN {{D1# w~a01::<<==K;;;;uS\\':':;;;L#A	$ 	$ 	$  F.//F __Fjnn##%%!%/9\\T  G ,T22  F
 '))&11FKGv&&&{{D"~~Ks   %I4 4JJr  c                  	 t          | |dd          \  }}                                |                                 		fd}t          j        |                                 |                                 ||          S )NFr  c                     j         t          j        k    r)t          j                    t          j                  }n            }t          j        | fd          S )Nc                 ,                           S r   r   )_unsafe_index_fnrO  self_loaders   rg   rj  z8_unsafe_masked_index.<locals>.inner_fn.<locals>.<lambda>I  s    KK8H8H8M8M,N,N ri   )r   r   r\   rR   r  r  )rO  mask_valrm  fillr  mask_loaderrn  s   ` rg   r  z&_unsafe_masked_index.<locals>.inner_fnD  sf    :##|KK$4$4ejAAHH"{3''Hz($N$N$N$N$N$NPTUUUri   r  )r)  ru  r@   r  rs   r   )
r/  r  r  rp  r  r"  r  rm  rq  rn  s
    ` `   @@@rg   _unsafe_masked_indexrr  <  s    "3gUU# # #FA ""$$K""$$KV V V V V V V V V   nn	   ri   c                    t          ||d          }|                                 fdt          t                              D             }t	          | ||d          S )Nr   c                r    g | ]3}|         r't          |         |          |         d z
            nd4S rV  rq  )r   rw   r  rL  s     rg   r   z7_unsafe_masked_index_put_accumulate.<locals>.<listcomp>W  sU        7>ajJgaj58)U1X\222d  ri   T)r@  )r  r   r+  r   rC  )r   r  r  r&  masked_valueclamped_indicesrL  s     `   @rg   #_unsafe_masked_index_put_accumulaterw  S  sv    vq))LJJLLE    s7||$$  O Q$OOOOri   c                R    t          j        |t          j        ||                     S r   )rR   r  r  r  minmaxs      rg   rr  rr  `  s     ;sCKQ//000ri   c                h    t          |           }t          ||||          }t          ||           |S r   )r  r  copy_)r/  rZ  r  r  rt  rV  output_views          rg   as_strided_scatterr  e  s5    4[[FVT6>BBK	+sMri   c                <    t          t          |           |||fi |S r   )scatter_r  )r   r   r_  rZ  r  s        rg   ri  ri  m  s$    E!HHc5#88888ri   r  include_selfr  r  r  c               b   t          |t                    }t          | ||                                t	          t
          j        |r|                                nt          |                    |r|                                j        nd|          rt          j
        | ||||||           |S d S )Nznot implr  )rp   rC   rQ   r   r
   r   r   r   rs   r2   ScatterFallback)r  r/  r   r_  rZ  r  r  src_is_tensors           rg   scatter_fallbackr  r  s     sI..MU[]I#--///S		JJ!.>J   	%	
 	
 	
 	
 4ri   r  c                   |dv sJ |Lt          t          j        t          j        j        j        j                  }t          || ||||          }||S |dk    rd}n|dk    rd}t          | ||||          S )N)Nr   multiplyr  r   r"  r  prod)
r   r  r  rS   r_   r`   rc   rL  r  scatter_reduce_)r/  r   r_  rZ  r  r  fallback_results          rg   r  r    s    .....~dmQW-A-H-VWW*sE3v
 
 
 &""	:		4eS&999ri   c                @    t          t          |           |||          S r   )scatter_add_r  r   r   r_  rZ  s       rg   scatter_addr    s    a#uc222ri   c                (    t          | |||d          S )Nr"  )r  r  s       rg   r  r    s    1c5#u555ri   c                >    t          t          |           ||||fi |S r   )r  r  )r   r   r_  rZ  reduction_typer  s         rg   scatter_reducer    s&    588S%nOOOOOri   )r  c          
         |dv sJ t          t          j                                                  dk    r dt          j                                        v s
J d            t	          t
                    rt                     t          t          j        j         |||          }|r|S t	           t                    sJ dt          |                                          v sJ t                                                     }|dk    rt           dg           t	          t                    r6t                                                    dk    rt          dg          t	          |t                    r6t          |                                          dk    rt          |dg          }|                                dk    r S t                                                       |                                t	          t                    r                                nd  fd} fd	}	d
 }
                                 }|J |st'          j        |                                  fd|                                |d           }t'          j        d t'          j                   |          }t.          j                            |          |_        t.          j                            |           t'          j        |                                 |	|                                | |
|                    }t'          j        d t'          j                   |          }t.          j                            |          |_        t.          j                            |           |dk    rt           g             S )N)Nr"  r  meanamaxaminr/   r  zKaten.scatter_reduce_.two is not the unique overload of aten.scatter_reduce_r  r   r   c                                                     }t          |          }t          |           }t          j         |           |dk    rdn|         d          |<   |S )Nr   r/   F)r  )r   r   rn   rR   r  )rO  rL  rt  indirect_idxr   r  r/  s       rg   r`  z'scatter_reduce_.<locals>.output_indexer  sk    5zzCyy1LDAIIqq5:
 
 
S ri   c                l    r |           S t          j                                                  S r   rR   r  r   )rO  r/  rZ  r  s    rg   r}   zscatter_reduce_.<locals>.fn  s7     	7:c??" <T^^%5%5666ri   c                    | dk    rdS | J d S )Nr"  r^  r   r  s    rg   backend_reduce_strz+scatter_reduce_.<locals>.backend_reduce_str  s     U??< >>>4ri   c                R    t          j        d                                          S r  r  )r_  r/  s    rg   rj  z!scatter_reduce_.<locals>.<lambda>  s    3<4>>3C3C#D#D ri   r_  rb  )r   r  r  r   rp   r!   r7  r  r  rC   r   r   r   r  r  rv  r  ru  rs   r2   re  r  rf  rS   r_   r~  r   r  )r/  r   r_  rZ  r  r  r  rt  r`  r}   r  rz   zero_outrb  ri  r  r  s   `` `           @@rg   r  r    s   BBBBBD **,,--22T)3355555T 65	6 #v #c""& !  O  dI&&&&&C))******t}}DqyyD1##y!! c#,,..&9&9Q&>&>3nn%## !ENN,<,<(=(=(B(BUQC  A
c
"
"CLLNNN$$&&L&0i&@&@J"""dJ      7 7 7 7 7 7 7   __F +:..""DDDD>>##)
 
 
 "066
 
 

 g--f55	""6***
 jnn~~%''//  G ,T22  F
 '))&11FKGv&&&qyyD"~~Kri   scales_xtuple[Optional[float], ...]exactc                &  
 |                                   |                                 |                                  d          
|                                 d           }d 
D             
t          |          k    sJ |}d t	          
|          D             t          |          D ]\  }}|d|z  |<   fd
fd}	t          j        |                                 | 	                                |	g ||          S )Nc                V    g | ]&}t           j        j                            |          'S r   r%  r  s     rg   r   z&upsample_nearestnd.<locals>.<listcomp>5  s+    >>>qw))!,,>>>ri   c                    g | ]
\  }}||z  S r   r   )r   rw   os      rg   r   z&upsample_nearestnd.<locals>.<listcomp>:  s     :::DAq!a%:::ri   r  c                z   t          j        | t          j                  } r2t          j        | t          j        dt          j                            } t          j        | t          j        |t          j                            } t          j        | t          j                  } t          j	        | |d          S )N      ?Fr6  )
rR   r  r   r  r   r  r  r  r  r  )r   r  r  r  s      rg   scale_fnz$upsample_nearestnd.<locals>.scale_fn?  s     N1em,, 	=3<U];;<<AGAs|E5=99::LEK(($QE::::ri   c           	         |  d          }| d           } g |fdt          |          D                       S )Nc                2    g | ]\  }}} |||          S r   r   )r   rw   r   r  r  s       rg   r   z2upsample_nearestnd.<locals>.fn.<locals>.<listcomp>N  s-    VVVJAq$88Aq$''VVVri   )r*  )rO  r   rU  i_sizes
inv_scalesr  r  r[  s      rg   r}   zupsample_nearestnd.<locals>.fnJ  sc    H!HxWaWVVVV#aW:U:UVVVW
 
 	
ri   r  )
realize_hintru  r   r   r*  ro   r@   r  rs   r   )r   r!  r  r  r  batcho_sizesrw   r  r}   r  r  r  r[  s      ``     @@@@rg   upsample_nearestndr  *  sZ    NN}}HjjllA233GJJLL1"E>>g>>>Gx==AG::C$9$9:::Jh'' ( (5%KJqM	; 	; 	; 	; 	;
 
 
 
 
 
 
 
 
 ||~~kkmm!!!	   ri   Optional[float]c                *    t          | ||fd          S )Nr/   r  r  r   r!  r  s      rg   upsample_nearest1dr  Y  s    avi1====ri   c                ,    t          | ||fdd          S )Nr/   Tr  r  r  r  s      rg   _upsample_nearest_exact1dr  ^  s    avi1DIIIIri   scales_hscales_wc                ,    t          | |||fd          S )Nr-   r  r  r   r!  r  r  s       rg   upsample_nearest2dr  c  s      ax.BaHHHHri   c                .    t          | |||fdd          S )Nr-   Tr  r  r  s       rg   _upsample_nearest_exact2dr  j  s"     ax.BatTTTTri   scales_dc                .    t          | ||||fd          S )Nr   r  r  r   r!  r  r  r  s        rg   upsample_nearest3dr  q  s#     ax8.LPQRRRRri   c                0    t          | ||||fdd          S )Nr   Tr  r  r  s        rg   _upsample_nearest_exact3dr  |  s-     	;8X6!4   ri   c                :     t           fd|D                       S )Nc              3  B   K   | ]}t          j        |          V  d S r   r  )r   r  r   s     rg   r   z$_create_constants.<locals>.<genexpr>  s/      66Aa''666666ri   )r   )r   rx   s   ` rg   r  r    s&    6666666666ri   c                    |                                  |                                 fd}t          j        |                                 |                                 |          S )Nc                    t          |           } t          |           t                    k    sJ D ]}|         dz
  | |         z
  | |<    |           S r  )rn   r   )rO  r   r+  r  r[  s     rg   r  zrev.<locals>.loader  sd    3ii3xx3u::%%%% 	3 	3Cc
Q#c(2CHHx}}ri   r  )ru  r   r@   r  rs   r   )r   r+  r  r  r[  s    ` @@rg   revr    s|     }}HJJLLE       ||~~kkmm	   ri   paddingSequence[int]r  c                l   d } |            rdS t          |          dk    s%t          |                                           dk    rdS |                                  t          | t          j                  rt          | j        t          j                  ret          | j        j        t          j                  s0t          j
        r5t          | j        j        t          j                  r| j        j        j        sdS |                                  t	          j        |           \  }}|j        }|d         dk    rdS |d         dk    s|d         dk    s|d         dk    rdS |d         }|dk    rdS |d         }|j        d         }	||	|z   k     rdS | j        j        j        }
|j        d         |j        d         |z   g}|t"          j        j        |
<   t)          | ||j        |j                  }t-          |d|	|	|z   d	          }t/          ||           t0          d
         dxx         dz  cc<   |S )z
    This optimization changes the semantics of padding from 'clone'
    style to 'view' style.

    Thanks to functionalization, this change can still maintain numerical
    correctness.
    c                     t           j        j        } | dS t          | j                  }t          |          dk    r1|d         j        t          j        j	        t          j
        j	        fv rdS dS )a  
        Conservatively check if padding can be fused with downstream op.
        1. if the downstream op is a sum, then there is little benefit to
           do inplace padding
        2. if the downstream op is a matmul, doing inplace padding can
           save membw.
        NTr/   r   F)rS   r_   r`   r   ra   r   rc   r  mmr  addmm)r`   ra   s     rg   _padding_can_be_fusedz6inplace_constant_pad_nd.<locals>._padding_can_be_fused  sj     w+4l())u::??uQxGOJ3
  
  
 5tri   Nr   r-   r/   r   r   F)r   r  r  rr  inductorinplace_padding)r   r   r  rp   r2   rC   rr   r   r  r0   can_inplace_pad_graph_inputInputBufferr   freeze_layoutr  r  r  rS   r_   buffer_to_padded_sizer  r  r  fill_r   )r   r  r  r  r"  r  r@  npadstride0rowsizebufnamepadded_size	resized_xsliced_xs                 rg   inplace_constant_pad_ndr    sB     (  t 7||qC

--22t
 IIKKK
 q",''!&"-00 qv{B$566	 2 qv{BN;; v{ tOO(++IAvmGqzQtqzQ'!*//WQZ1__t1:DqyytajGk!nG4tfkG;q>6;q>D#89K-8AG!'*		 I iQg7T>QVWWWH	(JZ*+++q0+++ri   c                   t          |          dz  dk    sJ t          d |D                       rt          |           S t          j        rt          | |          }|r|S |                                 }t          t          t          t          |d d d         |dd d                                                 t          |          t                    z
  g D ]>\  }}
                    t          j        j                            |          |f           ?t          |d                    }g t          |d                    D ]J\  \  }}	}

                    |
           |
                    t          j        |
|z   |	z                        Kt          |          t          |          k    sJ  t#          |                                                     fdfd}|                                 t)          j        |                                 |                                 ||          S )Nr-   r   c              3  "   K   | ]
}|d k    V  dS r   Nr   r   r  s     rg   r   z"constant_pad_nd.<locals>.<genexpr>  s&      
#
#a16
#
#
#
#
#
#ri   r/   c                f    g }t           	d                    D ][\  }\  }}}|dk    r#|                    t          |d                     |dk    r#|                    t          ||                     \t	          j        t          j        |          }t          j        | 
fd          S )Nr   c                                 S r   r   )r_  r[  s   rg   rj  z/constant_pad_nd.<locals>.mask.<locals>.<lambda>(  s     ri   )	r*  rt   range_mask_lowrange_mask_highr6  r  rR   r  r  )r_  r  rO  r  r  r  boundsr  
mask_sizesr  r[  s   `     rg   r  zconstant_pad_nd.<locals>.mask   s    (+E!""Ivz(J(J 	: 	:$C#tfaxxN322333qyyOC88999$//z$ 7 7 7 7 7DDDri   c                    t          | d                    }t          | d                    D ] \  }\  }}|                    ||z
             !t          |          t          |           k    sJ  |          S r   )rn   r*  rt   r   )r_  r  rO  r  _highbounds_precompr  r  s        rg   	offset_fnz"constant_pad_nd.<locals>.offset_fn*  s    rrOO	!$U122Y!?!? 	( 	(C#uS3Y''''9~~U++++tIri   r  )r   rC  r  r0   r  r  r   rn   rN  r*  rt   rS   r_   r   lookup_precomputed_sizer   rS  r   r   ru  r@   r  rs   )r   r  r  ru   r  lhr!  r  r  r  r  r  r  r  r  r  r[  s     `         @@@@@@rg   r+  r+    s   LL1""""

#
#7
#
#
### Qxx %a*== 	J JJLLE(4GCCaCL'!$Q$- @ @AABBCCFE

S[[ A 68N P P1qw/GGJJANOOOOuRaRy//KJ qrr33 < <dT$5<s
T(9::;;;;{s5zz))))-q{{}}--j99JE E E E E E E E E       }}H||~~kkmm	   ri   rw   
sympy.ExprUnion[sympy.Expr, int]c                    t          j        t          j        | t          j                  t          j        t          j        |          t          j                            S r   )rR   r  r  r   r  r   r  )rw   r  s     rg   r  r  :  sB    6q%+&&u}S))5;77  ri   c                    t          j        t          j        | t          j                  t          j        |t          j                            S r   )rR   r  r  r   r  )rw   r  s     rg   r  r  A  s8    6q%+&&tU[))  ri   c                d    t          j        t          | |          t          | |                    S r   )rR   r  r  r  )rw   r  r  s      rg   
range_maskr  H  s0    8q#4    ri   r  c                                                        d                                           pdgz   fd}|S )Nr   c                   | d           |  d          t          j        t          j        	fdt	                    D                       }rt          j        |
fd          nt          j        |fd          S )Nc                j    g | ]/}t          |         |         |         z   |                    0S r   )r  )r   rw   r  ih	padding_hs     rg   r   z=constant_boundary_condition.<locals>.load.<locals>.<listcomp>^  s<    WWWqZ1qtil2Yq\MBBWWWri   c                 B     t                     g           S )Nr   )constant_boundary_condition)r   r  pad_fill_valueprefixr   s   rg   rj  z;constant_boundary_condition.<locals>.load.<locals>.<lambda>c  s/    O3A~3OOO"fNrN  ri   c                 "     g            S r   r   )r  r  r[  s   rg   rj  z;constant_boundary_condition.<locals>.load.<locals>.<lambda>i  s    ((>V>b>*B*B ri   )r6  r  rR   r  r+  r  )r_  r  r  r  r   r  r  r  r  r   r   r[  s     @@rg   rz  z)constant_boundary_condition.<locals>.loadW  s    uuC455\HWWWWWWERUJJWWW
 
 PCJ           D"B"B"B"B"B"BJOO
	
ri   )r   ru  )	r   r  r  r  r   rz  r  r   r[  s	   ````` @@@rg   r  r  O  s     	


cTUUA}}H$A39I
 
 
 
 
 
 
 
 
 
 
 
* Kri   dilationc                  |dgt          |          z  }t          | d||         z  z   ||         ||         dz
  z  z
  ||         dz
  z   ||                   }|rt          | d||         z  z   ||         ||         dz
  z  z
  d||         dz
  z  z   ||                   }t          j        j                            |dz
  ||         z  | z
  ||         z
            dk    r?|dz  }t          j        j                            d|||         z  | z
  ||         z
             t          j        j                            ||z
            dk    r(t          j        j                            ||           d}n|}||fS )Nr/   r-   r   F)r   r)   rS   r_   r   	size_hintrq  rR  )	r   rw   kernel_sizer  r  	ceil_moder  x_outx_alts	            rg   pooling_sizer  o  s   3W%	A
NXa[KNQ,>??6!9q=Qq	 E
  '!*nqk[^a/01 6!9q=!" 1I
 
 7%%uqyF1I&=&AGAJ&NOOSTTTQJEG&&q%&)*;a*?'!**LMMM7%%eem4499G))%777IIE)ri   c               l    t          | |          } t          j        t          j        |           }|dk    S )N   )rM   r6  r  r?  r  )r  n_dimwindow_sizes      rg   %should_fallback_max_pool_with_indicesr    s0    {E22K"8<==Kri   assert_fallbackc               .   |dk    rdg|z  }|dk    rdg|z  }|s|}t          ||          }t          ||          }t          ||          }t          ||          }t          | t                    sJ t          |          |k    sJ t          |          |k    sJ t          |          |k    sJ t          |          |k    sJ t          |                                           |dz   |dz   fv sJ t          ||          }|||k    sJ |||||fS )Nr   r/   r-   r  )rM   rp   rC   r   r   r  )r   r  r  r  r  r  r  use_fallbacks           rg   max_pool_checksr    sQ    !||#+1}}3; {E22K&%((F7E**GHe,,Ha#####{u$$$$v;;%w<<5    x==E!!!!qzz||EAI 666668ERRRL"....<??ri   c          
        |                                   | j        d           }| j         d          t          fdt                    D              \  }| j        }	|	t
          j        u rdn.|	j        rt          d          nt          j	        |	          j
        }
t          |          t          |          z   }t                    s(t                    st          d D                       rt          | |
          n|                                 fd}t          j        d| |                                 |	|	||          }t          j        d	| |                                 t
          j        |	||          }t'          |j        j        t                    r|                                 t'          |j        j        t                    r|                                 ||fS )
Nc                H    g | ]}t          |         |           S )r  r  )r   r(  r  dhwr  r  r  r  s     rg   r   z*_max_pool_with_offsets.<locals>.<listcomp>  sM     

 

 

  A;X  

 

 

ri   F-infc              3  "   K   | ]
}|d k    V  dS rV  r   r'  s     rg   r   z)_max_pool_with_offsets.<locals>.<genexpr>  s&      ,E,EqQU,E,E,E,E,E,Eri   r  c                    | d           }|  d          fdt                    D             } 	g ||          S )Nc                h    g | ].}|         |         z  |         |         z  z   |         z
  /S r   r   r   rw   bhr  r  reduction_idxr  s     rg   r   z<_max_pool_with_offsets.<locals>.fn_inner.<locals>.<listcomp>  sR     
 
 
 UVAY=#3hqk#ABWQZO
 
 
ri   r+  )
rO  r%  r  r  r$  r  r  r  r  r[  s
    `  @rg   fn_innerz(_max_pool_with_offsets.<locals>.fn_inner  s    WufW%\
 
 
 
 
 
 
 
5\\
 
 
 x&2'''ri   r{  r  
input_noderz   	dst_dtyper  r  r  reduction_rangesargmax)r  rL  r*  r+  r   r   r\   r  r[  r  rz  rn   rD  r  ru  rA   r  rs   r  rp   rr   r  )r   r  r  r  r  r  r  r  dhw_outr   	min_valuer]  r'  r  offsetsr  r[  s    ``````        @@rg   _max_pool_with_offsetsr0    sK    NNGGeVGE
'5&''
C

 

 

 

 

 

 

 

 

 5\\	

 

 

GY GE EJ 	$6ReFmmmEK<N<N<R  E{{T']]*H
7|| #s9~~ #,E,EH,E,E,E)E)E #.q)GGG==??( ( ( ( ( ( ( ( ( ||~~$	 	 	F ||~~+$	 	 	G &+"I.. ',#Y// 7?ri   c           
         t          |          }t          | |||||d          \  }}}}}t          j        d          5  t	          | ||||||          \  }}	|t          |	t          j                  fcd d d            S # 1 swxY w Y   d S )NFr  r  unroll_reductions_thresholdr  )r   r  r0   r   r0  r  r   r
  )
r   r  r  r  r  r  r  r"  r  r/  s
             rg   !_low_memory_max_pool_with_offsetsr4    s    E 1@	1 1 1-K(A 
"	5	5	5 
5 
50
 
 
 x444
5 
5 
5 
5 
5 
5 
5 
5 
5 
5 
5 
5 
5 
5 
5 
5 
5 
5s   5BB
Br/  r  "Sequence[Union[int, torch.SymInt]]
input_sizeincrements_to_indexxCallable[[Sequence[Union[int, torch.SymInt]], Sequence[Union[int, torch.SymInt]]], torch._inductor.virtualized.OpsValue]c                Z   t                    |                                 t          j        t	          j        t          j                            fd}t          j	        | 
                                t          j        ||                                           }|S )Nc                     	|           }t          j        |
          }t          j        |          } | |          }t          j        t          j        | d                    t          j                  S r   )rR   r  r1   _flattened_index_to_ndr  _flatten_indexr   r  )rO  r  offset_sympyr%  idhwr7  r6  r  r  offsets_loaderr  s        rg   offsets_to_indicesz4_pool_offsets_to_indices.<locals>.offsets_to_indices+  s{    $$,V[AA&=lKXX""366~)$
E6770CDDek
 
 	
ri   r  )r   ru  r   r~  r6  r  r?  r  r@   r  rs   r   r  r   )	r/  r  r6  r7  r@  r  r  r?  r  s	    ```  @@@rg   _pool_offsets_to_indicesrA    s     E((**N-	 0{ K KLLK
 
 
 
 
 
 
 
 
 
 !!##k#!!	  G Nri   c                \    t          |          fd}t          | |||          S )Nc                \    |  d          fdt                    D             S )Nc                h    g | ].}|         |         z  |         |         z  z   |         z
  /S r   r   r#  s     rg   r   zX_low_memory_max_pool_offsets_to_indices.<locals>.increments_to_index.<locals>.<listcomp>H  sR     
 
 
 UVAY=#3hqk#ABWQZO
 
 
ri   r&  )rO  r%  r$  r  r  r  r  s    `@rg   r7  zD_low_memory_max_pool_offsets_to_indices.<locals>.increments_to_indexF  sY    %\
 
 
 
 
 
 
 
5\\
 
 
 	
ri   )r   rA  )r/  r  r6  r  r  r  r7  r  s      ``` @rg   '_low_memory_max_pool_offsets_to_indicesrE  =  s[     E
 
 
 
 
 
 
 
 $j*=  ri   c           	         t          | |||||          \  }}}}}t          | ||||||          \  }}	t          |	|| j        | d          |||          }
||
fS )Nr  )r  r0  rE  rL  )r   r  r  r  r  r  r  r"  ru   r/  r  s              rg   _max_pool_with_indicesrG  R  s     1@	;1 1 1-K(A *	;9E  LC 6	 G <ri   c           	     .    t          | |||||d          S Nr-   r  rG  r   r  r  r  r  r  s         rg   max_pool2d_with_indicesrL  p  *     "	;9A   ri   c           	     .    t          | |||||d          S Nr   r  rJ  rK  s         rg   max_pool3d_with_indicesrP    rM  ri   c                    dk    rddg|dk    rddg}st          |t                    sJ t                    dk    sJ t                    dk    sJ t                    dk    sJ t          |          dk    sJ t          |                                          dv sJ |                                  |                                 }t          |t                    rt          |j        j        t                    r|j        j        }	|	                                }
|
J t          j
        d t          j        |
|	                                |	                                          |	          }|                                 |                                }n|                                }|d ur|d         dk    p|d uo|d         dk    }t          d |D                       rt!          | ||||          S |                                ^ }}|                                 ^ }|                                |                                 t%          |                                          }t'          fdt)          d         dz            D                       t'          fd	t)          d         dz            D                       z  }|d
k    rt!          | ||||          S |                                fd}t          j        |                                 |                                 ||          }|rt          j                            |          S |S )Nr   r/   r-   ro  )rz   r   r  rb  c              3  "   K   | ]
}|d k    V  dS rV  r   r'  s     rg   r   z3max_pool2d_with_indices_backward.<locals>.<genexpr>  s&      
$
$a16
$
$
$
$
$
$ri   c              3     K   | ]V}t          t          |d                    t          d t          |d          z
  d                              z
  d          V  WdS r   r/   Nr{  r)   r   r  r  r  s     rg   r   z3max_pool2d_with_indices_backward.<locals>.<genexpr>  w         	HQq	""SHQQ5GPQ,S,S%T%TTVWXX     ri   c              3     K   | ]V}t          t          |d                    t          dt          |d          z
  d                              z
  d           V  WdS r/   r   NrU  r   wr  r  s     rg   r   z3max_pool2d_with_indices_backward.<locals>.<genexpr>  rW  ri   r  c                d   | ^ }}}t          j        |z  |z   t          j                  }|d         z   }|d         z   }t          j        t	          |d         z
  d         z   d                   t          j                  }t          j        t	          |d         z
  d         z   d                   t          j                  }t          j        t	          |d                   dz   t          j                  }t          j        t	          |d                   dz   t          j                  }t          j        |t          j        dt          j                            }t          j        |t          j        dt          j                            }t          j        |t          j        t          j                            }t          j        |t          j        t          j                            }d }	t                    D ]}
t                    D ]}t          j	        |t          j        |
t          j                            }t          j	        |t          j        |t          j                            }g |t          j
        t          j        |t          j        |t          j        dt          j                                      d         d          t          j
        t          j        |t          j        |t          j        dt          j                                      d         d          } |          } |          }t          j        ||          }|	5t          j        ||t          j        dt          j                            }	t          j        t          j        t          j        ||          t          j        ||                    |          }t          j        |t          j	        |	|          |	          }	|	J |	S )Nr   r/   Fr6  rr  r  )rR   r  r   r  r)   r  r  r  r+  r   r  r  r  r  r  r  r  )rO  r  r  r[  
index_testphstartpwstartphendpwendgradientph_pw_phpw
grad_indexindex_actual	grad_partr   r  grad_loaderh_window_sizer   indices_sizer  r  pooled_heightpooled_widthr  w_window_sizewidths                      rg   r}   z,max_pool2d_with_indices_backward.<locals>.fn  s   A^AIM5;??

N
N.QQ'&)3VAY??
 
 .QQ'&)3VAY??
 
 x6!95595;GGx6!95595;GG+gs|Au{'C'CDD+gs|Au{'C'CDDE3>-#M#MNNE3>,#L#LMM'' #	W #	WC]++ "W "WWWcl3&D&DEEWWcl3&D&DEE)Bs|Au{7S7S(T(TUU$R(#   )Bs|Au{7S7S(T(TUU$R(#  
  .~j99'K
33	|Z88#"yy#,sEM*J*J   HH 8F2u--F2u--   D  #yswx/K/KXVVHHE"WF ###ri   r  )rp   rC   r   r   r  ry  rr   r@   rs   r2   r  r  r   decide_layoutr  rD  )fallback_max_pool2d_with_indices_backwardru  rn   r{  r+  r  r  r8  )grad_outputr   r  r  r  r  r  r  	gO_striderr   rz   x_bufferx_strideis_channels_last_batch_heightr"  r]  r  r}   ru   rk  rl  r   rm  rn  ro  rp  rq  s     ```                @@@@@@@@rg    max_pool2d_with_indices_backwardr{    s>    !||a&1}}q6 a#####{q    v;;!w<<1x==Aqzz||&&&& ,,..I!Y (Jqv{I$F$F (v{""!!!$$nn&&]]__  
 
 
 
 	   &&((%%'' ,A!1A 3)A,!"3  
$
$8
$
$
$$$ 
8K(Iw
 
 	
  jjllVWe&1&:&:&<&<#Q|((**N))++KAJJLL!!H     {1~)**    M      {1~)**    M
  -/KR8K(Iw
 
 	
 ##%%L9 9 9 9 9 9 9 9 9 9 9 9 9 9 9v 
%%''##%%	  C  44S999
ri   r  c                >    |                                  fd}|S )Nc           
        	
 |\  	|\  
|\  }}t          j        t          j        t          j        z   t          j                  t          j        |t          j                            t          j        t          j        
	z   t          j                  t          j        |t          j                                      }t          j        |	 
fd          S )Nc                 2     g  z   z             S r   r   )h_start_indexr  iwr  w_start_indexr[  s   rg   rj  z3pad_adaptive_loader.<locals>.load.<locals>.<lambda>:  s+    HHNvN}r'9N=2;MNOO ri   )rR   r  r  r  r   r  r  )r  
incrementsstart_indicesend_indicesh_end_indexw_end_indexr  r  r  r  r  pad_valr[  s   `      @@@@rg   rz  z!pad_adaptive_loader.<locals>.load(  s    B'4$}#. [xF}r15;??{EK88  F}r15;??{EK88 	
 	
 zOOOOOOOOO
 
 	
ri   rt  )r   r  rz  r[  s    ` @rg   pad_adaptive_loaderr  %  s6    }}H
 
 
 
 
 
, Kri   c                    t          j        | ||          }t          j        |||          }t          j        | ||          }t          j        |||          }	||||	fS )N)out_diminp_dim)r6  rI  )
r  r  h_inw_inh_outw_outr  r  r  r  s
             rg    compute_indices_adaptive_poolingr  A  sm    %k5$OOOM#IudKKKK%k5$OOOM#IudKKKK+}kAAri   c                l    |\  }}|\  }}	t          | |||||	          \  fd}
|
S )Nc                <   | ^ }}} |          } |          } |          } |          }d }	t          j        t          d                   t          d                             D ]*\  }
} |||
|g||g||g          }|	|}	 ||	          }	+|	S r
  )r(  productr+  )rO  r  r  r$  bwr  r  r  r  r  r  r  r2  h_end_index_fnh_start_index_fnkernel_maxes
pooling_fnw_end_index_fnw_start_index_fns                rg   r}   z _adaptive_pooling_fn.<locals>.fnZ  s    R((,,$nR((((,,$nR(('l1o(>(>lSTo@V@VWW 
	1 
	1FB&R.k*	 C ~#C00ri   r  )r  r  r  in_sizes	out_sizesr  r  r  r  r  r}   r  r  r  r  s     `  `     @@@@rg   _adaptive_pooling_fnr  K  s     JD$LE5 	)YdE5	 	
         . Iri   c                p   
 |\  }|\  }}t          | ||||          \  

fd}	|	S )Nc                   | ^ }}} |          } |          } |          } |          }d }	d }
t          j        t          d                   t          d                             D ]\  }} ||||g||g||g          }t          j        ||z   z  |z   |z   t
          j                  }|
|}
n)t          j        t          j        ||	          ||
          }
|	|}	w ||	          }	|
S r
  )	r(  r  r+  rR   r  r   r  r  gt)rO  r  r  r$  r  r  r  r  r  maxvalmaxindexr  r  r2  r_  r  r  r  r  r  r  r  s                  rg   r}   z)_adaptive_pooling_fn_with_idx.<locals>.fn  s:   R((,,$nR((((,,$nR(('l1o(>(>lSTo@V@VWW 	1 	1FB&R.k*	 C N#t+m;b@%+ E  9SVC%8%8%JJ~#C00ri   r  )r  r  r  r  r  r  r  r  r  r}   r  r  r  r  r  s     `  `    @@@@@rg   _adaptive_pooling_fn_with_idxr  t  s     JD$LE5 	)YdE5	 	
! ! ! ! ! ! ! ! ! ! !F Iri   c                                                      t          j        k    rt          d          t	           t
                    sJ t          |          dk    sJ                                                                    ^ }}}t          j
        j                            |          }t          j
        j                            |          }|\  }}||k    r||k    rt                     S |dk    s|dk    r>g |||}t          |                                                                             S ||z  dk    r9||z  dk    r0t!          ||          t!          ||          g}t#           |          S t%          ||z   dz
  |          }	t%          ||z   dz
  |          }
t'          |          ||gz   }                                  }|	|
z  }|dk    rt)           |          S d }d }t+          |||	|
g||g||gt,          j        	          t1          t3                                fd
}t5          j                                         |||          }|S )Nz0'adaptive_avg_pool2d' not implemented for 'Long'r-   r   rT  r/   r  c                (    t          | |z  |          S r   r)   r_  r  r  s      rg   r  z)_adaptive_avg_pool2d.<locals>.start_index      7333ri   c                :    t          | dz   |z  |z   dz
  |          S r  r  r  s      rg   r  z'_adaptive_avg_pool2d.<locals>.end_index  %    g-7!;WEEEri   r  r  r  r  r  r  c                p    t          j         | t                               |                     S r   )rR   truedivr  )rO  fn_sumones_loaderr   s    rg   r}   z _adaptive_avg_pool2d.<locals>.fn  s=    {F3+A..//[1I1I
 
 	
ri   r  )r   r   r  r  rp   rC   r   r  r   rS   r_   r   r&  r  r\  rs   r)   
avg_pool2drF   rn   fallback_adaptive_avg_pool2dr  rR   r   r  	ones_liker@   r  )r   r!  r  r  r  r  r  o_sizer  h_kernel_maxw_kernel_maxr]  r   r  r  r  r}   rvr  r  s   `                 @@rg   _adaptive_avg_pool2dr    s   {{}}##MNNNa#####{q    NNUD$7%%d++D7%%d++DLE5 u}}QxxzzUaZZ'5'%''V1;;==HHHHe|qTE\Q..e,,htU.C.CD![)))D5L1,u55LD5L1,u55LE{{eU^+HKKMME-KR+A{;;;4 4 4F F F ""L1%.7  F &ill33K
 
 
 
 
 
 

 
	||~~	
 
 
B Iri   c                0                                      t          j        k    rt          d          t	           t
                    sJ t          |          dk    sJ                                                                    ^ }}}t          j
        j                            |          }t          j
        j                            |          }|\  }}|dk    s|dk    rlg |||}t          |                                                                             t          |t          j                                                   fS ||z  dk    r||z  dk    rt          t!          ||z   dz
  |          }t!          ||z   dz
  |          }	t#          |          ||gz   }
                                  }||	z  }|dk    rt%           |          S d }d }t'          ||||	g||g||gt(          j        	          t-          ||||	g||g||gt(          j        	           fd
} fd}t/          j                                         |||
          }t/          j                                         t          j        ||
          }||fS )Nz,adaptive_max_pool2d not implemented for Longr-   r   rT  r/   r  c                (    t          | |z  |          S r   r  r  s      rg   r  z(adaptive_max_pool2d.<locals>.start_index  r  ri   c                :    t          | dz   |z  |z   dz
  |          S r  r  r  s      rg   r  z&adaptive_max_pool2d.<locals>.end_index!  r  ri   r  c           	     R     | t          t          d                              S Nr  r  r[  )rO  inner_func_max_valr   s    rg   inner_fn_max_valz-adaptive_max_pool2d.<locals>.inner_fn_max_val6  '    !!#':1eFmm'L'LMMMri   c           	     R     | t          t          d                              S r  r  )rO  inner_func_max_idxr   s    rg   inner_fn_max_idxz-adaptive_max_pool2d.<locals>.inner_fn_max_idx9  r  ri   r  )r   r   r  r  rp   rC   r   r  r   rS   r_   r   r&  r\  rs   
ValueErrorrF   rn   fallback_adaptive_max_pool2dr  rR   r  r  r@   r  )r   r!  r  r  r  r  r  r  r  r  r]  r   r  r  r  r  r  r  rir  r  s   `                  @@rg   adaptive_max_pool2dr    s   {{}}##IJJJa#####{q    NNUD$7%%d++D7%%d++DLE5zzUaZZ'5'%''V1;;==HHH%%+allnnK
 K
 K
 
 	
 e|qTE\Q..D5L1,u55LD5L1,u55LE{{eU^+HKKMME-KR+A{;;;4 4 4F F F ."L1%.;   7"L1%.;  N N N N N NN N N N N N 
	||~~!	
 
 
B 
	||~~k!	
 
 
B r6Mri   c                                                                                   fd}|S )Nc                r                                    }t          |          dk    rwt          |           dk    r d| d         dz
  
z
  g          }n^t          |           dk    r! | d         | d         dz
  
z
  g          }n* dddz
  
z
  g          }n g | dz
  
z
            }t          j        |                                          }t          j        z
  t
          j                  }t          j        dz
  t
          j                  }t          j        t          j        |t
          j	                  t          j        |t
          j	                            }t          j
        t          j        |d          d|          }t          j        ||z   |z            t          j        ||z            z
  }t          j        |t
          j                  }t          j        ||          }	t          j        t          j
        |	||          t          j                            S )Nr   r/   r   r-   )r   r   rR   r  r   r   r  r  r  float64r  r  rE  r  r  r   r~  )r  rw   samples_shapesamplei_exprdiffout_sz_exprrm  seq_ir  r   in_sz	kernel_szndimsout_szsamplessamples_loaders             rg   rz  z)_fractional_pooling_offsets.<locals>.loadQ  s     ((**}""6{{a (F1Iuqy3(GHHV!! (F1Iuqy3(OPP (Auqy3(?@@ $^$>f$>eai#o$>??F7#4#4#6#677~ei/==nVaZ==Lu}--s|K/W/W
 
 	#&a00!U;;	6F?e344sy%7P7PPUEK00vfk**$SYtUD%A%A5=QVCWCWXXXri   rt  )r  r  r  r  r   r  rz  r  s   `````` @rg   _fractional_pooling_offsetsr  K  s    C[F#JE#I((**N!Y !Y !Y !Y !Y !Y !Y !Y !Y !Y !YF Kri   c                *    t          | |||d          S rI  _fractional_max_poolr   r  r!  random_sampless       rg   fractional_max_pool2dr  w      ;^STUUUUri   c                *    t          | |||d          S rO  r  r  s       rg   fractional_max_pool3dr  |  r  ri   c                
   |                                   | j        d           | j         d          c}t          j        d          5  fdt	                    D             |                                 fd}fdt          |          t                    z   }|                                 }t          j	        d| | 
                                ||||          }	t          j	        d| | 
                                t          j        |||          }
t          |	t                    s
J |	            t          |	j        j        t                    r|	                                 t          |
t                    s
J |
            t          |
j        j        t                    r|
                                 t#          |
| j                  }|	|fcd d d            S # 1 swxY w Y   d S )	Nr  r2  c           
     :    g | ]}t          |           S ))r  r  r  r  r  r   )r  )r   r(  inp_dhwr  r  r!  r  s     rg   r   z(_fractional_max_pool.<locals>.<listcomp>  sL     

 

 

  (&"%  

 

 

ri   c                L    | d           } g | | |                    S r   r   )rO  r%  r  r7  r  r[  s      rg   r'  z&_fractional_max_pool.<locals>.fn_inner  s;    5&\F8OfO':':3'N'NOPPPri   c                r    | d           |  d          fdt                    D             S )Nc                R    g | ]#} |         |                   |         z   $S r   r   )r   r(  bdhwdhw_index_fnr  r%  s     rg   r   zE_fractional_max_pool.<locals>.increments_to_index.<locals>.<listcomp>  sG         QQ00=3CC  ri   r&  )rO  r%  r  r  r  r  s    `@@rg   r7  z1_fractional_max_pool.<locals>.increments_to_index  sd    5&\Fvww<D      u   ri   r{  r(  r,  )r  rL  r0   r   r+  ru  rn   r   rA   r  rs   r   r  rp   rC   rr   r  rA  )r   r  r!  r  r  r  r'  r]  r   r  r/  r  r  r7  r  r[  s    ````       @@@@rg   r  r    s   NNWWufW%qwvww'7NE7	"	5	5	5 = =

 

 

 

 

 

 

 

 5\\

 

 

 ==??	Q 	Q 	Q 	Q 	Q 	Q 	Q	 	 	 	 	 	 ;;k!2!22! <<>>(	
 	
 	
 "#<<>>k(	
 	
 	
 &),,44f44,fk&	22 	NN'9--66w66-gl'33 	OO*[!'+>
 
 w{= = = = = = = = = = = = = = = = = =s   FG88G<?G<c                                                                                        ^ }}}t          j        j                            |          }t          j        j                            |          }|^ }}}	||z  dk    r9||	z  dk    r0t           t          ||          t          ||	          gd          S t          ||          }
t          ||	          }d fd}t          ||
|g||g||	gt          j                   fd}t          j                                                                          |t!          |                    }|S )	Nr   r/   )divisor_overridec                L    t          | |z  t          j        |                    S r   )r(   r   r~  r  s      rg   r  z0upsample_nearest2d_backward.<locals>.start_index  s     uwg(>(>???ri   c                $     | dz   ||          S r  r   )r_  r  r  r  s      rg   r  z.upsample_nearest2d_backward.<locals>.end_index  s    {EAI999ri   r  c                6     | t                              S r   )r  )rO  r  r   s    rg   r}   z'upsample_nearest2d_backward.<locals>.fn  s    vc.q11222ri   r  )r  r   rS   r_   r   r&  r  r)   rF   r  rR   r   r@   r  rs   r   rn   )r   r!  r6  r  r  ry  inp_hinp_wout_hout_wr  r  r  r}   r  r  r  s   `              @@rg   upsample_nearest2d_backwardr    s    NNJJLLVUEG&&u--EG&&u--E 'VUEu}eemq00&&(>(>?RS
 
 
 	
 5%((L5%((L@ @ @: : : : : ""L1%.7  F3 3 3 3 3 3 
	||~~kkmmJ
 
 
B Iri   r   c           
     0    t          | ||||||d          S )Nr-   r  _avg_poolndr   r  r  r  r  count_include_padr  s          rg   r  r    3     		 	 	 	ri   c           
     0    t          | ||||||d          S )Nr   r  r  r  s          rg   
avg_pool3dr    r  ri   c                   ssdgz  t                    t                    t                    t          | t                    sJ t                    k    sJ t                    k    sJ t                    k    sJ t          |                                           dz   dz   fv sJ |                                  |                                 d           }|                                  d          t          fdt                    D              \  }	}
t                    st          |
          rt          | d          d}n| 
                                d}t          |          t          |	          z   }|                                 }t          | t          j        d	          }fd
}t!          j        t$          j                  }|dk    rHt          d t                    D                       r!t(          dz
           } || |          S |dk    rt+          j        d          nt/          j                    }|                                 }|J |5  t5          j        d| |||||          }d d d            n# 1 swxY w Y   t9          |j        d          r3t          |j        j        t4                    r|                                 |r|r|r|n|}t?          ||          }nDfd}tA          j        |                                 |||          }t?          ||          }tC          ||          S )Nr   r/   r-   c           
     D    g | ]}t          |         |          S r   r  )r   rw   r  r  r  r  r  s     rg   r   z_avg_poolnd.<locals>.<listcomp>G  s?     

 

 

 1q+vw	JJ

 

 

ri   r  r  TF)r   r   c                    | d           }|  d          |fdt                    D              g |          S )Nc                V    g | ]%}|         |         z  |         z   |         z
  &S r   r   )r   rw   r$  r  r  r  s     rg   r   z1_avg_poolnd.<locals>.fn_inner.<locals>.<listcomp>a  s8    IIIbefQi"Q%''!*4IIIri   r&  )	rO  r%  r  r$  r  r   r  r  r[  s	      @@rg   r'  z_avg_poolnd.<locals>.fn_inner]  sh    UsdU#ZIIIIIIIeCjjIIIx&2'''ri   r  c              3     K   | ]>\  }}t           j        j                            t	          j        ||                    V  ?d S r   )rS   r_   r   statically_known_truer   Ne)r   r  r   s      rg   r   z_avg_poolnd.<locals>.<genexpr>f  sW          Aq 	
..ux1~~>>           ri   r2  r"  r(  rr   c                   |  d          }g }t                    D ]}||         |         z  |         z
  }t          j        |
|         z   	|         |         z             }s0t          j        |d          }t          j        |	|                   }t	          j        ||z
  t          j                  }|                    |           t          j
        t          j        |          S r  )r+  r   MinMaxrR   r  r   r  rt   r6  r  r  )rO  r$  divide_factorsrw   hstarthendfactorr  r   r  r  r  r  s          rg   fn_countz_avg_poolnd.<locals>.fn_count  s    cTUUBN3ZZ . .A*WQZ7y+a.!8!A$:KLL( 1"Yvq11F 9T1Q400Dvu{CC%%f----#CG^<<<ri   r  )"rM   rp   rC   r   r   r  r*  r+  rD  r  ru  rn   r   r   r   rd  r6  r  r?  r  fallbacks_avg_poolndr0   r   
contextlibnullcontextrs   rA   r  r  rr   r  div_primr@   r  )r   r  r  r  r  r  r  r   r  r  
ceil_modeshad_paddingr]  r   output_dtyper'  r  fallbackcontextrz   r  divisorr  r
  divide_factorr  r[  s    ````` `                 @@rg   r  r  *  sx      #){C00K&#&&F7C((Ga#####{s""""v;;#w<<3qzz||q#' 22222NNJJLL3$E	

cTUUA

 

 

 

 

 

 

 

3ZZ

 

 

E: 7|| s: .q#3???==??E{{T%[[(HKKMME%	;C!  L( ( ( ( ( ( ( ( "8<==KRC    V,,        (a0x
 
 	
 " 	4444#%%  \\^^F	 

 

 "(	
 	
 	


 

 

 

 

 

 

 

 

 

 

 

 

 

 

 rw Jrw|Y$G$G 


 -* -&6G""K"g&&	= 	= 	= 	= 	= 	= 	= 	= 	= 	= "(<<>>	
 
 
 "m,,FE"""s   J==KKc                   dk    s
J d            ssddgt          | t                    sJ t          |t                    sJ t                    dk    sJ t                    dk    sJ t                    dk    sJ t          |                                          dv sJ |                                  |                                ^ }t          d|          \  }	}
t          d|          \  }}|                                 d         pd         p|
p||                                 ^ }t          |                                          }|                                }t          fdt          d         dz            D                       t          fdt          d         dz            D                       z  }|dk    rt          | ||          S fd	fd
}t          j        |                                 |||          }|S )Nr   divisor must be not zeror-   ro  r/   c              3     K   | ]V}t          t          |d                    t          d t          |d          z
  d                              z
  d          V  WdS rT  rU  rV  s     rg   r   z&avg_pool2d_backward.<locals>.<genexpr>  rW  ri   c              3     K   | ]V}t          t          |d                    t          dt          |d          z
  d                              z
  d           V  WdS rY  rU  rZ  s     rg   r   z&avg_pool2d_backward.<locals>.<genexpr>  rW  ri   r  c           	        t          j        d         t          j                  }t          j        d         t          j                  }t          j        d         t          j                  }t          j        d         t          j                  }t          j        d         t          j                  }t          j        d         t          j                  }t          j        t          j        | |          |          }t          j        t          j        ||          |          }	t          j        t          j        ||          t          j        t          j        t          j                  |                    }
t          j        t          j        |	|          t          j        t          j        t          j                  |                    }t          j	        |t          j        dt          j                            }t          j	        |	t          j        dt          j                            }	t          j        |
t          j        t          j                            }
t          j        |t          j        t          j                            }t          j        t          j        |
|          t          j        ||	                    }|S )z{
        This computes the scaling factor that we will divide an element
        by when `count_include_pad=False`
        r   r/   )
rR   r  r   r  r  r  r  r   r  r  )rf  rg  stride_hstride_wpad_hpad_wkernel_hkernel_wr  wstartr  wendr  heightr  r  r  rq  s                rg   !compute_pool_size_without_paddingz>avg_pool2d_backward.<locals>.compute_pool_size_without_padding  s   
 <q	5;77<q	5;77WQZ55WQZ55<A<<<A<<X..66X..66{GFH%%GCN65;77??
 
 {GFH%%GCN5%+66>>
 
 VS\!U[%A%ABBVS\!U[%A%ABB{4!D!DEE{4u{!C!CDDf 5 5swtV7L7LMMri   c                   | ^ }}}|d         z   }|d         z   }t          j        t          |d         z
  d         z   d                   t          j                  }t          j        t          |d         z
  d         z   d                   t          j                  }t          j        t          |d                   dz   t          j                  }t          j        t          |d                   dz   t          j                  }t          j        |t          j        dt          j                            }t          j        |t          j        dt          j                            }t          j        |t          j        t          j                            }t          j        |t          j        t          j                            }d }t                    D ]}	t                    D ]}
t          j	        |t          j        |	t          j                            }t          j	        |t          j        |
t          j                            }}n"ssd         d         z  }n ||          }t          j
         g |t          j        t          j        |t          j        |t          j        dt          j                                      d          t          j        t          j        |t          j        |t          j        dt          j                                      d                    |          }t          j        t          j        ||          t          j        ||                    }|5t          j        ||t          j        dt          j                            }t          j        |t          j	        ||          |          }|J |S )Nr   r/   Fr6  r  )rR   r  r)   r   r  r  r  r  r+  r   r  r  r  r  r  r  r  )rO  r  r  r[  r_  r`  ra  rb  rc  rd  re  rf  rg  r  partr  r$  r  r  rk  rl  r  r  r  rn  ro  r  rp  s                   rg   r}   zavg_pool2d_backward.<locals>.fn  s   A
N
N.QQ'&)3VAY??
 
 .QQ'&)3VAY??
 
 x6!95595;GGx6!95595;GG+gs|Au{'C'CDD+gs|Au{'C'CDDE3>-#M#MNNE3>,#L#LMM'' *	R *	RC]++ )R )RWWcl3&D&DEEWWcl3&D&DEE#/,EE& Fk F'N[^;EE==b"EEE{K#1 #$&s|Au{7S7S(T(T!" !" !.&+    1 #$&s|Au{7S7S(T(T!" !" !-&+   & ) . xF2u%%F2u%%  #"ytS\#u}5U5UVVHH"yswx/F/FQQHHS)RT ###ri   r  )rp   rC   r   r   r  r  ru  rn   r   r{  r+  fallback_avg_pool2d_backwardr@   r  rs   )rt  r   r  r  r  r  r  r  r"  _h_out
ceil_mode1_w_out
ceil_mode2r]  r   r  r}   r  r$  rk  rl  r  r#  rn  ro  rp  rq  s     ``` ``          @@@@@@@@@rg   avg_pool2d_backwardr,    s?    #'71'<'<'<>X'<'<<  a&k9-----a#####{q    v;;!w<<1qzz||&&&&

Q%; FJ &eQVWiXXFJ))++K!*F
FjFJK&1&:&:&<&<#Q|AJJLL!!HKKMME     {1~)**    M      {1~)**    M
  -/KR+	
 	
 		
        8? ? ? ? ? ? ? ? ? ? ? ? ? ? ? ?B 
	%%''	
 
 
B Iri   c                j    dk    s
J d            ssg dt          | t                    sJ t          |t                    sJ t                    dk    sJ t                    dk    sJ t                    dk    sJ t          |                                          dv sJ |                                  |                                ^ } t          d|          \  }	}
t          d|          \  }}t           d|          \  }}|                                 t                    p|
p|p||                                 ^ }t          |                                          }|	                                }fdt          d          D             \  z  z  }|d	k    rt          | ||          S  fd
fd}t          j        |                                 |||          }|S )Nr   r  )r   r   r   r   )r   r   r/   r-   c              3     K   | ]7t          fd t                   dz            D                       V  8dS )c           
   3     K   | ]@}t          |         z  t          d |         z
           z            z
  d          V  AdS rT  )r{  )r   r(  rw   r  r  s     rg   r   z0avg_pool3d_backward.<locals>.<genexpr>.<genexpr>  sj       
 
 VAYQ[^);q	(I!J!JJANN
 
 
 
 
 
ri   r-   N)r{  r+  )r   rw   r  r  s    @rg   r   z&avg_pool3d_backward.<locals>.<genexpr>  s       3 3
 	 	 
 
 
 
 
 
;q>A-..
 
 
 	
 	
3 3 3 3 3 3ri   }   c           	     @   d D             \  }}}d D             \  }}}d D             \  }	}
}d t          | ||g|||g|||g          D             \  }}}d t          |||g|	|
|gg|||g          D             \  }}}d |||fD             \  }}}d t          |||gg          D             \  }}}t          j        t          j        t          j        ||          t          j        ||                    t          j        ||                    }|S )Nc              3  T   K   | ]#}t          j        |t          j                  V  $d S r   rR   r  r   r  r   s     rg   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  s0      'U'UQ(D(D'U'U'U'U'U'Uri   c              3  T   K   | ]#}t          j        |t          j                  V  $d S r   r3  r  s     rg   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  s0      MMs|Au{;;MMMMMMri   c              3  T   K   | ]#}t          j        |t          j                  V  $d S r   r3  )r   r  s     rg   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  sA       (
 (
-.CLEK(((
 (
 (
 (
 (
 (
ri   c              3  n   K   | ]0\  }}}t          j        t          j        ||          |          V  1d S r   )rR   r  r  )r   r  r   pads       rg   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  sR       "
 "
1c GCGAqMM3''"
 "
 "
 "
 "
 "
ri   c           
   3     K   | ]a\  }}}}t          j        t          j        ||          t          j        t          j        |t          j                  |                    V  bd S r   )rR   r  r   r  r   r  )r   r  r  r   r7  s        rg   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  sw       

 

 #q#s Kq!!373>#u{+K+KS#Q#Q 

 

 

 

 

 

ri   c              3  z   K   | ]6}t          j        |t          j        d t          j                            V  7dS r  rR   r  r  r   r  )r   r  s     rg   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  sP       "
 "
 Ks|Au{;;<<"
 "
 "
 "
 "
 "
ri   c              3     K   | ]9\  }}t          j        |t          j        |t          j                            V  :d S r   rR   r  r  r   r  )r   r  r   s      rg   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  sT       
 
S KS^C==>>
 
 
 
 
 
ri   )r*  rR   r  r  )pdrf  rg  stride_dr  r  pad_dr  r  kernel_dr  r   dstartr  r!  dendr  r"  r  depthr#  r  r  r  rq  s                      rg   r$  z>avg_pool3d_backward.<locals>.compute_pool_size_without_padding  s   'U'Uf'U'U'U$(HMMWMMMue(
 (
2=(
 (
 (
$(H"
 "
 Rx8<ueU>S "
 "
 "


 

 '*(8X.&u%	' '	

 

 

dD"
 "
 &&1"
 "
 "

 
tT 2UFE4JKK
 
 
dD GCGD&))374+@+@AA374QWCXCX
 
 ri   c                X   | ^ }}}}d t          |||g          D             \  }}}d t          |||g!          D             \  }}}d t          |||g!          D             \  }}	}
d |||fD             \  }}}d t          ||	|
g g          D             \  }}	}
d }t                    D ]o}t                    D ][}t          "          D ]G}d t          |||g|||g          D             \  }}}}n,ssd         d         z  d	         z  }n |||          }t          j         g |t          j        t          j        |t          j        |t          j        dt          j	                                      d
          t          j        t          j        |t          j        |	t          j        dt          j	                                      d
          t          j        t          j        |t          j        |
t          j        dt          j	                                       d
                    |          }t          j
        t          j
        t          j        ||          t          j        ||	                    t          j        ||
                    }|5t          j        ||t          j        dt          j                            }t          j        |t          j        ||          |          }I]q|J |S )Nc              3  &   K   | ]\  }}||z   V  d S r   r   )r   r  r7  s      rg   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>  s*      AAvq#1s7AAAAAAri   c              3     K   | ];\  }}}t          j        t          ||z
  |z   |          t          j                  V  <d S r   rR   r  r)   r   r  )r   r  r  r   s       rg   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>  s\       %
 %
1a N8AEAIq115;??%
 %
 %
 %
 %
 %
ri   c              3  |   K   | ]7\  }}t          j        t          ||          d z   t          j                  V  8dS rV  rG  )r   r  r   s      rg   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>  sU       
 
1 N8Aq>>A-u{;;
 
 
 
 
 
ri   c              3  z   K   | ]6}t          j        |t          j        d t          j                            V  7dS r  r:  )r   pstarts     rg   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>  sP       %
 %
 KQ < <==%
 %
 %
 %
 %
 %
ri   c              3     K   | ]9\  }}t          j        |t          j        |t          j                            V  :d S r   r<  )r   pend
pooled_dims      rg   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>  sT       
 
 j KcnZEEFF
 
 
 
 
 
ri   c              3     K   | ]9\  }}t          j        |t          j        |t          j                            V  :d S r   )rR   r   r  r   r  )r   rJ  p_s      rg   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>  sT       " "&FB R(E(EFF" " " " " "ri   r   r/   r-   Fr6  r  )r*  r+  rR   r  r  r  r  r  r   r  r  r  r  r  r   )#rO  r  r(  r  r[  pdstartr_  r`  pdendra  rb  rc  pd_rd  re  r=  rf  rg  r  r&  r  r$  r  d_window_sizer  rk  rl  r  r  r  pooled_depthrn  ro  r  rp  s#                        rg   r}   zavg_pool3d_backward.<locals>.fn  s   AqAAaAY)@)@AAA1a%
 %
1ay+v>>%
 %
 %
!'

 
Q1Iv..
 
 
ue
%
 %
"GW5%
 %
 %
!'
 
$'u%m\'R% %
 
 
ue '' 8	V 8	VC]++ 7V 7V // 6V 6VC" "*-$gw7#sC+ +" " "JBB (3 0* N+ N +AQ ?+a. P A A"b" M M;#!' # 5$'K(*CGE3<5;;W;W,X,X%& %& %1*/!" !" !" !$ 5$'K(*CGE3<5;;W;W,X,X%& %& %2*/!" !" !"  !$ 5$'K(*CGE3<5;;W;W,X,X%& %& %1*/!" !" !"! 4 7 D< 8E!2!2CF2u4E4EFFr5)) D  '#&9 $S%-(H(H$ $ $'9T378T3J3JH#U#Um6V7Vp ###ri   r  )rp   rC   r   r   r  r  ru  rD  rn   r   r+  fallback_avg_pool3d_backwardr@   r  rs   )!rt  r   r  r  r  r  r  r  ry  _d_outceil_mode_dr(  ceil_mode_hr*  ceil_mode_wr"  r]  r   r  r}   r  r$  rS  rC  rk  rl  r  r#  rT  rn  ro  rp  rq  s!     ``` ``             @@@@@@@@@@@@rg   avg_pool3d_backwardrZ  ^  s9    #'71'<'<'<>X'<'<<  ))k9-----a#####{q    v;;!w<<1qzz||&&&&$%JJLL!VUFE&q+vw	 FK '; FK 'q+vw	 FK ))++Kg,,K+KKK4?4H4H4J4J1Qm\AJJLL!!HKKMME3 3 3 3 3
 q3 3 3/M=-  -/-?KS+	
 	
 		
# # # # # # # # # #JU U U U U U U U U U U U U U U U U Un 
	%%''	
 
 
B Iri   c                   |                                  }t          |t                    r|g}n|st          t	          |                    }t	          |          dk    r t          |          dv sJ d|             g S t          |          }t          t	          |                    D ]}||         dk     r.||xx         t	          |          rt	          |          ndz  cc<   d||         cxk    rt	          |          k     s#n t	          |          dk    r||         dk    sJ t	          t          |                    t	          |          k    s
J d            |S )Nr   )r   r  r  zinvalid axis: r/   zreduction axis not unique)r   rp   r   r+  r   r   rn   r'   )r   r  r  rw   s       rg   _validate_reduction_axisr\  &  sb   ::<<D$  v  SYY
4yyA~~T{{////1H$1H1H///	::D3t99 M M7Q;;GGGCII4s4yyy14GGGDG''''c$ii''''CIINNtAw!|||Lz$  CII---/J---Kri   )r  c                  |t          | |          } |                                 t          t                   t	          | |                    }g }g g g t          t                              D ]g}||v r1                    |                               |                    7                    |           |                    |                    hd|dv rt                    dk    rt          |           rt          | j
        t                    rdnt          | j
        t          j                  sCt          | j
        t          j                  rat          | j
        j
        t          j                  r=|                                 }|                                p|                                 fd}	r)t'                    }
D ]}t(          j        j        |
|<   n|}
|                                 t1          |                                 |p|                                 |                                 |	|
          S )NF)r,  argminr/   Tc                    t          |          t                    k    sJ 
r0t                     t                    k    sJ  fdD              t                     t                    k    sJ d gt                     t          |          z   z  }t          j        t                     t          |                    D ]
\  }}|||<    	|          }rid |D             }|d         }t	          dt          |                    D ]}||         z  ||         z   }|t          j        |t          j                  fS |S )Nc                     g | ]
}|         S r   r   )r   rw   r_  s     rg   r   z9_make_reduction_inner.<locals>.loader.<locals>.<listcomp>a  s    000!U1X000ri   c                6    g | ]}t          j        |          S r   r  r  s     rg   r   z9_make_reduction_inner.<locals>.loader.<locals>.<listcomp>l  s     ???!el1oo???ri   r   r/   )	r   r(  r)  r*  r+  rR   r  r   r  )r_  reduction_indexr  rO  varr  rindex
linear_idxrw   inner_loaderkeepdimskept_idxreduced_idxreduced_sizesshould_compute_logical_indexr  s   `        rg   r  z%_make_reduction_inner.<locals>.loader]  su   ?##s;'7'77777 	1u::T****0000x000E5zzS]]****Fc%jj3+?+??@	!%  #k?"C"C
 
 	! 	!HC !IcNNY'' ( 		D?????F  J1c&kk** G G'-*::VAYF

3>*ekBBCCri   )rz   r*  r  r  r  r+  )r  r   r'   r   r\  r+  r   rt   r;   rp   rr   r?   r2   r   r   Bufferr  is_transposedis_contiguousrn   r   rO  rP  ru  r8  rs   r   )r   r  rg  r   rh  r  
kept_sizesrw   r  r  r]  rf  rh  ri  rj  rk  r  s     `        @@@@@@rg   _make_reduction_innerrp  8  sx    Q::<<Dc?3At<<==DJHMK3t99 ' '99q!!!  a))))OOAd1g&&&& $) ...""aLL # afk** 	+/(( 233 	qvr}--	2<QV[")2T2T	 \\^^F$$&&Df.B.B.D.D*D )          6  :: 	& 	&A'+HQKK	& ==??L||~~'81;;==++--&   ri   r  rT   c                     dd d fd}|S )NFr   c                   t          | |||          }t          j        d| d|}t          |j        j        t                    r|                                 |S )N)r  rg  r   rh  r  )r  r)  r   )rp  rA   r  rp   rr   r  )r   r  rg  r   r  r  rh  r  s         rg   r  zmake_reduction.<locals>.inner  s}    &"7)
 
 
 !XAXXQWXXK
 
 	 NNri   r   r   )r  rh  r  s   `` rg   make_reductionrs    s:    T        " Lri   c                  |t          | |          } t          | |          }t          |                                 |                                 f|                                 f|                                 |          S )N)rz   dtypes	inner_fnsr  r  )r  rv  r8  rs   r   ru  r   )r   r  r   s      rg   _make_scan_innerrw    sr    QD!!D||~~==??$ZZ\\   ri   r   c               r   |t          | |          } |                                 t          | |          }|                                 }|t          j        t          j        fv rt          | t          j                  } t          | ||          }t          fd|D                       }t          j        ||                                 |                                           }t          j        |t          |                                                    }t          t!          ||          |          S )Nc              3  (   K   | ]}|         V  d S r   r   r   rw   r  s     rg   r   zmean.<locals>.<genexpr>  '      00a$q'000000ri   r^  )r  r   r\  r   r   r  r  r[  sum_rP   r2   r9   rs   r8   r  rn   div)r   r  keepdimr   r  
sum_resultdenomr  s          @rg   r  r    s    Q::<<D#At,,D;;==Lu~666Q$$aw''J0000400000Ee1;;==XXXEeT**=*=*?*?%@%@AAEC
E**L999ri   c                  
 |d}|                                  
t          | |          }t          | |d          }|r|                                 t	          t          | |                    }t          |||          }t          
fd|D                       }|rt          j	        ||z
  d          }t          j        ||                                 |                                           }t          j        |t!          |                                                     }t#          ||          }	|s|	fS |r|nt%          ||          }|	|fS )Nr/   T)r~  c              3  (   K   | ]}|         V  d S r   r   rz  s     rg   r   z var_mean_sum_.<locals>.<genexpr>  r{  ri   r   r^  )r   r\  r  r  squarer  r|  rP   r   r  r2   r9   r   rs   r8   r  rn   r}  r-  )r   r  
correctionr~  return_meanx_meandiffsr  r  x_varr  s             @rg   var_mean_sum_r    sI   
::<<D#At,,D!T4(((F 3q&>>""EeT7++J0000400000E 1	%*,a00e1;;==XXXEeT**=*=*?*?%@%@AAE
E""E x9VVGFD$9$9F&=ri   c                   t          | |          }t          | ||d d           }|d         }t          |d                   }t          |t          j                  o/t          |          t          j        k     ot          |          dk    S )Nr  rg  r   rh  r  r+  r/   )	r\  rp  rP   rp   r   r  r   r0   r3  )r   r  r~  r  r  reduction_numels         rg   use_two_step_variancer    s    #At,,D"	wd$  F HF#F+=$>??O?EM22 	'  6#EE	'&!!Q&ri   c               j   dt          | ||d d           }|                    d          }|                    d           |                    d           t          j        j        d|fd|                                 d|\  }}}	|                                 |                                 |                                 t          | |          }t          fd|D                       d	 fd
}
 t          |
          |          }|r|                                 ||fS |fS )Nr/   r  r  r*  r  welford_reduce)rv  r  r   c              3  (   K   | ]}|         V  d S r   r   rz  s     rg   r   z$var_mean_welford_.<locals>.<genexpr>  s'      11q47111111ri   c                    t          | t          j                  r9| j        s2t	          j        t	          j        | t          j                  |          S t	          j	        | |          S r   )
rp   r   r   	is_numberrR   r  r  r   r  r  r`  s     rg   get_constant_or_index_exprz5var_mean_welford_.<locals>.get_constant_or_index_expr   sS    a$$ 	GQ[ 	G<q%+ > >FFF|Au%%%ri   c                               }           }t          j        d          }| t          j        |||z
            z  S r  )rR   r  r  )rr   cNzeror  r   r  rnumels       rg   r  z#var_mean_welford_.<locals>.scale_fn  sS    &&z599&&vu55|Au%%ck$A....ri   r   )rp  r  r2   WelfordReductionr  r   r  r   r\  rP   r  )r   r  r  r~  r  r  r  r  m2r"  r  rc  r   r  r  r  s     `         @@@@rg   var_mean_welford_r    sz   
"	wd$  F ZZ
##F
JJ{
JJ{%, )'kkmm  	 KD"a JJLLLKKMME::<<D#At,,D1111D11111F& & &
/ / / / / / / / #.
"
"2
&
&C Dy6Mri   c               8   |                                  t                    }t          | |d          } t          | ||||          }t	          | ||          rt          di |nt          di |}t          fd|D                       }|s|d         n|S )NFr  )r   r  r  r~  r  )r  r~  c              3  <   K   | ]}t          |d           V  dS )Fr  Nr  )r   r   r  s     rg   r   z#var_mean_helper_.<locals>.<genexpr>#  s2      FF!8Ayu555FFFFFFri   r   r   )r   r   r  r8  r  r  r  r   )	r   r  r  r~  r  r   r  rV  r  s	           @rg   var_mean_helper_r    s    I))44MM...A
  F !w???	)(((( 
 FFFFvFFFFFF'36!99V3ri   )r  r~  c               *    t          | |||d          S )NFr  r  r~  r  r  r   r  r  r~  s       rg   var_r  '  s$    	W%   ri   c               *    t          | |||d          S )NTr  r  r  s       rg   var_meanr  .  s$    	W$   ri   c                .   |dk     r$t          t          j        |           | |          S |dk    rt          j        d|          S |dk    r| S t          | |dz  |          }t          j        ||          }|dz  dk    rt          j        ||           }|S )Nr   r/   r-   )pow_recursiverR   r  r  r  )r   rW  r   r  s       rg   r  r  5  s    1uuS^A..E:::Avv|Au%%%Avv1a1fe,,FWVV$$F	A!||##Mri   c                ,    t          j        | |          S r   )rR   r  r  rU  s     rg   
pow_nativer  D  s    71a==ri   )r  c                Z    t          t                    r1                                rt           t	                              S t          t                    rdk    rt                     S t          t                    rdk    rt                     S t          d  fD                       }t          |          }t          t                    odcxk     odk     nc p|odk    }|ri 	                                 fd}t          j                                                                          |                                           S t           t                    rL dk    rt!          d          S  d	k    r0t#                                                    rt%                    S |rZt           t                    rt'                     S t          t                    rt)                     S t+                     S t-                     S )
Nr  r/   c              3  r   K   | ]2}t          |t          j                  |                                V  3d S r   )rp   r2   rC   r   r   s     rg   r   zpow.<locals>.<genexpr>\  s;      NN1*Q2M2MNNNNNNNri   i    r   c                \    t           |                                                     S r   )r  r   )rO  r  rU  r  s    rg   r}   zpow.<locals>.fnf  s%     a???ri   r  r-   )rp   r[  r   r  r   sqrtr  re  r    ru  r@   r  rs   r   r   r!   r7  r   exp2fallback_pow_scalarfallback_pow_tensor_scalarfallback_pow_tensor_tensorr  )r  rU  r   is_integer_powembed_exponentr}   r  s   ``    @rg   r  r  R  s1   !U  1c!ff~~	Au		 !s((Aww	As		 QQxx NNANNNNNE%e,,N  3'' a"32AF   
	@ 	@ 	@ 	@ 	@ 	@ 	@ <<>>++--::<<	
 
 
 	
 !V 66Q??"66nQ[[]]33677N 4a   	4&q!,,,6"" 	4-a333-a333ari   c                <   t          | t                    r| j        }n| }t          |t                    r|j        }t          |t          j                  st          j        |                                 |                                 |	                                | 
                                          }t          |t          t          f          sJ |j        }t          |t          j                  sJ t          |t          j                  ri|                                sU|                                sAt          |j        t          j                  s"|                                 |j        |_        | S t          j                            |||           | S )Nr  unsafe_alias)rp   rC   rr   r2   r   r@   r  rs   r   ru  r   r6   r<   is_input_bufferis_module_buffer	NopKernelr  rf  realize_into)changedr2  r  changed_datare   s        rg   rY  rY    s   '9%% |#y!! hc2=)) 
.%%''##%%__&&##%%	
 
 
 $: 677777i#r}-----,.. 	$$&&	 ((**	 l'66		 	H!..\ /    Nri   c                >    t          | t          | |                    S r   )rY  r7  )r   r  s     rg   r  r    s    Q	!Z00111ri   c                    | |u r| S t          ||                                           }t          ||                                           }t	          ||                                           }t          | |          S r   r  rs   r  r   rS  r   rY  )r  rZ  r  s      rg   r}  r}    sg    
czz

C))
*
*C
3
(
(C
cllnn
%
%CS#ri   c                ,    t          j        | |          S r   )rR   floordivr  s     rg   r  r        <1ri   c                ,    t          j        | |          S r   )rR   truncdivr  s     rg   r  r    r  ri   c                   t          |           ot          |          }t          |           ot          |          }|dk    r;|r
J d            |rt          | |          nt          t	          | |                    S |dk    r;|r
J d            |rt          | |          nt          t	          | |                    S t	          | |          S )Nr@  z5floordiv operands can not be boolean at the same timerE  z5truncdiv operands can not be boolean at the same time)r   r   r  r@  r}  r  rE  )r  rU  rounding_modeboth_integerboth_booleans        rg   div_moder    s    "1%%</!*<*<L"1%%</!*<*<L XX!XXX!-Cx1~~~5Q3C3CCXX!XXX!-Cx1~~~5Q3C3CCq!99ri   c                    t          |           ot          |          }|rt          | |          S t          t          j        j                  } t          |          | |          S r   )r   logical_andr>   r  r  r  r  )r  rU  	both_boolr}   s       rg   r  r    s`    ""9q'9'9I (1a   *++!~b!!!Q'''ri   r  Optional[ir.Constant]c                L   t          | t          j                  rt          | j                  S t          | t          j                  r!t          |                                           S t          | t          j                  r| S t          | t          j                  sdS t          j
        j                            |                                           }t          j        |          5  t!          j        t          j        dd          5   | j        |                                  }ddd           n# 1 swxY w Y   ddd           n# 1 swxY w Y   t          |t          j
        j        j                  sJ t          |j        t          j                  r|j        S dS )z:Try convert an arbitrary IR node into an ir.Constant valueNallow_indexingT)rp   r2   r<   get_constant_valuerr   r6   r  r  Loopsr   	_inductorops_handlerExtractConstantsHandlerrs   rS   set_ops_handlerr   objectr  r  inner_fn_argsvirtualizedOpsValuer  )r   r  ru   s      rg   r  r    s   
 !R]## *!!&)))!R[!! 3!!--//222!R[!!  a"" to)AA!,,..QQG	'""- -R&(8$??- - aj!//++,	- - - - - - - - - - - - - - - - - - - - - - - - - - - - - - c5?6?@@@@@#)R[)) y4s6   ,!ED6*E6D:	:E=D:	>EEEc                   t          d | |fD                       }|rt          | |          S t          |          x}j|                                 j        dk    rM|j        dk    r(t          j        t          d          |j                  }n
d|j        z  }t          | |          S d } t          |          | |          S )Nc              3  R   K   | ]"}t          |          pt          |          V  #d S r   )r   r   r   s     rg   r   zdiv_prim.<locals>.<genexpr>  s7      OO1oa((>OA,>,>OOOOOOri   r   r   infr  c                     t          j        |  S r   )rR   r  r  s    rg   r}   zdiv_prim.<locals>.fn  s    {D!!ri   )rC  r  r  rs   r   r  mathcopysignr[  r  r  )r  rU  is_integralr  r  r}   s         rg   r  r    s    OOAOOOOOK 1~~ &a(((5!,,..:MQV:V:V =AuU||W]CCJJw},J1j!!!" " " >"a###ri   c                b    t          | |ft          j                  \  } }t          | |          S r  )rj  r   INT_TO_FLOATr  r  s     rg   r}  r}    s9     	
A$C$P  DAq Aq>>ri   c                    t          |           pt          |           }|rd }nd } t          |          | |          S )Nc                ,    t          j        | |          S r   )rR   modr  s     rg   r}   zfmod.<locals>.fn&  s    71a== ri   c                ,    t          j        | |          S r   )rR   fmodr  s     rg   r}   zfmod.<locals>.fn+  s    8Aq>>!ri   )r   r   r  )r  rU  r  r}   s       rg   r  r     sc    !!$$:(:(:K "	! 	! 	! 	!
	" 	" 	" >"a###ri   c                   t          |                                           s!t          |                                           r|t          j        }t          d|          } || |||          S )Nr"  r  r   r    r   r   r   r  rs  r   r  rg  r   r}   s        rg   r|  r|  1  sk     	''+;AKKMM+J+J
-	U	;	;	;B2axu----ri   c                   t          |                                           s!t          |                                           r|t          j        }t          |                                           dk    r.|dv sJ |p|                                 }t          | |d          S d }t          | ||          }t          j
        j        d	i |d|i\  }|t          | ||          S |S )
Nr   r   rr  Tr  c                >    | \  }|\  }t          j        ||          fS r   )rR   r   a_tupleb_tupler  rU  s       rg   
combine_fnzcumsum.<locals>.combine_fnO  #    1ri   r  r   r  r   r   r   )r    r   r   r   r  r   r   r  rw  r2   Scanr  fallback_cumsumr   r  r   r  r  r  s         rg   cumsumr  C  s     	''+;AKKMM+J+J
-
1::<<Aw&5t,,,,     
 ad%888F????J???IV~qd%8888Mri   c                   t          |                                           s!t          |                                           r|t          j        }t          |                                           dk    r.|dv sJ |p|                                 }t          | |d          S d }t          | ||          }t          j
        j        d	i |d|i\  }|t          | ||          S |S )
Nr   r  Tr  c                >    | \  }|\  }t          j        ||          fS r   )rR   r  r  s       rg   r  zcumprod.<locals>.combine_fng  r  ri   r  r  r  r   )r    r   r   r   r  r   r   r  rw  r2   r  r  fallback_cumprodr  s         rg   cumprodr  [  s     	''+;AKKMM+J+J
-
1::<<Aw&5t,,,,     
 ad%888F????J???IV~t59999Mri   c                *   d }|                                  }t          |                                           dk    r|dv sJ t          |           S t	          | ||          }t          j        j        di |d|i\  }|t          | |          S |S )Nc           	         | \  }|\  }t          j        ||          }t          j        ||          }||k    t          j        |           z  }t          j        |t          j        t          j        ||z
                      |z   |          fS r   )rR   r  r  r5  r  log1pexp)r  r  r  rU  min_vmax_vr  s          rg   log_add_exp_helperz(logcumsumexp.<locals>.log_add_exp_helperu  s    Aq!!Aq!!CIe$4$4#45	$	#'%%-*@*@ A AE I1MMOOri   r   r  r  r  r  r   )	r   r   r   r  rw  r2   r  r  fallback_logcumsumexp)r   r   r  r   r  r  s         rg   logcumsumexpr  s  s    P P P KKMME
1::<<Ag~~~~Qxxac777FGGGG4FGGGIV~$QC0000Mri   c                   t          |                                           dk    r0dv sJ t          |           t          | t          j                  fS |                                 }t          j        d|d          }t          | |          }|t          j        f|d<   | 
                                fd	f|d
<   t          j        j        di |d|i\  }}|t          |           S ||fS )Nr   r  r   r,  Fr   arg_break_ties_leftr  ru  c                N    t          j        |          t          j                  S r   rR   r  r   r  rO  r  s    rg   rj  zcummax.<locals>.<lambda>      CN3t9ek:: ri   rv  r  r  r   )r   r   r  r  r   r  r   r2   get_reduction_combine_fnrw  ru  r  r  fallback_cummaxr   r  r   r  r  r&  r  s    `     rg   cummaxr    
   
1::<<AwQxxAU[99999KKMME,5  J ad%888Fu{+F8	::::F; gnEEvEE*EEEOFG~qd++++7?ri   c                   t          |                                           dk    r0dv sJ t          |           t          | t          j                  fS |                                 }t          j        d|d          }t          | |          }|t          j        f|d<   | 
                                fd	f|d
<   t          j        j        di |d|i\  }}|t          |           S ||fS )Nr   r  r   r^  Fr	  r  ru  c                N    t          j        |          t          j                  S r   r  r  s    rg   rj  zcummin.<locals>.<lambda>  r  ri   rv  r  r  r   )r   r   r  r  r   r  r   r2   r  rw  ru  r  r  fallback_cumminr  s    `     rg   cumminr    r  ri   c                   t          |                                           s!t          |                                           r|t          j        }t          d|          } || |||          S )Nr  r  r   r  r  s        rg   r  r    sk     	''+;AKKMM+J+J
-	e	<	<	<B2axu----ri   c                l    t          | t          j                  }  t          d          | ||          S )NrD  r  rg  )r  r   r\   rs  r   r   r~  s      rg   
reduce_anyr    s2    EJA >%  w????ri   c                r    |$t          | ||          t          | ||          fS t          | d |          S Nr  )reduce_amaxreduce_argmaxr  s      rg   
reduce_maxr!    L    
g666!#888
 	

 qtg6666ri   c                r    |$t          | ||          t          | ||          fS t          | d |          S r  )reduce_aminreduce_argminr  s      rg   
reduce_minr&    r"  ri   xor_sumr{  rz  r,  r  r^  
logical_or)r  r  stabler   
descendingc          	        |d}|                                  }|                                 }t          t          |          |          }t          |          dk    r+t	          |           t          d|t          j        |          fS t          |          r||         nd}t          j	        j
                            |t          j        t          j                  j                  st          | |||          S t!          |ddt          j        |d          }dgt          |          z  }t          |          r|||<   t#          ||          }t%          ||          }t&          j                            || j        |j        f|                                 |                                f||||          \  }	}|	t          | |||          S |J |	t1          |t          j                  fS )NFr   r/   r)  )r  r  r   rz   r  )rz   ru  rv  r  r  r*  r+  )r   rs   r   r   r  r  r   r  rS   r_   r   statically_known_ltr  int16r{  sort_fallbackr  r  rS  r2   Sortr  r   ru  r  )
r   r*  r   r+  rL  rz   ru  r  
view_shaper&  s
             rg   sort_stabler2    s   ~JJLLE\\^^F
3u::s
+
+C
5zzQQxxq&%+u==== ZZ.uSzzQH7//%+ek:R:R:VWW OQv3:NNNNVSX  G s5zz!J
5zz #"
37J''GWe$$Ggnn'==??G$7$7$9$9: %  OFG ~Qv3:NNNN8GU[1111ri   c                (    t          | d||          S )NFr)  )r2  )r   r   r+  s      rg   sortr4    s    qCJGGGGri   c                <    t          | |t          j        |          S )Nr   r   r  )r  r   r  )rb   r   r  s      rg   register_pointwise_numericr7    s(    
;H'	   ri   torch._ops.OpOverloadPacketc                `    t          | j                   t          | t          j                  S r  )rO   r  r  r   r  r,  s    rg    register_pointwise_numeric_ldf64r:  %  s2    '444
;H   ri   r  )r  c                 	 t           j        st          S t          | ||t          j                  |                                 |                                |                                j        ot          j	        j
         		fd}t          j        |                                 ||                                           S )a  
    Computes self + value * tensor1 * tensor2 using FMA for better precision.

    Matches eager CUDA kernel order: self + value * (tensor1 * tensor2)
    This is computed as: fma(value, tensor1 * tensor2, self)

    Note: FMA is only used for floating-point types on non-AMD GPUs. For integer types,
    we fall back to regular arithmetic since FMA doesn't support integers.

    For floating-point types, we use mul_rn (round-to-nearest multiplication)
    to force rounding of the product before the FMA. This prevents Triton's
    compiler from fusing the multiplication with the FMA, matching eager's
    rounding behavior.

    When emulate_precision_casts is False, we return NotImplemented to use the
    decomposition instead.
    r  c                    |           } |           } 	|           }dk    r
rt          j        |||          S 
rt          j        ||          }nt          j        ||          }t	          t
          j                  rt          j                  }nt          j                  }
rt          j        |||          S t          j	        |t          j        ||                    S r  )
rR   fmamul_rnr  rp   r   r   r  r  r   )rO  self_valt1_valt2_valt1_times_t2
value_exprr   rn  	t1_loader	t2_loaderuse_fmar  s         rg   r  zaddcmul.<locals>.inner_fn\  s    ;s##33A::':7668444  	2 *VV44KK'&&11K eU[)) 	4u55JJeU33J 	G7:{H=== 78SWZ%E%EFFFri   r  )r0   r  NotImplementedr   r   rd  ru  r  r   versionhipr@   r  rs   r   )
r/  tensor1tensor2r  r  r   rn  rD  rE  rF  s
      ` @@@@@rg   r  r  8  s    & ) ;C	  E ""$$K##%%I##%%I %?em.?*?GG G G G G G G G G G<   }}	   ri   c                R   t           j        st          S t          t          j        j        j                  dk    p*t          j        j        j        t          v pt                      }t          t          | ||                    }fd}t          |t          |           ||          S )z
    Foreach version of addcmul with scalar value parameter.
    Uses foreach_group_loop for consistent grouping behavior.

    When emulate_precision_casts is False, we return NotImplemented to use the
    decomposition instead.
    r   c                    t          | diS )Nr  )r  )rx   r  s    rg   r  z)_foreach_addcmul_scalar.<locals>.apply_fn  s    *E***ri   )r0   r  rG  r   rS   r_   r`   ra   rc   r  rh   r|   r*  r  )r/  rJ  rK  r  r  r  r  s      `   rg   _foreach_addcmul_scalarrN    s     )  	AG &''1, 	,7&*==	,)++   D'7 ; ;<<F+ + + + + fc$ii?KKKri   logical_not)r  )r   r
  rh  identity)r  pointwise_overrides_datac              #  H  K   t           |         t          | j        d           }|d S fd}t          |t          j        j                  r>|                                D ]'}t          ||          }|j         ||          fV  (d S |j         ||          fV  d S )Nc                4    j         t          |           S d S r   )tritonr  )rb   rr   s    rg   make_triton_fallbackz6_get_pointwise_overrides.<locals>.make_triton_fallback  s!    ;#B''' ri   )	rQ  r   r   rp   r   r   r   r   r   )nsr   rb   rU  olnamer  rr   s         @rg   _get_pointwise_overridesrX    s      #D)D	TY	%	%B	z( ( ( ( ( "ej122 Ellnn 	I 	IFV$$Bd.0D0DR0H0HHHHHH	I 	I $*,@,@,D,DDDDDDDri   r6  c                z    | t           |<   t                              |            fd}t          | |           d S )Nc                      | i |}g }t          | d         |          D ]*\  }}|                    t          ||d                     +|S )Nr   Tr  )r*  rt   rY  )rx   r  resultsmut_resultsr   r  outplace_ops         rg   r}   z$register_foreach_inplace.<locals>.fnF  sk    +t.v..tAw00 	J 	JKCyf4HHHIIIIri   )rZ   r  r   r;  )aten_opoutplace_aten_opr]  r}   s     ` rg   register_foreach_inplacer`  B  sT    07,-G$$$     w+++++ri   c                @    t          | d           fd            }|S )Nr  c                      | i |}t          || d                                                   }t          | d         |          S r  )r  r   rY  )rx   r  r  r]  s      rg   r}   zregister_inplace.<locals>.fni  sJ    d-f--&$q'"3"3"5"566a&)))ri   )rJ  )r^  r]  r}   s    ` rg   register_inplacerc  h  s;    wD999* * * * :9*
 Iri   c                    d S r   r   ry  s      rg   sym_constrain_rangere        4ri   c                    t           j        j        j        d         }t	          |t
          j                  r|j        j        S t          |          S Nr2  
rS   r_   r`   r  rp   r   rm  re   rn  r   r  r   r2  s      rg   r  r    =    
'

#E
*C#u|$$ x}3xxri   c                    t           j        j        j        d         }t	          |t
          j                  r|j        j        S t          |          S rh  ri  rj  s      rg   
sym_striderm    rk  ri   c                *    |                                  S r   )r  r  s    rg   	sym_numelrp    s    ;;==ri   c                    t          j        |  S r   )r   Addr  s    rg   sym_sumrs    s    9dri   c                     t          d          )NzHelpful for debuggingr   )r/  rx   r  s      rg   foobarru    s    
5
6
66ri   c                H    |                                   t          |           S r   )r  r  r   s    rg   _realizerw    s    IIKKK88Ori   c                X    |                                   t          j        | |           | S r   )r  r2   ResizeStorageBytes)variabler]  s     rg   resize_storage_bytes_r{    s,    (H---Ori   c                    |                                   |                                  t          j        t          j        | |                    S r   )r  rC   r  r2   SetSourceTensorKernel)r/  source_tensors     rg   set__source_tensorr    s?    LLNNNB4T=IIJJJri   r}  c                    | |u r| S t          ||                                           }t          ||                                           }t	          ||                                           }t          | |          S r   r  )r  rZ  s     rg   
fsdp_copy_r    sg    #::JS^^--..sCMMOO,,S#,,..))c"""ri   c                 	
 t          | t                    sJ t          |t          t          f          sJ |t          j        }|t          j        k    rt          d|           |t          j        k    rt          |          dk    sJ |t          j
        k    rt          |          dk    sJ |                                 
|                                 }|                                 }t          | j        t          j                  r| j                                        | _        t	          j                    rat          j        j        j        rKt-          |          rt/          d          n.t1          |          rt	          j        |          j        ndndt6          j        j                            
d          rt?          |||          S tA          | 
gd	g          }|!                                	t          j"        #                    ||          }t          j$        ||||          %                                	
fd
}tM          j'        |||t          |                    }|S )Nzunsupported memory format: r   r   nanTr  r   rT  r/   c                     |           t          j        t          j                  }t          j        t          j                  }t          j        ||          }t          j        |fd          S )Nc                       g          S r   r   )
flat_indexflat_loaders   rg   rj  z*resize.<locals>.inner_fn.<locals>.<lambda>  s    ZL(A(A ri   )rR   r  r   r  r  r  )	rO  flat_index_exprlimitr  r  r  	old_numelout_indexeruninitialized_vals	       @rg   r  zresize.<locals>.inner_fn  sl     [%%
.U[AAy%+66vou--z$ A A A A ACTUUUri   r  )(rp   rC   rn   r   r   contiguous_formatpreserve_formatr  channels_lastr   channels_last_3dr  r   rf  rr   r2   r6   r  rd  rJ  deterministicfill_uninitialized_memoryr   r[  r    r  r{  rS   r_   r   rz  r  r  ru  r   stride_ordered_for_memory_formatr  r  r@   r  )r   r  r  r   rz   x_flat
out_strider  ru   r  r  r  r  s            @@@@rg   resizer    s   a#####dT5M*****/---HHHIII+++4yyA~~~~...4yyA~~~~IKKMME""$$F!&"+&& &##%% 	244 K%?  %   	% %ee$$ 	% %E 2 2 6 $  w//	1== ID+5HHHH		
 	
 F $$&&K"CCD-XXJ.jAANNPPKV V V V V V V V 
UXd4jj  C Jri   )auto_functionalizedc                    ddl m} |                    |          }t          j        | ||i ||           d |                                D             S )Nr   )kernel_side_table)
kernel_idxgridtma_descriptor_metadatakernel_argsc                D    i | ]\  }}t          |t                    ||S r   r  )r   rU  r2  s      rg   r   z'triton_kernel_wrap_.<locals>.<dictcomp>6  s-    RRRcz#y7Q7QRCRRRri   )*torch._higher_order_ops.triton_kernel_wrapr  get_constant_argsr2   UserDefinedTritonKernelr$  )r  constant_args_idxr  r  r  r  constant_argss          rg   triton_kernel_wrap_r  $  sz     MLLLLL%778IJJM 7/v//	    SRV\\^^RRRRri   3list[Union[ir.TensorBox, ir.ShapeAsConstantBuffer]]c                X   t          d | g|D                       rFd}t          j        j        j                            dd           x}r| d| }|t          j        _        t          j        	                    | |||          }t          t          t          j	        |                    S )Nc              3  ^   K   | ](}t          |t                    ot          |          V  )d S r   rp  r   s     rg   r   zcond.<locals>.<genexpr>?  s7      
M
Ma:a  1Yq\\
M
M
M
M
M
Mri   z"control flow operator: torch.cond.rH  rI  )rD  rS   r_   r`   r  r  rN  r2   Conditionalr  rn   maprC   )predtrue_fnfalse_fnoperandsr   rH  r  s          rg   r   r   9  s     
M
MD;L8;L
M
M
MMM 02'.377tLLL; 	988;88C,/)^""4(HEEFI$f--...ri   c                   t           j        sbt          d ||z   D                       rFd}t          j        j        j                            dd           x}r| d| }|t          j        _        t          j
                            | ||||          }t          |t                    sJ t          t          t          j
        j        |                    S )Nc              3  ^   K   | ](}t          |t                    ot          |          V  )d S r   rp  r   s     rg   r   zwhile_loop.<locals>.<genexpr>M  sK       * * 	1f.)A,,* * * * * *ri   z(control flow operator: torch.while_loop.rH  rI  )r0   rM  rD  rS   r_   r`   r  r  rN  r2   	WhileLoopr  rp   r   rn   r  _maybe_wrap_as_tensor_box)cond_fnbody_fncarried_inputsadditional_inputsstack_outputr   rH  r  s           rg   
while_loopr  I  s     ! 0c * *"33* * * ' ' 0 9'.377tLLL; 	988;88C,/)\  .*;\ F fh'''''BL:FCCDDDri   )r  subgraph_fnir.Subgraph
identifierc                ~    t          j        j        | g|R  }t          t	          t
          j        |                    S r   )r2   InvokeSubgraphr  rn   r  rC   )r  r  r  r  s       rg   invoke_subgraphr  b  s8    %k=H===FI$f--...ri   r  torch.fx.GraphModulec                   d}t          | j        j                  D ]0\  }}|j        dk    r0|t          j        j        vsJ ||         t          j        j        |<   A|j        dk    rTt          j                            |          \  }}t          j        j	        
                    t          j        |||          }|t          j        j        vsJ t          j        j        }	 |t          j        _        t          j                            |          t          j        j        |<   |t          j        _        # |t          j        _        w xY w|t          d          |S )a  Process nodes from a FX graph by executing them through V.graph.

    This is a common pattern for executing a subgraph's nodes:
    - Placeholder nodes are mapped to the provided args
    - Output nodes return their result
    - Other nodes are executed via V.graph.run_node

    Nr  rV  zNo output node found in graph)ro   r_   nodesrb   rS   envfetch_args_kwargs_from_envr   r  InterpreterrV  r`   run_noder  )r  rx   rV  rw   re   output_argsr  saved_current_nodes           rg   process_subgraph_nodesr  h  s4    F\/566 : :47m##qw{**** $QAGKW  "#'"D"DT"J"JKX)00$VTTFFqw{****!"!5:'+$$%G$4$4T$:$:D!'9$$'9$9999~:;;;Ms   #AD88E)control_depsc                   g }| D ]S}t          |t                    s|                                 |                    |                                           Tt
          j        j        j        }d}t          |          |z   t          |          k    sJ t          t
          j        j
                  }t          |j        j                            d                    t          |          k    sJ t          |j        t          |                    }|| sJ t
          j        j
        |d         D ]<}	|D ]7}
|	j        }|J t
          j        j        |                             |
           8=|S )aS  
    Lower control_deps_op by ensuring dependencies are realized and tracking them.

    The control_deps_op HOP makes dependencies explicit in the graph. During lowering:
    1. Realize all additional dependencies to ensure they're computed
    2. Execute the target operation normally
    3. Track the dependencies for the scheduler
    r-   r  r,  N)rp   r:   r  rt   r  rS   r_   r`   rx   r   
operationsr  
find_nodesr  rn   operation_nameadditional_buffer_depsr   )additional_depsr  rx   	dep_namesdeporiginal_args
arg_offsetoperation_lenrV  rb   dep_nameop_names               rg   control_deps_op_loweringr    sw    I ) )#v&& 	((((G(-MJt99z!S%7%77777*++M{'-88M8JJKKsSWyyXXXX $K$<d4jjIIF/1 g 0 B B! 	B 	BH'G&&&G*7377AAAA	B
 Mri   )schemec                  d }t           j        j        j                            dd           }|J t          | j        j        j                  D ]o\  }}|j        dk    r||         t           j        j	        |<   ,|j        dk    rt           j        
                    |          \  }}t          j        ||                                          D ]}	|	                                 |j        r6t           j        j                            |	                                           t           j        j                            |	                                           t(          j        j                            t           j        |||          }>t           j                            |          t           j        j	        |<   q|S )Nquant_optionsr  rV  )rS   r_   r`   r  r  ro   r  r  rb   r  r  r(  r)  r&  r  codegen_low_precisionlow_precision_codegen_opsr   r  invoke_quant_opsr   r  r  rV  r  )
r  r  r  rV  r  rw   re   rx   r  r  s
             rg   invoke_quant_tracerr    so   FG(-11/4HHM$$$[5;ABB 7 747m## (AGKW  7==dCCLD&_T6==??;; E E		 6 RG599!:N:N:P:PQQQ(,,Q-A-A-C-CDDDDX)00$fMMFF ! 0 0 6 6AGKMri   r  r  tuple[torch.Tensor]c                  	 ddl mm} t          |          dk    rt	          d          fdt          j        ||          D             } || |          		fd}t          |d         dd           }t          d |D                       |d	<   t          d
 |D                       |d<   t          j
        j        d|dd|}|d         t	          d          |S )Nr/   )InputDescriptorlower_pointwise_subgraphr   zSUnable to generate code for associative_scan op, because there are lifted argumentsc                r    g | ]3} |                                 |                                           4S )rT  )r   rs   )r   r   r  s     rg   r   z$associative_scan.<locals>.<listcomp>  sH        	akkmmALLNNCCC  ri   c                `     g t          j        |           t          j        |          R  S r   )r{  r  )lhsrhslowered_combine_fns     rg   wrapped_combine_fnz,associative_scan.<locals>.wrapped_combine_fn  sC    !! 
$$
$$
 
 
 	
ri   r  c              3  >   K   | ]}|                                 V  d S r   r  r   s     rg   r   z#associative_scan.<locals>.<genexpr>  s*      77qQ[[]]777777ri   ru  c              3  >   K   | ]}|                                 V  d S r   rt  r   s     rg   r   z#associative_scan.<locals>.<genexpr>  s*      <<A<<<<<<ri   rv  F)r  can_fallback_to_atenz/Unable to generate code for associative_scan opr   )r  r  r  r   r  r(  r)  rw  r   r2   r  r  )
r  xsr  r  subgraph_inputsr  r  r  r  r  s
           @@rg   associative_scanr    sS    MLLLLLLL
!!a
 
 	
   R((  O 21*oNN
 
 
 
 
 be!4888F77B77777F8<<<<<<<F;W^ %"   F
 ayLMMMMri   c                    d S r   r   )tokenss    rg   _sink_tokensr    rf  ri   c                     d S r   r   r   ri   rg   _make_tokenr    rf  ri   c                R   ddl m}m}  ||          }||t          j        j        j        u rddlm}m	} |
                                }	|	r|	j                            t          j        j        j                  }
|
rht          |
|          sJ |
                    |d                   }|r9t          |          dk    s
J d            t!          t#          |                    }t          t$          j        j                  }|t*          v r/t+          |         |i |}t-          j        t0          d |           n0d }t-          j        |t5          j        j        |g|R i |          }t          t$          j        j        |d                   dk    sJ d	| d
            |rt$          j        j                            |          }t$          j        j        |d         D ]\}d |_        |rP|                                 }t$          j        j!        |         "                    |                                            ]|t$          j        j        |<   	 d }t-          j        |||f          \  }} ||||          }nm# tF          $ r`}tI          |          }tJ          &                    d||           t          |tN          tP          f          r| g|R cY d}~S | |fcY d}~S d}~ww xY wt          |j)                  dk    r| |fS t          |j)                  dk    r| |fS | g|R S )z
    We lower the operator directly, and then we add StarDep dependencies to all
    the newly created nodes in the graph.
    r   )_get_effect_get_schemaN)InvokeSubgraphCacheTracingContextr/   zMultiple effects NYIc                *    |                                  S r   )r  ro  s    rg   rj  zwith_effects.<locals>.<lambda>*  s    !))++ ri   c                b    t          | t          j                  rt          j        |           n| S r   r  r   s    rg   r  z"with_effects.<locals>.wrap_tensors-  r  ri   zCNo operation nodes were generated when lowering effectful operator .c                     dS r   r   r   ri   rg   rj  zwith_effects.<locals>.<lambda>=  s    d ri   c                L   t          | t          j                  r|                                 S t          | t                    r`	 | j        }t          |d          r.t          |j        d          r|j                                        S n# t          t          f$ r Y nw xY w| S | S )Nrr   get_example)
rp   r2   TorchBindObject	get_valuerC   rr   r  r  AttributeErrorr   )r  r  s     rg   convert_ir_to_valuez)with_effects.<locals>.convert_ir_to_valueH  s    !R/00 {{}}$Ay)) fGw// :Gm5 5 :  '|77999&(;<   D Hs   AB BBz5Failed to get schema for %s: %s. Assuming list output)*torch._higher_order_ops.effectsr  r  r   rR   higher_orderr  torch._guardsr  r  try_gethop_dispatch_set_cache	get_cacherp   get_effectsr   re  r]  rS   r_   r  rX   r{  rI  rC   r  r2   r  r  effectful_opsr  has_side_effectsr  additional_star_depsr   r  r   r  r  r   rn   returns)tokenrb   rx   r  r  r  effect_typer  r  tracing_ctxinvoke_subgraph_cacheeffectsr  r  r  prev_effect_buffernew_opr  r	  schema_argsschema_kwargsschemae	error_msgs                           rg   with_effectsr	  	  s    IHHHHHHH +b//KrUY%;%KKKEEEEEEEE$,,.. 
	6$/$F$P$P	&6% %! % 6!"79LMMMMM/;;DGDD 6w<<1,,,.D,,,"&tG}}"5"5K *++M 
Y2///Y(=(=vFFFF	J 	J 	J "+22GGGGGG
 
 qw!-..122Q666SbSSS 766  
W266{CCg(8 	Y 	YF&2lF#! Y //++,W599:L:U:U:W:WXXX  	
k* #	 	 	$ &,_$&
 &
"] Rm<< # # #FF	CR	
 	
 	
 fudm,, 	##F########6?""""""# 6>av	V^			!	!vs+   *I> >
K(AK#K(K#K(#K()register_comm_loweringsregister_symm_mem_loweringsc                p   t          | |ddd          }|d         }t          j        j                            t          |                    }t          j        j        di |d|d\  }}|dk    rKt          j        j        	                    |t          j                  rt          j        d| d|d	|\  }}||fS t          j        t!          j        d
                     t%          | |d          }	t'          t(          j                 t-          | |	                    }
t/          |
|d          }|	|fS )zn
    Lowering inductor_prims.prepare_softmax_online to compute max/sum in one pass if no split is needed.
    TNr  r+  online_softmax_reduce)r  r  r/   r-   )r)  
num_outputreduction_hintz
            Online softmax is disabled on the fly since Inductor decides to
            split the reduction. Cut an issue to PyTorch if this is an
            important use case and you want to speed it up with online
            softmax.
            )rg  r   )rp  rS   r_   r   simplifyrP   r2   rA   
num_splitsstatically_known_geqr0   r3  r=   r  r  r  textwrapdedentr  rX   r  r  r  r|  )r   r   r  r+  r  hint	num_split
max_tensor
sum_tensorr  r  xsums               rg   prepare_softmax_onliner*	  w  sq   
 #	d$d  F 01W&&}5E'F'FGGFl-  
.   OD) A~~!'*??2 ~ "8!> "
Qt"
 "
?E"
 "

J :%% 	O 		
 		
 		
 1cD111!#a,,//Ct,,,Tzri   c                     t           j                                        o!t           j                                        dk    S )z.Check if we're on SM100+ hardware (Blackwell).)r   r   )r   cudais_availableget_device_capabilityr   ri   rg   _is_sm100_or_laterr/	    s.    :""$$V)I)I)K)Kw)VVri   c                   t                      st          d          |                                 }|t          j        t          j        t          j        fvrt          d|           |t          j        k    rt          | t          j                  } t          j
        t          j        ddt          j        dd          } t          |          |           }t          |t          j                  S )z
    Lowering for cvt_e8m0_rceil. Uses PTX cvt.rp.satfinite.ue8m0x2.f32 on SM100+.

    The PTX instruction takes 2 float32 and outputs 2 e8m0 packed in uint16.
    Currently we pass 0.0 as the second input and only use the low byte result.
    zFcvt_e8m0_rceil requires SM100+ (Blackwell) for PTX instruction supportzAcvt_e8m0_rceil requires float32, float16, or bfloat16 input, got z)cvt.rp.satfinite.ue8m0x2.f32 $0, 0.0, $1;z=h,rTr/   )asmconstraintsr   is_purepack)r/	  r   r   r   r  r  r  r  r  r6  rI  rR   inline_asm_elementwiseuint16r  r  )r   r   r}   r  s       rg   cvt_e8m0_rceil_loweringr7	    s      
!T
 
 	
 MMOOEU]EM5>BBBWPUWW
 
 	

 sEM**		"7l
 
 
B  ^B$$FFEK(((ri   )r  )quantized_lowerings)mkldnn_lowerings)jagged_loweringsc              #    K   t          | t          j        j                  s
J d            t                              |           }	  t          |           t          |                      dV  |r|t          | <   dS t                              |            dS # |r|t          | <   w t                              |            w xY w)z^
    A context manager to force fallback an op. Used in unit test
    for FallbackKernel.
    z+Only OpOverload to make the clean up easierN)	rp   r   r   r   rX   r  rJ  r  r  )rb   old_handlers     rg   force_fallbackr=	    s       b%*/00  5 0 --##K".r22333 	'IbMMMMM"  	'IbMMMM"s   )B )Cr6  )rj   rk   r[   rl   )r}   r~   r[   r   )r   r   r[   r   )r   r\   r   r   r[   r   )r}   r   r[   r   )r}   r   r   r   r[   r   )r   r   r[   r   )r   r	   r[   r   )r   r	   r[   r   )rx   r	   r   r   r   r\   r[   r   )rb   r   r   r   r[   r\   )r   rC   rz   r   r[   rC   )rx   r  r  r  r  r\   r   r	  r
  r\   r[   r  )r   r1  r2  r~   r[   r~   )
r2  r~   r  r\   r   r	  r
  r\   r<  rW   )r   r	  r[   rG  rp  )NNNFN)F)r   rC   r   r   r  r\   )r   rC   r   r   )r   rC   rz   r   )r   rC   r  rC   r   )r   rC   r  ra  r[   rC   )r   r   re  r/   Tr  )r  rC   r  rC   r  rC   r  r   r  r   r  r   r   r   r[   rC   )r  rC   r  rC   r  rC   r  r   r  r   r  r   r   r   r  r  r[   rC   )r  rC   r  r[  r  r   r  r   r  r   r   r   r[   rC   )r  rC   r  r[  r  r   r  r   r  r   r   r   r  r  r[   rC   )r  rC   r  rC   r  rC   r  r   r  r   r   r   r[   rC   )r  rC   r  rC   r  rC   r  r   r  r   r   r   r  r  r[   rC   )r   r   r/   )r  r   rA  r   rB  r   r  )T)r{   r  )re   r  )NTF)rz   r   )r  r  r  rC   r  r   r  r   )
r  r   r  r   r  r  r  rC   r  r   )r  rC   r[   r  )r  rC   r[   r%  )r.  rC   r/  rC   r*  r\   r+  r\   r,  r0  r-  r1  r[   rC   )r  rC   rE  rC   r*  r\   r+  r\   )r   r   r_  r   )r   NNr/   )NNN)rr  FF)r   r   )r  r1  r   r   r  r0  r  r\   )r   r   r  r0  )r   r   r  r\   )r-   F)r  r  r  r   r  r\   )r  r  )r  r  r  r  )r  r  r  r  r  r  )r   rC   r  r  r  r[  r[   r1  )rw   r  r  r  )rw   r  r  r  )rw   r  r  r  r  r  )Nr  N)
r/  rC   r  r5  r6  r5  r7  r8  r[   rC   )Nr   r/   F)r  )NNNN)r   r   FTN)r  rT   r   )r   r  r[   r  )rr  F)rb   r8  )r/   )r[   r  )r  r  r  r   )r  r  rx   r  )r  r  )r  r  r  r  r7  (  
__future__r   r  r  r6  r(  loggingr  r?  r  r#	  r  collectionsr   collections.abcr   r   r   r   typingr	   r
   r   r   r   r   r   typing_extensionsr   unittest.mockr   r   r   $torch.ao.quantization.fx._decomposedtorch.fxtorch.utils._pytreerJ  _pytreer{  torch._dynamo.utilsr   (torch._higher_order_ops.associative_scanr   r  r   "torch._library.fake_class_registryr   torch._library.utilsr   torch._prims_commonr   r   r   r   r   r   r   r   r   r    r!   torch.fx.experimental.sym_noder"   r#   ru  r$   r%   r&   torch.utils._ordered_setr'   torch.utils._sympy.functionsr(   r)   r*   r+   r,   _dynamo.utilsr.    r0   r1   r2   r3   decompositionr4   r5   r6   r7   r8   r9   r:   r;   r<   r=   r>   r?   r@   rA   rB   rC   rD   rE   rF   rG   rH   rI   rJ   rK   rL   rM   rN   rO   rP   rQ   r  rR   rS   r  rT   rU   rV   FALLBACK_ALLOW_LIST	getLoggerr  r  rX   __annotations__rY   r   r   rA  r  tr_c10dr  r   _higher_order_opsr  rd   r  rZ   quantized_decomposedrh   r|   r   r   r   r   r   r  r  r  r,  bmmconvolutionconvolution_backwardrL  rP  r{  r  r  r  _int_mmr  r
  r.  r  r  r  r  r  	complex32	complex64r\   r  r   r   r   r   r   r   r   r  r0  r;  rF  rd  rJ  rX  rj  r  r  r  r  r  r  r  r  r   r  r  
device_putr  r  r  r  r  r  r  r  r'  aliasdetachdetach_liftview_ofr!  r  r"  r-  r/  r3  r5  r;  r=  r@  rC  r  rE  rS  rP  rR  r`  _unsafe_viewreshaperd  slicer  r  r  quantize_per_channelr  r  r   r  _functional_assert_asyncr  dequantize_per_channelr  quantize_per_tensorr  dequantize_per_tensorr  r  r  r  r*  rU  rW  r[  rc  rh  rj  rm  rx  rK  r|  rv  r  r  cacher  r  r  r  r  r  rngprimsr  r  r  	bernoullir  r  r  	lru_cacher  r  r  r  r  r  r  r  r  randintstreamsrecord_event
wait_eventforce_stride_orderr  r  r  r  r
  lookup_seedr  randomr  r  r$  r(  rB  r  r9  	NO_OPMATHrL  rN  r   r8  rT  r   r   r  _adaptive_avg_pool3dadaptive_max_pool3d*_scaled_dot_product_attention_math_for_mpsuniformexponential_pdist_forwardsoft_margin_loss_backward_fused_rms_normxpur-	  embedding_dense_backwardmtia_is_compilednative_layer_norm_cdist_forward_cdist_backward
_trilinearsegment_reduce_segment_reduce_backwardhistc	histogrambin_ct_histogramdd_bin_edges_histogramdd_from_bin_ctsaddbmm_addmm_activation_grouped_mm
_cudnn_rnn_cudnn_rnn_backward
miopen_rnnmiopen_rnn_backward_embedding_bag_embedding_bag_forward_only_embedding_bag_backward*_embedding_bag_per_sample_weights_backward_fused_moving_avg_obs_fq_helper*_fused_moving_avg_obs_fq_helper_functional max_pool3d_with_indices_backward_adaptive_avg_pool2d_backward_adaptive_avg_pool3d_backwardadaptive_max_pool2d_backwardadaptive_max_pool3d_backwardfractional_max_pool2d_backwardfractional_max_pool3d_backwardreplication_pad1d_backwardreplication_pad2d_backwardupsample_linear1d_backwardupsample_bicubic2d_backwardupsample_trilinear3d_backwardgrid_sampler_2d_backward_pdist_backwardr4  r*  kthvaluetopkr  median	nanmedianrX  resize_
resize_as__linalg_detlinalg_householder_productlinalg_inv_exlinalg_ldl_factor_exlinalg_ldl_solve	linalg_lulinalg_lu_factor_exlinalg_lu_solvelinalg_matrix_exp	linalg_qr_linalg_slogdet_linalg_solve_exlinalg_solve_triangular_linalg_svd	lu_unpackormqr_linalg_check_errorslinalg_pinvatol_rtol_tensor_linalg_eightriangular_solvelinalg_cholesky_excholesky_inversecholesky_solvegeqrf_fft_r2cnonzerogcd_thnn_fused_lstm_cell_prims	rng_primsrun_and_save_rng_staterun_with_rng_stategraphsafe_run_with_rng_statemasked_scattermasked_scatter_backwardr  angle_efficientzerotensor(_sparse_coo_tensor_with_dims_and_tensors	to_sparse
_to_sparser   r  ru  rs  #_scaled_dot_product_flash_attention	quantized,_scaled_dot_product_flash_attention_backward#_scaled_dot_product_cudnn_attention,_scaled_dot_product_cudnn_attention_backward+_scaled_dot_product_flash_attention_for_cpu4_scaled_dot_product_flash_attention_for_cpu_backward0_scaled_dot_product_fused_attention_overrideable9_scaled_dot_product_fused_attention_overrideable_backward_flash_attention_forward_flash_attention_backward_efficient_attention_forward_efficient_attention_backwardindex_reducerepeat_interleave_weight_norm_interface_backwardr  r  r  r  r  r  r  r  scalar_tensorr  
LongTensorr  r  r  r  r  r7  r  r\  r  r  r  r  
zeros_liker  r  r  r  r  r  r  r  r  r'  r*  r)  r_  r9  rA  rC  rE  rP  rK  rU  r>  rr  fallback__unsafe_masked_indexrw  ,fallback__unsafe_masked_index_put_accumulaterr  r  ri  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r+  r  r  r  r  r  r  r  r0  r4  rA  rE  rG  rs  r  r  r  r  r  r  r  r  r  r  r  r  r  r  
avg_pool1dr  r  r'  rZ  rU  r\  rp  rs  rw  r  r  r  r  r  rc  r  r  r  r  r  Tensor_Tensorr  r  r  Tensor_Scalarr  rY  r  r}  r  r  r}  r  r  r  r  true_divider  r"  r|  r  r  r  r  r  r  r  r  r  r  r  rD  r  r{  r!  rz  r&  r'  r  r  r  r$  r,  r   r^  r%  r   r/  r2  r7  r:  rsqrtr  r  expm1relur  r  r  r  r  rN  r  cossinabsbitwise_andbitwise_left_shiftbitwise_not
bitwise_orbitwise_right_shiftbitwise_xorlgammaerfspecial_erfr  tantanhr  rO  r(  logical_xorr  r  	clamp_min	clamp_maxnegr  	remaindersignsignbit	_neg_viewler  r  r  r  necoshsinhacosacoshasinasinhatan2atanatanhr  erfcerfinvhypotlog10log2	nextaftercodegen.commonr  rQ  rX  r   rb   r   r  _foreach_addListforeach_add_listforeach_add_scalar_foreach_mulforeach_mul_listforeach_mul_scalar_foreach_sub_foreach_neg_foreach_abs_foreach_powScalarAndTensor_foreach_divforeach_div_listforeach_div_scalar_foreach_sqrt_foreach_rsqrt_foreach_maximum_foreach_minimum_foreach_clamp_min_foreach_clamp_max_foreach_reciprocal_foreach_sign_foreach_copyforeach_copyr`  _foreach_add__foreach_mul__foreach_div__foreach_copy_rc  add_bitwise_and_bitwise_left_shift_bitwise_not_bitwise_or_bitwise_right_shift_bitwise_xor_mul_div_Tensor_modelogical_and_logical_not_logical_or_logical_xor_sub_relu_sigmoid___and__
__lshift____or__
__rshift____xor____iand____ilshift____ior____irshift____ixor__re  r  r   rm  rp  r$  methodfuncrs  ru  _inductor_testr  rw  r  r{  set_source_Tensorr  fsdpr  r  *torch._higher_order_ops.auto_functionalizer  r  r	  r   r  while_loop_stack_outputrI  r  r  .torch._inductor.fx_passes.control_dependenciesr  r  invoke_quantr  r  r  r  r	  comm_loweringr	  r	  r*	  r/	  cvt_e8m0_rceilr7	  r  r8	  register_quantized_opsregister_woq_mm_opsr9	  register_onednn_fusion_opsr:	  register_jagged_opscontextmanagerr=	  r   ri   rg   <module>rh
     sJQ   " " " " " " "                    				   # # # # # # D D D D D D D D D D D D P P P P P P P P P P P P P P P P P P ' ' ' ' ' '         + + + +  $ $ $ $ $ $ $ $ $ ( ( ( ( ( ( H H H H H H U U U U U U ? ? ? ? ? ? : : : : : :                          M L L L L L L L         
 0 / / / / /              - , , , , , 8 8 8 8 8 8 8 8 8 8 8 8 = = = = = = = =                                   $                                      +****** WT]]Yt__ !j    g!!FH	 H H H H       .Juz,-//	y~
)
	9
5:#89;; /j./
)* 
 8j!6799 NP  P P P Py5       0   B B B BF F F F
   &3 3 3 3   !$$-&  , {z{{{}}}
 .   " " " "# # # # "'C C C C C C(         (O O O Od   42 2 2 2n  	(/    ,# # #42 2 2 2n  $V V V Vr$ $ $ $N  @E E E E E 5*7TRRR) ) SR)X 5-4HHH) ) ) IH) @E 5 5 5 5 5 5( 49?===& & & >=& ;@e K K K K K K 5#>>>F F F F ?>F 
7? $* * * *Z $ #7D    4:4HHH   IH>  >         4:DIII  JI0 4)UPTUUU  VU. DJT\49emTUU  VU 74 ,&do&&s+++ 4<T:::B B B ;:B, 4$$???" " " @?" DM?##   $# 4:C C C 4:C C C 49! ! ! 4:! ! ! 4:%&&% % '&% 4:! ! ! 4;D999> > :9>4 5)tDDD
 
 ED
 4>t<<<# # =<# 4;1 1  1h 4$$???49$7774<T:::1 1 1 ;: 87 @?1 4<T:::> > ;:> 4:4888uE uE uE 98uEp 4?===J J J >=J8 4#>>>   ?> 4'TBBB   CB
B B B BJ '<RVWWW/ / / XW/d  $ 4%)**$ $ +*$ 40455$ $ 65$ /T   (,/ / / / / /d ,4$  ! ! ! !H .6D   (,! ! ! ! ! !H ,3  * * * *Z .54   (,, , , , , ,^ 48y: y: y: y:x 4=d;;;3E 3E 3E 3E <;3El 4%4@@@6 6 6 6 A@6 4(dCCC    DC 4;D999=C =C :9=C@ 4:4888   984 4(dCCC      DC  4;D999   :9 4;D999D D :9D8 4>t<<<  =< 4?===  >=    48      $       D
3 
3 
3 
3 
3%6 %6 %6 %6 %6P7N 7N 7N 7Nt	, 	, 	, 59%1tLLL#+ #+ ML#+L 4&DAAAG G BAG 4?===  >= 4>#>>>' ' ?>' 4<    !  QB B B/ / /
 )():;; **49+>?? ))$**<== ++DJ,@AA  dl    ei,4 5 5 5 ei*2 3 3 3 49J J J 4:J J J >4$OOOL L POL
 >&DAAAK K K BAK >'TBBBJ J CBJ
 >-4HHH	 	 IH	 >(dCCCRS      DC6 >)tDDDLM     ED8    . . . . 4$+FFF
 "&^ ^ ^ ^ ^ GF^B N/N/X   7 7 7 7 7 7t               $F F FT d' ( ( ( d& ' ' ' d= > > > dl ' ' ' ' d&U 3 3 3 3 d!#5 6 6 6 d,5 9 9 9 9 d" / / / /9 M%E    	: MU   
 d! " " " d" # # # do    d!) * * * d+3 4 4 4 dj    dn# $ $ $ d)1 2 2 2 d,4 5 5 5 dk    d$5 1 1 1 1 d . . . d')@ A A A do} - - - d&(: ; ; ; do} - - - d&(: ; ; ; d!#5 6 6 6 d.0B C C C d* + + + d= > > > d= > > > d2 3 3 3 d= > > > d3 4 4 4 d0- @ @ @ d0 1 1 1 d/ 0 0 0 d/ 0 0 0 d1 2 2 2 d1 2 2 2 d- . . . d- . . . d- . . . d.0B C C C d0 1 1 1 d+ , , , d"$6 7 7 7 di    di    dm    di    di    dk    dn    dm    dl    do    d    d- . . . d  ! ! ! d' ( ( ( d# $ $ $ dn    d& ' ' ' d" # # # d$ % % % dn    d" # # # d# $ $ $ d* + + + d    dn    dj    d' ( ( ( d/ 0 0 0 d       d# $ $ $ d% & & & d# $ $ $ d! " " " dj    dm    dl" # # # dhU + + + + d(- 8 8 8 el$; < < < el$7 8 8 8 el$A B B B
 d! " " " d* + + + d"$6 7 7 7 dj    d' ( ( ( d; < < < dn    do    dj    08	   
 9A	   
 ,4	   
 ,6	    5=	   
 ,4	   
 5=	   
 4<	   
 =E	   
 9A	   
 BJ	   
 d+3_ E E E d+5 6 6 6 d,4o F F F d/7 I I I d08/ J J J d       d$+T B B B B d2:<N O O O
 49$777   87 4:"       * 74"## 3+d*++E222 5:  ( 4&DAAA      BA F 4%4@@@A A A A@AH   EL$"4566tDU 5 5 5 5 765p 5?##4 4 4 $#4 5#$$+ + %$+ 4+,,)" )" -,)"X 4&''  (' 4/00%)59$    10  <K K K  8 EK,-- 
    .-$  .> > > 0t//0B0B50I0IJJ
11!4455	 2 21 5 566
    4>"" $T$4     #" 4%&&Td    '&@ 4)**"4    +*" 5%-..A A /.A EJ	*++: : ,+:
 4;D999! ! ! :9!H 4>t<<<   =<B# # #F F F F FR  "1 "1 "1 "1J 4:4888
 
 98
 4%4@@@/ / A@/ 4>t<<<   =< 4)**   +*; ; ;  . 4?===   >= >4$OOO   POt t t tn !1 0%5! ! !  0@/?,4%0 0 0 ,
 4,$GGG  HG, 4;QUVVV	P 	P WV	P 1 1 1 4*EEE   FE 4<T:::9 9 9 ;:9 !     D 4=d;;;DH : : : : : <;:" 4#>>>3 3 3 ?>3 4$$???6 6 6 @?6 4&DAAAP P P BAP 4'TBBBPT q q q q q CBqp , , , , ,^ 4*233> > > > 43> 419::J J J J ;:J 4*233RVI I I I 43I 419::RVU U U U ;:U 4*233 !% $ $S S S S 43S 419:: !% $ $	 	 	 	 ;:	7 7 7 59$%%  &%*\ \ \ \~ 4'TBBB4 4 4 CB4n          :>   @ MQ     <   JN@ @ @ @ @:H H HV 5:PTUUU 5 5 5 VU5D   > 	1t    $  < 4/TJJJ 
 
 
 KJ
 4/TJJJ 
 
 
 KJ
 -=,<)1- - - ) 48dSSSN N TSNb   8B B B& & &R2 2 2j  0/%5      
 4,--B B .-BJ  0/$%      
 4+,,N N -,Nb) ) )X 4-..V V /.V 4-..V V /.VA A AH 43;<<BF. . . =<.b 4?===    >=* 4?===    >=, T_,%HHHT_,%HHHT_,%HHH A# A# A#H  0/$%      
 4+FFF b b b GFbJ  0/$%      
 4+FFF D D D GFDN  & HLO O O O Od    *   49:t : : : : :   4   ( ( (V4 4 4( DHei())T5     *) 4=!!u     "!      .-H    '&txERRR --H   
 48t,,,- - -,-`" " " "J 4:2 2 2 4:4888   98       48t,,,   -, DH:...( ( /.(   @ EI;$///$ $ 0/$. 	tx'7D  
  
 DIuz*d;;;$ $ <;$  DHei()). . . . . *). #"4;#677##DL$899 (():)BCC ""4;#677""4;#677 4;    . 4<     ! . 4$%%  &%* 4;D999   :9, 4;D999   :9, 49. . . . . . 48@ @ @ @
 486667 7 7 767 486667 7 7 767 !  %-    	!:!: ; ; ;*	**>>%+@+@AA*	**>>%+@+@AA.!!$+..N85;???  /!!$+..N85;???  H$L   ! !1uMMM 49#>>>!re $2 $2 $2 $2 ?>$2N 49$$???H H H @?H       	#"4:..&&tx00!!$),,""4:..$)$$
*
*4<
8
8''	22		DK	(	(t444 4<4000-. F F F F 10FRL L L L2  4079P Q Q Q !   * * *     * * *""  !122''(?@@   -    00
(()ABB   !122  4; ' ' '  **  *I*V      4: & & &  48 $ $ $  49 % % %     * * *  *	   ! *	    O*	  
 ! *	   
T\
*
*

T\
*
* !  $. ! !' * * * !  $. ! !' * * *""""''88
  4> " " "$)LLL  49     4<uz B B B B !  $. ! !# & & &  47%* = = = =  47%* = = = =  47%* = = = =uzBBB  47%* = = = =  47%* = = = =  49 % % %  49 % % %  49 % % %  4: & & &  49 % % %  4: & & &  4: & & &  49 % % %  4: & & &  4= ) ) )  49 % % %  4; ' ' '  4: & & &  4: & & &  49 % % %  4> * * * D D D D D D D DE E E$ % 
 
D4L4Ld5 5 
 
0 	 3+		
 	
 	
 	
 	
 5M4Lt5 5 
 
0 	 3+		
 	
 	
 	
 	

 .-CT    0/ct     4,3Sd K K K K--d.?.DcJJ   4,3S 9 9 9//0A0H#NN   4,13 7 7 7  4,3S 9 9 9  4,4c : : :  4,4c : : :  4,3S 9 9 9  4,13 7 7 7  4,<c B B B--d.?.DcJJ   4,3S 9 9 9//0A0H#NN   4-t 4 4 4  4. 6 6 6  405w ? ? ?  407 A A A  405w ? ? ?  407 A A A  427 A A A  4297 C C C  427 A A A  4297 C C C  43Z @ @ @  4-t 4 4 4))$*<dCC, , ,  T.35E    t079K    T.35E    t079K    T.35E    t079K    !3!;\  
    C        "K 0 0 0  )+= > > >  "K 0 0 0  !: . . .  *,? @ @ @  "K 0 0 0  C        !3 ' ' '  & 1 1 1  "K 0 0 0  "K 0 0 0  !: . . .  "K 0 0 0  C        T " " "   ( ( (    $,   , , , "  $/ " "#5 6 6 6   $+  z * * * "  $/ " "#6 7 7 7   $,   , , ,   - - -  !4? 3 3 3  t{ + + +  !4? 3 3 3   - - - 4+,,   -, 4=$%%  &% 4?&''  (' 4>""  #" (M')) 8 8LFD1((0011$7777 5=!!  "! 4<  7 7 ! 7 59+344  54
 59%;<<  =< 59>&455K K 65K 759>7## 
#uy~+344# # 54# 59>())%) = = = = *)=@ K J J J J J ! " " " 122S S 32S( 59).DIII/ / / JI/ 59)4$OOOE E E POE&  	I2  )JT2224 4 4
 59)9tTTT/ / / UT/
       H H G G G G G <T:::+ + ;:+\ 5*7TRRRDH      SR8 &DAAA      BA F 59?/788  98 59?.677  87 59)6DQQQc  c  RQc L P O O O O O O O           >8dSSS1 1 TS1hW W W
 >0dKKK") ") LK")L            ! ! ! ! ! ! +  * , , , '  ' ) ) )       ,  + - - -       %  $ & & &      ri   