
    J-ja                   *   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 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-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j                  e{      Z|i Z}d e~d!<   i Zd"e~d#<   i Zd$e~d%<    eEej                  j                            Zej                  j                  Zej                  j
                  Zej                  j                  Z eEej                  j                            Z eEej                  j                     ej                  j                  g      Z eEej                  j                            Zi Zd&e~d'<   ej                  j                  Zdd(Z	 	 	 	 dd)Zdd*Z	 	 	 	 dd+Zdd,Z	 	 	 	 dd-Z	 	 	 	 	 	 dd.Z eej*                  ej,                  ej.                  ej0                  ej2                  ej4                  ej6                  ej8                  ej:                  ej<                  ej>                  ej@                  ejB                  ejD                  g       ejF                  ejH                  ejJ                  ejL                  ejN                  ejP                  ejR                  ejT                  ejV                  ejX                  ejV                  ejZ                  ej\                  d/Zdd0Zdd1Zdd2Zd3d4	 	 	 	 	 	 	 dd5Zd6 Z	 	 	 	 	 	 dd7Zdd8Z	 	 	 	 	 	 	 	 	 	 	 	 dd9Z	 	 	 	 	 	 dd:Z	 	 	 	 	 	 	 	 	 	 dd;Zd3e7jt                  d3e}f	 	 	 dd<Zd= Zdd>Zd? Z	 	 	 	 	 	 dd@ZddAZdB Z	 d	 	 	 	 	 	 	 ddDZ eej                  j                  dE      dF        Z eej                  dE      ddG       Zd3dHddIZ eej                  j                  dE      ddJ       Zd3d3dKddLZ eej                  dE      dddM       ZddCe7jt                  d3ddd3d3df	dNZ endOe7j                  dP        eej                  dCdQ      ddR       ZdS Z eЫ        	 	 ddTZ eej                  d3dQ      dU        Z eej                  d3dQ      dV        Z eej                  ej                  ej                  ej                  ej                  g      dW        Z eedX      r  eej                        e٫        eej                  dE      ddY       Z eej                  dE      ddZ       Z eej                  g      dd[       Z eej                        d\        Z eej                        d]        Z eej                        d^        Z eej                        d_        Z eej                  j                        d`        Z eej                        da        Z eej                  dE      db        Z eej                  dE      dc        Z eej                  dE      dd        Z eej                        de        Z eej                  dE       eej                  dE       eej                  dE      ddf                     Z eej                  dE      dg        Z eej                  dE      d d e
j                  ddCfdh       Z eej*                  dE      ddi       Z eej                  dE      ddj       Z eej,                  dE      ddk       ZddlZ eej                  dE      	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddm       Zdn Z eej                  j                        do        Z eej                  j                        dp        Z eej                  dE      ddq	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddr       Z eej                  j                  dE      	 	 	 	 	 	 	 	 	 	 	 	 	 	 dds       Z eej                  j                  dE      ddq	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddt       Z eej                  j                  dE      	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddu       Z  eej                  j                  dE      ddq	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddv       Z eej                        ddw       Z eej                  dE      dddx       Z eej                  dE      dddy       Z eej
                  dE      dddz       Z eej                  dE      d{        Z eej                  dE      dd|       Z eej                  dE      dd}       Z eej                  dE      dd~       Z	 eej                  dE      d        Z
 eej                  dE      d        Z eej                  dE      d        ZddZ eej                        dd       ZddZej                   d        ZdddZdddZdddZ	 	 	 	 ddZd Z eej                  j.                  j0                  dE      d        Z eej2                  dE      d        Z eej4                  dE      d        Z eej6                  j8                  dE      d        Z eej<                        d        Z ej>                  d      d        Z d Z! eejD                  j                        Z# eejD                  jH                        Z% eejL                  j                        Z' eejL                  jH                        Z( eejR                          eejT                  dC        eejV                  dC        eejX                  dC        eej                  jZ                  j\                  j                          eej                  jZ                  j^                  j                          eej                  jZ                  j`                  j                          eej                  jZ                  jb                  j                          eejD                        d        Z" eejL                        d        Z& eePjd                  dE      d        Z3 eePjh                  dE      dd       Z5 eePjl                  dE      d        Z7 eePjp                  dE      d        Z9ddZ: eePjv                  dE      d ejR                  d	 	 	 	 	 	 	 	 	 dd       Z< eePjz                          eePj|                          eePjR                  dE      d d	 	 	 	 	 	 	 	 	 dd       Z?ddZ@ddZA eej                  j                  dE      d3d3ddd	 	 	 	 	 	 	 	 	 	 	 	 	 dd       ZB eej                  j                  e7j                  E      d3d3d	 	 	 	 	 	 	 dd       ZDd ZFd ZGd ZHd ZId ZJd ZKd ZLd ZM eej2                  j                  eM       d ZN eej                          eej                          eej                          eej                  d3        eej                  j                  d3        eej                  eH        eej                  d3        eej                  d3       ej                  j                         r eej                  d3       ej                  j                         r eej                  d3        eej                          eej                          eej                          eej                  j                          eej                  j                          eej                          eej                  j                          eej                  j                          eej                  j                          eej                          eej                  d3        eej                  eG        eej6                  eM        eej                  eG        eej                  eH        eej                  eG        eej                  eH        eej                  eH        eej                  eH        eej                          eej                          eej                          eej                          eej                          eej                          eej                  eG        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                  eH        eej                          eej                          eej                   eH        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j0                          eej2                          eej4                          eej6                          eej8                          eej:                          eej<                          eej>                  j                          eej@                  j                  d3        eejB                  eG        eejD                  jF                  jH                          eejD                  jF                  jJ                          eejD                  jF                  jL                          eejD                  jF                  jN                          eejP                          eejR                          eejT                  eH        eejV                          eejX                          eejZ                          eej\                          eej^                          eej`                  jb                          eejd                  j                  eNd3        eejf                  j                  eNd3        eejh                  j                  eNd3        eejh                  jj                  d3        eejl                  j                  eNd3        eejn                  j                  eNd3        eejp                  j                  eNd3        eejr                  j                  eNd3        eejt                  j                  eNd3        eejv                  j                  eNd3        eejx                  j                  eNd3        eejz                  j                  eN        eejz                  jj                          eej|                  j                  eN        eej~                  j                  eN        eej                  j                  eN        eej                          eej                  j                  dC        eej                  j                  eH        eej                  dE      dd       Z eej                        ddd       Zd Z eed      r  eej                        eū        eej                        d        Z eej                  j                  dE      	 dddddd3dd       Z eej                  dE      dd       Z eej                  dE      dd       Zd Z eej                  ej                  ej                  g      dddd3dd       Z eej                        dd       Z eej                        d        Z eej                        d        Z eej                        d        Z eej                        	 ddddd       Zd ZdÄ ZdĄ Z eej                  ej                  g      dddddddŜdƄ       ZdǄ ZdȄ Z  eej                         eڐe٫            Z e ed            Z e ed             ZdɄ 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E      ddЄ       Z eej                  dE      ddф       Zd҄ Z	 ddӄZdԄ ZddՄZ eej                  dE      dք        Z eej                  dE      dׄ        Z eej                  dE      dd؄       Z eej                        ddل       Zdڄ Zdۄ Z eej                  dE      dd܄       Z eePj                  dE      dd݄       ZddބZ eej                  j                  d3߫      Z eej                  j                  d3߫      Z eej                  dE      d        Z eej                  dE      d        Zed        Z eej                  dE      dd       Z eej                  dE      dd       ZddCd	 	 	 	 	 	 	 ddZ eej                  dE      dddd       Z eej                  dE      dd       Z eej                  dE      dd       Z eej                  dE      dd       Z eej                   dE      dCddd       Z 	 	 d	 	 	 	 	 ddZ eej                  j                        ddd       Z eej                  j                        ddd       Z eej@                  j                        	 d	 	 	 dd       Z eejB                  j                        	 d	 	 	 dd       Z eej                  j                        	 	 	 d	 	 	 	 	 dd       Z eej
                  j                        	 	 	 d	 	 	 	 	 dd       Zd Z eej                  j                        d        Z	 	 	 	 	 	 	 	 ddZ	 	 	 	 	 	 	 	 ddZ	 eej                  dE      dd       Z
ddZddZddZ	 ddZdddZd  ZdddZd Z eej&                  dE      	 dd       Z	 	 	 	 	 	 	 	 	 	 dΐdZ eej*                  dE      d        Zd Z eej8                  dE      	 	 	 	 dϐd       Z eej:                  dE      	 	 	 	 dϐd	       Z eej<                  j                  d3߫      Z eej<                  dE      d
        ZdАdZd Zd Zd Z eej8                  j                  d3߫      Z eej8                        d        Z eej<                  j                  d3߫      Z eej<                        d        Zd Z  eejB                        d        Z! eejD                        d        Z"d Z# eejH                  j                        	 dѐd       Z$ eej.                  dE      	 	 	 	 	 dҐd       Z eejJ                  dE      	 	 	 	 	 dҐd       Z% eejL                  j                  d3߫       eej.                  j                  d3߫       eejJ                  j                  d3߫      gZ'd Z( eej0                  j                  d3߫      Z) eej0                  dE      	 dd       Z eejT                  j                  d3߫      Z+ eejT                  dE      	 dd       Z*d Z,dddZ-ddӐdZ.d Z/ eej`                        ddd d!       Z0d" Z1d# Z2d$ Z3d% Z4 eejj                  ejj                  g      ddd3d&d'       Z6 eejn                        ddd3d&d(       Z7d) Z8ed*        Z9 eejt                  jv                  d3߫      Z< eejt                  jz                  d3߫      Z> eejt                  j~                  d3߫      Z@ eejt                  dC+      d,        Z:dd-ZA eej                        d.        ZB eej                  dE      dd/       ZCed0        ZDed1        ZEed2        ZF eej                  dC+      dd3       ZH eej                  gdC+      d4        ZIdՐd5ZJ eej                  gdC+      d6        ZK eej                  ej                  j                  gdCe7j                  Q      d7        ZG eej                  ej                  gdC+      d8        ZM eej                  ej                  g      ddd d9       ZO eej                  j                        ZQ eej                  j                        ZS eej                  j                        ZU eej                  j                        ZW eej                  j                        ZY eej                        dd:       ZP eej                        dd;       ZR eej                        d<        ZT eej                  dE      dd=       ZV eej                  dE      dd>       ZX eej                        ddd d?       ZZ eej                        dԐd@       Z\ eej                  dE      dԐdA       Z^ eej                  dE      dԐdB       Z`  eej                         e.dC               eej                         e.dD            Zc  eej                         e.dE            Ze  eej                         e.dFejN                  G            Zg  eej                         e.dHejN                  G            Zi eej                  dCdCdIJ      Zj eej                  j                  d3߫      Zm eej                  j                  dE      ddd3dKdL       Zn eej                  j                  dE      d֐dM       Zk eej                  j                  d3߫      Zp eej                  j                  d3߫      Zr eej                  j                  d3߫      Zt eej                  j                  d3߫      Zv eej                  j                  d3߫      Zx eej                  j                  dE      dN        Zy eej                  j                  dE      ddO       Zz eej                  j                  dE      d֐dP       Z{ eej                  j                  dE      dאdQ       Zo eej                  j                  dE      d֐dR       ZqddSZ|dؐdTZ} e|ej                        Z~ e}ej                        Z e|ej                         Z e|ej                        Z eej                        Z e}ej                        Z e}ej                        Z eej
                        Z eej                  dCU      Z eej                  dC+      ddVdW       Z eej                  dC+      ddVdX       Z eej                  jz                  edCdYZ      Z eej                  jz                  edCdYZ      Z e}ej                          e}ej                          eej                        Z eej                         Z eej"                        Z eej$                  d[\      Z eej&                        Z eej(                        Z eej*                        Z e|ej,                          e|ej.                        Z  eej0                  e7j                  E      e        e|ej2                          e|ej4                          e|ej6                          e}ej                          eej8                  ddCejZ                  ]      Z eej:                  ddCejZ                  ]      Z eej<                  ddCejZ                  ]      Z eej>                  ddCejZ                  ]      Z eej@                        Z eejB                        Z  eejD                        e         eejF                        e        eejH                        Z eej                        Z e|ejJ                        Z eejL                          eejN                  d^\      Z eej                          eejP                  ejZ                  G         eejR                        e        eejT                  ejZ                  G        eejV                  ejZ                  G        eejX                  ejZ                  G        eejZ                  ejZ                  G      Z eej\                  ejZ                  G        eej^                  ejZ                  G        e|ej`                          e|ejb                          e|ejd                          e|ejf                          e|ejh                          e|ejj                          e|ejl                          e|ejn                          e|ejp                          e|ejr                          e|ejt                          e|ejv                          e|ejx                          e|ejz                          e|ej|                          e|ej~                         dd_lmZmZ d` ZÐeD ]S  Z eeeī      D ]  \  ZŐZƐZ e̐eŐeĐeƐeǐa         eeeī      D ]  \  ZŐZƐZ e̐eŐeĐeƐeǐa        U  eej                  j                  ejdCU      Z eej                  jz                  ejdCU      Z eej                  j                  ejdCU        eej                  j                  eI      Z eej                  j                  eI        eej                  jz                  eI      Z eej                  j                  e        eej                  jz                  e        eej                  j                  e        eej                  j                  e        eej                  jz                  e:        eej                  j                  e:        eej                  j                  e:        eej                  j                  eG      Z eej                  j                  eG        eej                  jz                  eG      Z eej                  e        eej                  e~        eej                  j                  e        eej                  jz                  e        eej                  j                  e        eej                  jz                  e        eej                  j                  e        eej                  jz                  e        eej                  j                  e        eej                  jz                  e        eej                  e        eej                  e        eej                  eū        eej                  eī      Zdb Z eej                  j                  ej                  j                  eʫ        eej                  jz                  ej                  jz                  e˫        eej                  j                  ej                  j                  eͫ        eej                  jz                  ej                  jz                  eΫ        eej                  j                  ej                  j                  eի        eej                  jz                  ej                  jz                  e֫        eej                  j                  ej                  j                  e        eej                  jz                  ej                  jz                  e        eej                  jz                  ej                  jz                  e       dc Z eej                  ej        eej                  e        eej                  e        eej                  e        eej                  e        eej                  e        eej                  e        eej                  eI        eej                  j                  eG        eej                  j                  eH        eej                  e        eej                  e        eej                  e        eej                  e        eej                  e        eej                  e        eej                  e         eej                        e         eej                        e         eej                        e         eej                        e         eej                        e        eej                   ej                          eej                  ej                          eej                  ej                          eej                  ej                          eej                  ej                          eej
                        ddd       Z eej                  j                        de        Z eej                  j                        df        Z eej                        dg        Z	 e>j                         D ]  \  ZZ  e e?e            e         eej                        dh        Z eej<                        di        Z eej                  j                  j                         dj        Z eej                  j$                  j&                        dk        Z eej                  j                  j(                  j*                        dl        Z eej                  j.                  dm      r; eej                  j.                  j                  j                        dn        Z eej                  j                  j2                        dddo       Zd dplmZ  ee        ee*      dq        Z eej                  j:                  j<                  dE      	 	 dِdr       Z eej                  j:                  j>                  dE      dds       Z  eej                  j:                  j@                  dE       ejB                  edCt              eej                  j:                  jD                  dE      dڐdu       Z" e#       Z$dېdvZ%d dwl&m'Z'  ee'dE      dx        Z( eej                  jR                  dE      ddydܐdz       Z* ee(dE      	 	 	 	 dݐd{       Z+ eej                  j                  jX                  j                        d|        Z, eej                  j                  jZ                  j                        d}        Z- eej                  j:                  j\                  dE      d~        Z.ddl/m0Z0m1Z1  e0         e1         eePjd                  dE      d        Z2d Z3 eePjh                  dE      d        Z5 eej                  jl                  dE      dCddd       Z7ddlNm8Z8  eMe8       ddlNm9Z9  e9jt                           e9jv                          ddlNm<Z<  e<jz                          ddlNm>Z>  e>j~                          ej                  dސd       ZAy(      )annotationsN)defaultdict)Callable
CollectionIterableSequence)AnycastTYPE_CHECKING	TypeGuardTypeVar)	ParamSpec)patch)counters)associative_scan_op)triton_kernel_wrapper_mutation)FakeScriptObject)is_opaque_value)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_add2dict[Callable[..., Any] | str, Callable[..., Any]]	loweringsz/dict[torch._ops.OpOverload, Callable[..., Any]]user_loweringsz6dict[torch._ops.OpOverload, Callable[..., Any] | None]_maybe_layout_constraintsz2dict[torch._ops.OpOverload, torch._ops.OpOverload]inplaceable_foreach_opsc                     t         j                  j                  j                  D ]7  } | j                  D ]&  }|j                  dk(  r|j
                  t        v r%  y 9 y)Ncall_functionTF)rR   graphcurrent_nodeusersoptargetforeach_ops)nodeusers     f/Users/jameslopez/projects/TradingBot25/.venv/lib/python3.12/site-packages/torch/_inductor/lowering.pycur_node_has_non_foreach_usersrf      sI    $$**JJDGG.DKK;4N  +
     c                f   t        t              }d}t        |       D ]  \  }}t        |t              sd}|f}t        |  xs t        j                  }d }|D ].  }t        |t              s|j                  j                         } n |J d       |r|\  }|||f   j                  ||f        |S )NFTz.foreach op should have at least one tensor arg)r   list	enumerate
isinstancer   rG   r/   #combo_kernel_foreach_dynamic_shapesrB   data
get_deviceappend)	arg_pairsoutunpack_argsiargsuse_foreachdevicets           re   group_foreach_argsrx      s     d
CKY'4$)K7DD!!OV%O%O 	 A!Y'**,  !S#SS!GTV[!"))1d)4 (  Jrg   c                    t        | t        j                  j                        syt	        | d      x}rt        |      S | t        v r	t        |    S y)zHGet layout constraints. Returns None if there are no layout constraints.NF)with_default)rk   torch_ops
OpOverloadr   tag_to_layout_constraintrY   )fnmaybe_layout_tags     re   maybe_layout_constraintsr      sO    b%**//04ReLLL'(899	&&(,,rg   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y 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)tags    re   r~   r~      s     ehhll...((
ehhll333))
ehhll333&&
ehhll***
:3%@
AArg   c                $    | st        d|       y )Nzinductor does not support NotImplementedErrorcondmsgs     re   
assert_nyir      s    !$>se"DEE rg   c                    t         t        t        t        t        f      r D cg c]  }t        |       c}S t         t        j                  j                        rt        j                          y t         t        j                  j                        r,t        j                   fd j                         D               y c c}w )Nc              3  6   K   | ]  }t        |        y wN)getattr).0overloadr   s     re   	<genexpr>z,add_needs_realized_inputs.<locals>.<genexpr>   s      %
2@hGB!.   )rk   ri   settupler&   add_needs_realized_inputsr{   r|   r}   needs_realized_inputsaddOpOverloadPacketupdate	overloads)r   xs   ` re   r   r      s    
 "tS%45689b)!,b99"ejj++,!!"%
 	 
B

33	4$$ %
24,,.%
 	
  :s   Cc                    t        | t        j                  j                        r)| j	                         D ]  }|t
        t        | |      <    y |t
        | <   y r   )rk   r{   r|   r   r   rY   r   )r   
constraintr   s      re   add_layout_constraintr      sG     "ejj112H?I%gb(&;< ' )3!"%rg   )r   r.   r,                     	   
         c                b    t        | t              s| S | t        v sJ d|  d       t        |    } | S )Nzid z missing from DTYPE_ID_LOOKUP)rk   intDTYPE_ID_LOOKUPdtypes    re   decode_dtyper     s=    eS!O#Os5'1N%OO#E"ELrg   c                    t        | t              r4t        | j                               xs t	        | j                               S t        | t
        j                        r| j                  du S t        | t              S NT)	rk   rB   r   	get_dtyper   sympyExpr
is_integerr   r   s    re   is_integer_typer     sX    !Y.Q2B1;;=2QQ	Auzz	"||t##!S!!rg   c                t    t        | t              rt        | j                               S t        | t              S r   )rk   rB   r   r   boolr   s    re   is_boolean_typer   &  s*    !Y..!T""rg   F)return_compute_dtypec                f    dd}|D cg c]
  } ||       }}t        |d| i\  }}|r|S |S c c}w )Nc                    t        | t        t        j                  f      r| S t	        | j                               }t        j                  dg|z  | j                               S )Nr.   r   )	rk   r    r   Basiclenget_sizer{   zerosr   )inpdims     re   construct_inputz+get_promoted_dtype.<locals>.construct_input2  sG    cFEKK01Jclln%C;;sSy@@rg   type_promotion_kind)r   r	   returnr	   )r   )r   r   rt   r   arginpscompute_dtyperesult_dtypes           re   get_promoted_dtyper   -  sQ    
A -11DSOC DD1"4	##6#M< 1=BlB	 2s   .c                0   t        | t        t        f      s| g} nt        |       } t        |       D ]b  }t        |t        j                  j
                        s(|j                         D ](  }t        ||      }|t        vs| j                  |       * d | S r   )
rk   ri   r   r{   r|   r   r   r   rW   ro   )aten_fnr   r   other_fns       re   get_overloadsr   A  sy    ge}-)w-7mb%**556LLN"2x09,NN8, +  Nrg   c                    t        | t        j                  j                        r|| j                  v S t        | t        j                  j
                        r|| j                         v S yNF)rk   r{   r|   r   _qualified_op_namer}   name)r`   	namespaces     re   in_namespacer   Q  sQ     "ejj112B1111	B

--	.BGGI%%rg   c           	         t        | j                  t        j                        rt	        | j                               r| S t        j                  j                  j                  | j                               }| j                         }|u|j                  dk(  rf||k7  rat        |      dk(  st        |      dk(  rE|d   dk(  r=t        t        j                  t        j                  j!                  | |d                  S | S )zB
    Copy cpu scalar if doesn't not match with given `device`
    cpur   r.   F)rk   rm   r1   ReinterpretViewr$   r   rR   r]   sizevarsguarding_hints_or_throwrn   typer   rB   
StorageBox
DeviceCopycreate)r   rv   size
cur_devices       re   maybe_copy_cpu_scalarr   [  s     affb0015N	

6 7733AJJLADJOOu$& Y!^D	Q47a<r}}';';Az5'QRSSHrg   c           	         t               D cg c]  \  }}t        |t              s| }}}j                         D 	cg c]  \  }}	t        |	t              s| }
}}	|s|
s fS |s|r|rt        j
                  nl D cg c]0  }t        |t        t        j                  f      st        |d      r|2 }}|j                  d j                         D               t        |d|i|r |d      n|
d      j                         |D ]  }t         |          |<    |
D ]  }t        |         |<    d	fd} D cg c]
  } ||       c} j                         D 	ci c]  \  }}	| ||	       c}	}|r1t        t!        t#        j$                   fd|D        fd|
D                     }t!        |d   j'                               }t)        ||dt+        |             D ]
  \  }}| |<    t)        |
|t+        |      d       D ]
  \  }}||<    t-        t+                     D ]<  }t         |   t.        j0                        s!t3        j4                   |   |       |<   > D ]<  }t        |   t.        j0                        s!t3        j4                  |   |      |<   >  fS c c}}w c c}	}w c c}w c c}w c c}	}w )
zB
    Transforms arguments for broadcasting and type promotion
    r   c              3  :   K   | ]  }t        |d       s|  yw)r   N)hasattr)r   as     re   r   z!transform_args.<locals>.<genexpr>  s     !T_7@S!_s   r   r   c                    t        | t              rt        |       S t        | t        j                        r"t        j                  | j
                        S | S )Nvaluer   rv   )rk   rB   to_dtyper1   Constantr   )r   rv   r   s    re   promoteztransform_args.<locals>.promote  sD    #y)U++C-{{%OO
rg   c              3  (   K   | ]	  }|     y wr    )r   rs   rt   s     re   r   z!transform_args.<locals>.<genexpr>  s     3lT!Wl   c              3  (   K   | ]	  }|     y wr   r   )r   kkwargss     re   r   z!transform_args.<locals>.<genexpr>  s     71VAYr   N)r   r	   r   r	   )rj   rk   rB   itemsr{   r   r    r   r   r   extendvaluesr   rn   r   broadcast_tensorsri   	itertoolschainr   zipr   ranger1   r   r7   r   )rt   r   	broadcastr   convert_input_to_boolrs   r   args_indicesr   vkwargs_indicesr   promoting_argsr   broadcastedr   rv   r   s   ``              @@re   transform_argsr
  o  s    #,D/N/$!QZ95MA/LN$*LLNONDAqjI6NaNNOV|3 JJE
 Aa&%++!6771g;N    !!!TV]]_!TT&$7E &2Da!vnQ>O7P
*, 	 A+DGV<DG   A-fQi@F1I  	 %))Dq
D),2LLN;NDAq!WQZ-N;'3l377
 KN++-.k2EC4E&FGDAqDG HC4E4G(HIDAqF1I J s4y!A$q'2;;/$++DGT:Q " A&)R[[1&--fQi>q	  <I OO> *;s(   KKK	K	
5KK4Kc                    t        j                        dfd       }t        |       }t        j	                  |       t
        j	                  t        j                  ||             |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
    c                 .     | i |}t        |       |S r   )rC   )rt   r   rq   	decomp_fns      re   wrappedz+_register_foreach_lowering.<locals>.wrapped  s    ((C
rg   )rt   r	   r   r	   r   r	   )	functoolswrapsr   rb   r   rW   dictfromkeys)r   r  r  aten_fnss    `  re   _register_foreach_loweringr    sX     __Y  
 W%Hx T]]8W56Nrg   c                     t        j                         fd       }t                |j                  t        j                   |             |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                 P   t        |       } t        |      }d}t        |       dk(  r)t        | 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 xs t        |d         yw)_c10d_functionalN)	fallbacksr   )r   r   s     re   r   z6_register_lowering.<locals>.wrapped.<locals>.<genexpr>  s)      
PW"R9_DR1C DDPWs    "c              3  &   K   | ]	  }|d k(    yw)rq   Nr   r   r   s     re   r   z6_register_lowering.<locals>.wrapped.<locals>.<genexpr>  s     6v!1:v   zout= ops aren't yet supported)	ri   r  r   rk   r   allanyr
  rC   )	rt   r   unpackedrq   r   r  r  r  r   s	       re   r  z#_register_lowering.<locals>.wrapped  s    t*!%ft9>ja4-@HQ=D 
PW
 
 6v66W8WW6%&)%8:O
f 6D((C
rg   )r  r  r   r   r  r  )r   r  r  r   r  lowering_dictr  s   `````  re   _register_loweringr!    sJ    & __Y  6 G$Gw89Nrg   c                @    t        j                  t        | ||||      S )z+
    Shim to support decorator syntax.
    )r  r   r  r   )r  partialr!  )r   r  r   r  r   s        re   register_loweringr$    s)     /3# rg   c                   t        |      }| r| |k(  r|S g }t        j                  t        |       t        |      t        j
                  j                        D ]  \  }}t        j                  j                  j                  |      r|j                  |       Bt        j                  j                  j                  |      r|j                  |       }t        j                  j                  j                  ||       t        t	        j                  |      j                        t        t	        j                  |      j                        k  r|j                  |       |j                  |        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OnerR   r]   r   is_size_one_or_falsero   check_equalsr   expandfree_symbols)r   boutputr   ys        re   broadcast_symbolic_shapesr2  %  s    	aAQF%%hqk8A;%''++VV177003MM!WW2215MM!GG))!Q/5<<?//03u||A7S7S3TTa a  W &!""rg   c                  	 |	|J d       ||t         j                  }t        d | D              s| S t        d | D              r*|xs t	        | d|ifd}| D cg c]
  } ||       c}S t        d | D              }|j                         	|t        j                  k(  r(	t        j                  t        j                  fv r	fd}nd }g }| D ]  }t        |t        t        f      rg|j                  t        j                   t#        j$                   ||      	|j'                         	      t)        |j+                                            t        |t,        j.                        rW|j                  t        j                   t1        |	|j'                         
      t)        |j+                                            |j                  |        |S c c}w )NzEonly one of override_return_dtype or type_promotion_kind may be givenc              3  f   K   | ])  }t        |t        j                  t        t        f       + y wr   )rk   r   r   r   floatr  s     re   r   z$promote_constants.<locals>.<genexpr>G  s#     HAz!ekk367   /1c              3  f   K   | ])  }t        |t        t        t        j                  f       + y wr   )rk   r   r5  r   r   r  s     re   r   z$promote_constants.<locals>.<genexpr>I  s#     
DV:a#uekk23Vr6  r   c                    t        | t        j                        r!t        j                  | t        d             S t        j                  | t        d             S )Nindexr   rv   r   )rk   r   r   r1   r8   rF   r   r   r   s    re   
const_funcz%promote_constants.<locals>.const_funcP  sJ    !U[[)**5t1D  {{%d@STTrg   c              3  l   K   | ],  }t        |t        t        t        j                  f      s)| . y wr   )rk   rB   r7   r1   r   r  s     re   r   z$promote_constants.<locals>.<genexpr>Y  s%     WA:a)Z1U#Vas   *44c                N    t        j                  |       j                         S Nr   )r{   tensoritem)r  tensor_dtypes    re   <lambda>z#promote_constants.<locals>.<lambda>a  s    %,,q"E"J"J"Lrg   c                    | S r   r   )r  s    re   rC  z#promote_constants.<locals>.<lambda>c  s    !rg   r   r9  )r   DEFAULTr  r  r   nextr   r{   r   bfloat16float16rk   r   r5  ro   r7   r   r1   r   get_device_or_errorri   r   r   r   r8   )
inputsoverride_return_dtyper   r<  r   ex_round_scalarrq   r   rB  s
           @@re   promote_constantsrN  ?  s    (,?,G OG $)<)D=EEHHH

DV
DD% 
);*
 !4*
	U (..v!
1v..	WW	WB<<>L 

*|@ 0 M#
Ca#u&JJ!!KK+A.*!557
 '	 5;;'JJ!!$|B<R<R<T '	 JJqM- 0 JM /s    Gc                    t        | |t        j                        | j                         |j                         fd}t	        j
                  | j                         || j                               S )zBCompute a + alpha * b using FMA for CUDA floating-point precision.r   c                     |       } |       }t        t        j                        rt        j                        }nt        j
                        }t        j                  |||      S r   )rk   r   r   rQ   
index_exprconstantfma)idxa_valb_val
alpha_expra_loaderalphab_loaderr   s       re   inner_fnz%_add_with_alpha_fma.<locals>.inner_fn  sV    eU[[)u5JeU3Jwwuj%00rg   rv   r   r\  ranges)r   r   rE  make_loaderr?   r   rn   r   )r   r/  rZ  r\  rY  r[  r   s     ` @@@re   _add_with_alpha_fmar`    sf    		;CCE
 }}H}}H1 ||~zz|	 rg   c           	     2     ddd fd}|S )zcWraps a pointwise fn and returns a function representing the pointwise in
    the define-by-run IR.N)rZ  c           	       	
 t        d D              r	rJ   S t              r| | dk7  rrt        d   t              rnd   j	                         }d   j                         j                  r>t        j                  j                  s$|"|j                  dk(  rt        d   d   |       S t              t        d   |       d<   n| J D cg c]  }|j                          c}d   j                         xs d   j                         dd  D ]`  }t        |t         j"                        rt%              t%        |j                               k(  rDJ d d d|j                                  t        j&                  t        j(                  ft*        j,                  d uxr{ t/        t*        j,                  dd       d uxr\ t*        j,                  j0                  j2                  d uxr4 t*        j,                  j0                  j2                  j5                  d	d
      

xr v 		
f	d}sSd }D ]7  }t7        |j	                         j                        s'|j	                         } n |sd   j	                         }xs }t9        j:                  ||      S c c}w )Nc              3  V   K   | ]!  }t        |t              xr t        |       # y wr   rk   r9   r:   r   r   s     re   r   z0make_pointwise.<locals>.inner.<locals>.<genexpr>  s'      /
BH3JsF#6	#6&   ')r.   r   cudazndim mismatch  r^   low_precision_pointwise_barrierFc                
  	 t        |       t              k(  sJ d|  d        t        j                  k(  r D cg c]
  } ||        c} S g }t              D ]e  \  }} ||       }|   j	                         }	r2|v r.t        j                  ||d      }t        j                  ||      }|j                  |       g  
| }r.t        j                  |d      }t        j                  |      S |S c c}w )Nzwrong ndim ri  F)use_compute_types)r   r{   r   rj   r   rQ   r   ro   )r:  loadinputs_loaded	inp_indexrq   	inp_dtypedowncastr   emulate_output_castemulate_precision_castsr   rJ  loaders	low_pr_fpoverride_fn_when_input_boolr^  s          re   r\  z/make_pointwise.<locals>.inner.<locals>.inner_fn  s	   u:V,LE7!F8.LL,

"'B'N2W4UWTT%[W4UVV "'0'9OItu+C &y 1 ; ; =I.9	3I#&<<YRW#X!ll8Y?!((- (: -(&  #||C%PH<<%88
# 5Vs   D r]  )r  rN  rk   r9   rn   r   is_floating_pointr{   versionhipr   r`  ri   mulr_  r   r1   BaseConstantr   rG  rH  rR   r]   r   r^   metagetrH   r?   r   )rZ  rJ  
inp_devicer   otherr\  rv   rs   r   rr  rs  rt  ru  r^  allow_alphar   override_devicerv  rK  triton_fallbackuse_fma_for_alphas    `      @@@@@@re   innerzmake_pointwise.<locals>.inner  s   &3 /
BH/
 ,
 #"?"F++"6+@A UaZ %F1Iv)F!'!5!5!7Jq	++-?? % 1 1&2&OOv526!9fQiOO f U3r
= =,23Fq1==?F3##%%>)<)<)>ABZEeR__5V I : At1VHAenn.>-?@A    ^^U]]3	GG4 X6dBX$$))5X $$))--.OQVW	 	  6L%9:L	 	, F!,,.--.\\^F  --/ !*F	
 	
o 4s   $K)rJ  rB   r   )r   rK  r  rv  r  r  r  r  s   ``````` re   make_pointwiser    s     )- X
 X
t Lrg   c                $     dddd fd}|S )Nr.   )rZ  r   c                  	 t        |      }dk(  r,|r*t        |d   t         t        f      s|j                         	n
dk(  r|	n| 	t	        t
        j                  j                  j                        dk(  xs6 t
        j                  j                  j                  t        v xs
 t               }d }|D ]  }t        |t         t        f      s|} n |J d       g }|D ]H  }t        |t         t        f      s|j                  |gt	        |      z         8|j                  |       J t        t        |       }
	fd}t        |t	        |      ||      S )Nr   rh  r   z1at least one input must be a list to a foreach opc                &    r
 | i iS  |  S r   r   )rt   r  pw_fnscalar_kwarg
scalar_vals    re   apply_fnz7make_foreach_pointwise.<locals>.inner.<locals>.apply_fn2  s%    dA|Z&@AAd|#rg   )ri   rk   r   popr   rR   r]   r^   r_   ra   inplace_foreach_opsrf   ro   rx   r  foreach_group_loop)rZ  r   rJ  realize_outputsa_list_inputinputbroadcast_inputsgroupsr  r  r  r  r  s            @re   r  z%make_foreach_pointwise.<locals>.inner  sO   
 fG#vbzD%=9  JW$JJ $$**+q0 0ww##**.AA0-/ 	 E%$/$  ' 	
?	
'
 EedE]3 ''#l2C(CD !''.  $C)9$:;	$ "&#l*;XWWrg   )rJ  zlist[list[TensorBox]]r   )r  r  r  r  s   ``` re   make_foreach_pointwiser    s    45Q 2X 2Xh Lrg   c                   dg|z  }| j                         D ]  \  \  }}}g }|D ]v  \  }	}
 ||
      }|||	<   t        j                  j                  |t        j
                        sB|sE|sH|j                          |j                  |j                                x |st        j                  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 
 y wr   r   r  s     re   r   z%foreach_group_loop.<locals>.<genexpr>Y       .gq}g   )r   rR   r]   has_featureBackendFeatureFOREACHrealizero   get_operation_nameregister_operation_listr  )r  num_outputsr  r  outputsrv   ru   groupoperation_list
output_indrt   r0  s               re   r  r  =  s     f{"G(.$u$& %Jd^F"(GJ ##FN,B,BC# %%f&?&?&AB !& GG++N; )7" .g....Nrg   Tc                    | j                         k(  r|rt        |       S | S fd} t        |      |       S )Nc                    t        j                  |       }t        j                  t        j                  f}s|v rt        j                  |      }|S )N)	src_dtyperl  )rQ   r   r{   rG  rH  )r   resultru  r   r  rl  s      re   	_to_dtypezto_dtype.<locals>._to_dtyped  sR    /	
 ^^U]]3	 Ui%7 \\&%0Frg   rK  )r   cloner  )r   r   copyrl  r  r  s    ` ` @re   r   r   ]  sE     IEuQx&Q& B>)5A!DDrg   rP  c                   ddl m} |}| j                  } ||t        j                        }t        j
                  |      5   |j                  |  ddd       |j                  }|sJ t        |      }dgt        |      z  }	|j                         D ]  \  \  }
}}g }|D ]k  \  }}||	|<   t        j                  j                  |
t        j                        s:|s=|j                          |j                  |j!                                m |s~t        j                  j#                  |        t%        d |	D              sJ |	S # 1 sw Y   xY w)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 
 y wr   r   r  s     re   r   z_foreach_map.<locals>.<genexpr>  r  r  )subgraph_loweringr  graph_modulerR   r]   set_graph_handlerrungraph_outputsrx   r   r   r  r  r  r  ro   r  r  r  )subgraphrt   r   r  rJ  gmpw_subgraphsub_outputsr  r  rv   ru   r  r  r  r0  s                   re   _foreach_mapr  v  s.    =F			B+BAGGLK	
		[	)  
* ++K;,Ffs;''G(.$u$& 
"(GJww""6>+A+AB{ %%f&?&?&AB  GG++N; )7 .g....N3 
*	)s   E		Ec                   |j                   s| j                         j                   rk| j                         r/t        | |      }t        j
                  j                  ||        |S  t        t        j                  j                  d      | |      S | j                         }t        j                  t        j                  f}t        j                  xr
 ||v xs ||v  }t!        | |d|      S )Nr   Fadd_to_fallback_setT)r  rl  )
is_complexr   r   
empty_liker1   InplaceCopyFallbackr   fallback_handlerprimsconvert_element_typedefaultr{   rG  rH  r/   rs  r   )r   r   dstr  ru  rl  s         re   _convert_element_typer    s    1;;=33::< Qe,C""))#q1J#**22  I/I&& 	;)#9u	'9 Au4;LMMrg   r  c                  | j                         }||k(  r|rt        |       S | S d } ||      } ||      }||k7  r* t        t        j                  j
                        | |      S t        t        j                  | |            S )Nc                    | j                   rt        j                  |       j                  S | t        j                  k(  ryt        j
                  |       j                  S )Nr   )rw  r{   finfobitsr   iinfor   s    re   _get_primitive_bitwidthz1to_dtype_bitcast.<locals>._get_primitive_bitwidth  sD    "";;u%***ejj ;;u%***rg   )	r   r  r  atenviewr   rB   r6   r   )r   r   r  x_dtyper  src_bitsdst_bitss          re   to_dtype_bitcastr    s~    kkmG%uQx&Q&+ 'w/H&u-H80		0E::))!U344rg   c                &   |j                   s| j                         j                   r`t        j                  t        j
                  j                  t        j                  j                  j                  j                  | |            S t        | |      S r   )r  r   rB   r   r1   ComplexViewr{   rQ   r  r  r   r  r;  s     re   _view_dtyper    se    1;;=33NN!!%)).."5"5";";QF
 	
 Au%%rg   r  non_blockingc                   t        |      }| j                         |k(  r|rt        |       S | S t        j                  t
        j                  j	                  | ||            S r   )rF   rn   r  rB   r   r1   r   )r   rv   r  r  s       re   	to_devicer    sO    6"F||~uQx&Q&BMM00FLIJJrg   c                     t        | |d|      S )NTr  )r  )r   rv   r  s      re   _device_putr    s    QTEErg   c
                    |xs | j                   }t        |      }
t        |||       |t        |      }t        |
|||||	      }
 t	        | |||      |
      }
t        t        |      r" t	        t        t        |      d|      |
       |
S )z3A pointwise function that maps ops.{name} to inputsN)rK  rv  r  r  r  )r  r   r  )r   r  )__name__r=   rM   r  r$  r   r  r   )r   r   r  r   r  rK  rv  r  r  r  r   s              re   register_pointwiser    s     #7##D	T	B'!#8 #.&12M&N#	
3$?+'
B
	/3	

 	

B ud	
E4  $"7	
 		
 Irg   ldexp)r   rK  )r  r   c                Z  	 t        d      | j                         }|j                         }|j                  }|j                   xr |t        j                  k7  }|r|rfd} t        |      | |      S t        |       rt        j                  n|		fd} t        |	      | |      S )Nr  c                     | |      S r   r   )r   nldexp_fns     re   compute_ldexpz%ldexp_lowering.<locals>.compute_ldexp'  s    Aq>!rg   c                    t        j                  |      }t        j                  d      }t        j                  ||      }t        j                  | |      S )Ng       @)rQ   r   rS  powrz  )r   r  
n_out_typetwo
pow_result	out_dtypes        re   compute_fallbackz(ldexp_lowering.<locals>.compute_fallback/  sD    a3J,,sI.Cj1J771j))rg   r  )r=   r   rw  r{   r   r  r   float32)
r   r  r  n_dtype
x_is_floatn_is_intr  r  r  r  s
           @@re   ldexp_loweringr    s    7#HkkmGkkmG**J,,,FEJJ1FHh	" -~m,Q22 &5Q%7EMMW		*
~"+
 Q 	rg   c                 .   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   rt   r   r  s     re   frexp0zregister_frexp.<locals>.frexp0@      d%f%a((rg   c                      | i |d   S Nr.   r   r  s     re   frexp1zregister_frexp.<locals>.frexp1C  r  rg   r  c                 0     d   | i | d   | i |fS Nr   r.   r   )rt   r   pw_fnss     re   r   zregister_frexp.<locals>.fnK  s.    vay$)&)96!9d+Ef+EEErg   NrP  )
r=   r  r{   int32r$  r  r  r   r  r   )r   r  r  r   r  r  s       @@re   register_frexpr   ;  s    D E)) 	vvU[[AF
F
	




B ud	
E4  $	
 	 Irg   c                :    t        |||      }t        | |      }|S )Nr  r  )r  r  )r   pointwise_lowering_fnr  r  r   s        re   register_foreach_pointwiser  ]  s(     
 ;\
B 
$GR	0BIrg   c           
        d }t        |t        t        f      r t        |      |      }t        |t        t        f      r t        |      |      }| ||g}t	        |d   |d   t
        j                        }t        |      D cg c]  \  }}t        |t              s| }}}t        |t        |D cg c]  }||   	 c}       D ]
  \  }}|||<    t        t        |            D ]Y  }t        ||   t        j                        s!t        j                   ||   t#        ||d      j%                                     ||<   [  t'        ||      |d   t)        |d   |      t)        |d   |            S c c}}w c c}w )Nc                 &    t        j                  |  S r   )rQ   wherert   s    re   r   zwhere.<locals>.fnl  s    yy$rg   r.   r,   rP  r   r  )rk   r5  r   constant_liker   r   rE  rj   rB   r  r   r  r   r1   r   r7   r   ri   r   r  r   )	r   r   r/  r   rt   r   rs   r   indicess	            re   r  r  j  si     !eS\"M!Q!eS\"M!Q!Q<DQa.M.U.UE 't_I_TQ
1i0Hq_GIG.'0J'Qa'0JKL1Q M3t9d1gr{{+ ''Qd71:6F6O6O6Q1RSDG  ;>"E:Q$q'5)8DGU+C  J0Js   E9 E96E?c                    t        |       dk(  r&t        | d   t        t        f      rt	        | d    S | S t        j                  t        d | D        d      }g }| D ]q  }t        |j                               x}|k(  rn?t        |      t        |      k7  st        d t        ||      D              rt        ||      }|j                  |       s |S )Nr.   r   c              3  <   K   | ]  }|j                           y wr   )r   r  s     re   r   z$broadcast_tensors.<locals>.<genexpr>  s     #A&QAJJL&   r   c              3     K   | ]Z  \  }}t         j                  j                  j                  |      t         j                  j                  j                  |      k7   \ y wr   )rR   r]   r   r+  r   r   r/  s      re   r   z$broadcast_tensors.<locals>.<genexpr>  sP      .
 +1 GG11!4ww44Q78*s   A A")r   rk   ri   r   r   r  reducer2  r   r  r  r-  ro   )rJ  ra   r  r   sizess        re   r   r     s    
6{afQi$/$fQi00(//!#A&#A2 F G1::<((EV3Z3v;&# .
 E6*.
 +

 q&!Aq  Nrg   c                    | S r   r   r   s    re   nopr    s    Hrg   
lift_freshc                   t        | t              sJ |(t        t        j                  | j                              S t        |t
        t        j                  f      r)t        j                  j                  j                  |      nt        d |D              }t        t        | j                               |      }t!        t        |t              s|fn|      }g }t#        | j                               D ]X  \  }}||v r>t        j                  j                  j%                  t        j&                  |d            rH|j)                  |       Z || j                         k7  rt+        | |      S | S )Nc              3  n   K   | ]-  }t         j                  j                  j                  |       / y wr   )rR   r]   r   	guard_intr   ds     re   r   zsqueeze.<locals>.<genexpr>  s&     >#Q177##--a0#s   35r.   )rk   rB   rA   r   rm   r   r   r   rR   r]   r   r  r   r   r   r   r&   rj   guard_or_falseEqro   r  )r   r   dims	new_shaper  ss         re   squeezer    s   a###
{++AFF344 cC,- 	
""3'>#>> 
 C

-s
3CJsE$:sfDDI!**,'1T	agg..==ehhq!nMQ (
 "+ajjl!:49AArg   c                ,    t        t        | |            S r   )r  r  )r   r   s     re   squeeze_copyr!    s    C!!rg   c                    t        | |      }t        | t              sJ t        |t              sJ |j                  | _        | S r   )r  rk   rB   rm   r   r   vals      re   squeeze_r%    s=    
!S/Ca###c9%%%XXAFHrg   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     re   r'  r'    ?    qE44	W	B?>"EJJ?BBrg   c                    t        |       rt        | dt        j                        S t	        d      } t        |t        j                        |       S )NFr   isnanr  r(  r*  s     re   r-  r-    r+  rg   c                f    t        |       rt        |       S t        d      } t        |      |       S )Nceilr   r  r=   r  r*  s     re   r/  r/    s/    qQx	V	B>"a  rg   c                f    t        |       rt        |       S t        d      } t        |      |       S )Nfloorr0  r*  s     re   r2  r2    /    qQx	W	B>"a  rg   c                f    t        |       rt        |       S t        d      } t        |      |       S )Nroundr0  r*  s     re   r5  r5    s/    qQx!!~b!!$$rg   c                f    t        |       rt        |       S t        d      } t        |      |       S )Ntruncr0  r*  s     re   r7  r7    r3  rg   c                   t        | g      \  } t        | t        j                        rt	        j
                  | t        |            S t        | t              sJ t        |t        t        f      sJ t        | j                               t        |      k(  r| S t        | j                               st        j                  j                  j                  t        | j                                     }|dkD  rOt        |      sD| j!                  t        j                  j                  j                  t        |            |z         t        t	        j
                  | j"                  t        |                  S r  )rN  rk   r1   r{  r7   r   r   rB   ri   r   r#   rR   r]   r   guarding_hint_or_throwrO   
mark_reuserm   )r   r  x_size_products      re   r-  r-    s   aS!DQ!R__%  E%L11a###edE]+++QZZ\eEl* .))@@!**,'
 A&;E&BLL  77e8LM!" Z&&qvvuU|<==rg   c                    t        |      }|D ]  }d||<   	 | }t        |      D ]  \  }}|dk7  st        ||      } t        ||      S Nrh  )ri   rj   	unsqueezer-  )r   shapebroadcast_dimensionsr  broadcast_dimensionr  rU  r   s           re   broadcast_in_dimrB    s\    UA3!#
  4 	
AA,Q7!S!A  !Urg   c                6    t        | |j                               S r   )r-  r   )r   r1  s     re   	expand_asrD  #  s    !QZZ\""rg   c                j   t        | j                               t              t              kD  rKt        j                  j
                  gt              t              z
  z  z   t        | t                    } t              t        | j                               k(  sJ t        | j                               }d}t        t                    D ]  }|   dk(  rd}||   |   z  ||<    |r*t        || j                         | j                               S t        d t              D              rt        t        | |            S fd}t              st        |      s{t         j"                  j$                  j'                  t)                    }|dkD  rD| j+                  t         j"                  j$                  j'                  t)        |            |z         | j-                         t/        j0                  | j                         | j                         |t        |            S )NFr   Tr   rv   c              3  :   K   | ]  \  }}|d k(  xs |d k(    ywr.   Nr   r  s      re   r   zrepeat.<locals>.<genexpr>:  s'     
A*@$!QAFa1f*@   c                   t        |       t              k(  sJ t        |       } t        t                    D ]G  }|   dk7  s|   dk(  rt        j                  j
                  | |<   2t        | |   d|         | |<   I  |       S r  )r   ri   r  r   r)  Zeror+   )r:  rs   old_sizerepeatsx_loaders     re   r\  zrepeat.<locals>.inner_fn?  s    5zS\)))Us7|$AqzQA;!#$ww||E!H.uQxHQKHE!H % rg   r]  )ri   r   r   r   r)  r*  r  r  emptyr   rn   r  r  r  r-  r#   rR   r]   r   r9  rO   r:  r_  r?   r   )	r   rM  new_sizezero_tensorrs   r\  old_size_productrL  rN  s	    `     @@re   repeatrS  (  s   AJJL!H
7|c(m#GGKK=CL3x=$@AHLDN#w<3qzz|,,,,AJJL!HK3w< 1:?KqkGAJ. !
 XQ[[]1<<>JJ

A#gx*@
AAVAx())	 !*3H3R77++BB(#
 a LL  77h8OP#$
 }}H||~kkmH~	 rg   c                T    t        t        j                  | j                  |            S r   )rB   rD   r   rm   )r   r  s     re   r  r  `  s     T[[/00rg   c                    t        | t              sJ t        |t        t        f      sJ t        t	        j
                  | j                  t        |                  S r   )rk   rB   ri   r   r>   r   rm   )r   r  s     re   permuterV  g  sF    a###dT5M***[''d<==rg   c           
     	   ddl m}m} t        | t              sJ t        | |d      }| j                         |   }t        j                  |      }t        |t        j                        s|dkD  sJ |       	 |dk(  r1t        j                  j                  j                  ||      r|dk(  r| S dd}	d\  }
}|}|r[ |	||d      }
|;t        j                  j                  j                  |t         j"                        r|}n
 |	|||      }|
||
|}}d}|sIt        j                  j$                  } |t        j                  j                  j&                  ||j(                  j+                  di       ni       }|r|j-                         D ]  \  }}| |d      t/        j0                  |      fk(  r]t3        j4                  |||||      }t        j                  j7                  |      |_        t        j                  j;                  |       | |d	      fk(  st=        d
       t	        t2        j>                  jA                  | jB                  |||||            S |sJ  |t        j                  j                  j&                  t        j                  j$                  j(                  d         }|J tE        |      dk  sJ |       d\  }}|j-                         D ]7  \  }}| |d      t/        j0                  |      fk(  r|})| |d	      fk(  s6|}9 |
|J t3        j4                  ||||| j                         |         }t        j                  j7                  |      |_        t        j                  j;                  |       |}| jG                         | jI                          |
6|J | jK                         jL                  |
| jO                         |   z  z   }nt3        jP                  ||| jK                         jL                  | jO                         |   | j                         |   d      }t        j                  j7                  |      |_        t        j                  j;                  |       |}tS        | j                               }tS        | jO                               }|||<   ||xx   |z  cc<   tU        | |||      S # t        $ r Y 5w xY w)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.   c           	        | |S d }t        j                  |       } t        j                  |      } |t        j                  t        j                  | d      t        j                  | |                  r| S  |t        j                  t        j
                  | d      t        j                  | |                   r| |z   S  |t        j                  | |            r|S  |t        j
                  | |             ry |t        j                  | d            rt        j                  | |      S  |t        j
                  | d            rt        j                  | |z   d      S y )Nc                T    t         j                  j                  j                  |       S r   )rR   r]   r   r  r   s    re   rC  z5slice_.<locals>.compute_slice_index.<locals>.<lambda>  s    qww''66q9rg   r   )	r   r-  AndGeLeLtGtMinMax)r:  r   r  r   s       re   compute_slice_indexz#slice_.<locals>.compute_slice_index  s   =N9U#||D!eii*EHHUD,ABCL		%((5!,ehhute.DEF4<%&K&'"#99UD))"#99UT\1--rg   NNFunbacked_bindingsr   storage_offsetzWUnexpected storage_offset unbacked binding when both start and end indices are resolvedclampr,   Tr   )+%torch.fx.experimental.symbolic_shapesrX  r%   rk   rB   _validate_dimr   r   r-  r   rR   r]   r   statically_known_leq	TypeErrorstatically_known_equalssysmaxsizer^   	shape_envr|  r}  r   pytreeSequenceKeyr1   DynamicSliceSizeregister_bufferr   register_operationr   	SliceViewr   rm   r   maybe_get_layoutr  
get_layoutoffset
get_strideDynamicSelectStorageOffsetri   
as_strided)r   r   startendsteprg  rX  r%   r   rb  start_index	end_indexambiguous_slicer^   node_unbacked_bindingssymkeypathb_sizerd  sym_sizesym_storagerP  new_storage_offset	b_storage	new_sizesnew_stridess                             re   slice_r  p  s   
 a###
3
"C::<D<<DdEJJ'4!89T93QJ  55dC@	H
. (K
 O)%q9 ?qww//GG 
 I+Ct<I"y'<$i3E#O ww++!:GG&&' !!"5r:	"
 " 6 < < >W}V4f6H6H6MNN00eS$MF"#''"9"9&"AFKGG..v6/?!@ BB )=  !? LLUCUK
 	
 L51	""AGG$8$8$=$=>Q$R ((( !Q&9(99&&Hk)//1W}V,f.@.@.EFFH'78::K	 2 )"333  	

SF ''))&1FKGGv&H#			"""\\^22[1<<>RUCV5VV11LLN!!LLN3JJL
	 00;		""9-(QZZ\"Iq||~&KIcNaK1CDDW  s   75S 	S#"S#c                   d }d }t        | t              rZt        | j                  t        j                        r6| j                         }| j                  }| j                  j                         } | j                          t        j                  |       st        d|  d      t        j                  |       \  }}t        j                  |r|n|j                  |r|n|j                  |D cg c]  }t        j                  |       c}|D cg c]  }t        j                  |       c}t        j                  |xs d            }	t        t        j                   ||	            S c c}w c c}w )Nzunrealized as_strided(z, ...)r   rm   layout)rk   rB   rm   r1   r5   rn   r   unwrap_viewr  is_storage_and_layoutr   as_storage_and_layoutFixedLayoutrv   r   r-  r   )
r   r   stridere  
new_device	new_dtypestorage
old_layoutr  
new_layouts
             re   r{  r{    s   JI!YJqvvr{{$C \\^
GG	FF IIK##A&!$:1#V"DEE2215GZ 
j&7&7	J$4$4"&'$Qa$'"()&Qa&)^(q)J R''WZHII	 	()s   ,E)E.c                d    t        | t              sJ t        | |||      j                  | _        | S r   )rk   rB   r{  rm   )r   r   r  re  s       re   as_strided_r  6  s/    a###48==AFHrg   c                4    t        | |||      }t        |      S r   )r{  r  )r   r   r  re  r  s        re   as_strided_copyr  =  s    48F=rg   c                    g d} D ]1  }j                  |||j                            z   f       d   d   }3  D cg c]  }|j                          c} fd}t         d   j                               }d   d   |<   t	        j
                   d   j                          d   j                         ||      S c c}w )Nr   rh  c           	       
 t        j                  |    t        j                        }g }g }t	        t                    D ]?  

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                  ||      }|j                  |       t        |       t           
   d   z
        <   |j                  t        j                  |
fdd             B |d   }	t	        t              dz
  dd      D ]  
t        j                  |
   |
   |	      }	! |	S )Nr   r.   c                                S r   r   )rs   idx_loadinputs_loaderss   re   rC  z1pointwise_cat.<locals>.inner_fn.<locals>.<lambda>o  s    -N1-h7rg           rh  r,   )rQ   rR  r{   int64r  r   rS  geltand_ro   ri   r)   maskedr  )rU  idx_dimmasksmasked_loadsr|  r}  
start_condend_condmasknext_valrs   r  r   rJ  r  inputs_rangess             @@re   r\  zpointwise_cat.<locals>.inner_fnM  s   ..S5;;7s6{#A 6 Q,^^M!$4Q$7E 
 ..q!1!!4ekkBC/Jvvgs+HAvc&kAo%!xx
H5LLCyH %Xc]]15Ea5H%HIHSM

75 $D  #Fq("b1AyyaQH 2 rg   r]  )ro   r   r_  ri   r?   r   rn   r   )rJ  r   prev_endr   r\  rP  r  r  s   ``    @@re   pointwise_catr  C  s    9;MHh3<<>#3F(FGH $R(  4::6Ccoo'6:N.` F1I&&()H!"%b)HSMay##%Qi!!#	 k ;s   Cc           	       	
 t        j                               dk(  sJ d       t        j                               dk(  sJ d       | j                         t        j                  k(  rt        | t        j                        } | j                         t        j                  k(  sJ d| j                                 t        | j                               k  s!J dt        | j                                       | j                         j                         	j                         
	
f	d}t        j                  | j                         || 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                b  	 | 
   f} |       } |      } |      }t        t        j                        \  }}j                  t        j                  k7  r$t	        j
                  |t        j                        }j                  t        j                  k7  r$t	        j
                  |t        j                        }t	        j                  |      }t	        j                  ||z        |z   }t	        j                  |t	        j                  ||            }	t	        j
                  |	      S r?  )_create_constantsr{   r  r   rQ   r   r  
reciprocalr5  maximumminimum)rU  channel_idxr  scale
zero_pointqminqmax	inv_scaler$  clampedaxisr   input_loader	quant_max	quant_minscalesscales_loaderzero_pointszero_points_loaders             re   r\  z;quantized_decomposed_quantize_per_channel.<locals>.inner_fn  s    4ylS!k*'4
&y)5==Q
d<<5==(LL6E+j%++>JNN5)	ii	)*Z7++dCKKc$:;||GU++rg   r]  )r   r   r   r{   rG  r   r  r_  r?   r   rn   )r  r  r  r  r  r  r   r\  r  r  r  s    `````` @@@re   )quantized_decomposed_quantize_per_channelr    s?    v !Q&=(==&{##%&!+G-GG+ENN*.??- 
FuGXFYZ- #enn&'' 
!#enn&6"7!89' $$&L&&(M$002, ," !~~	 rg   c           	     "     j                          t         t        j                          fd}t	        j
                   j                          j                         |t         j                                     }|j                          |S )Nc                    t         j                  j                         5  t        j                   j                         |             cd d d        S # 1 sw Y   y xY wr   )r1   ComputedBufferforce_realizerQ   device_assert_asyncr_  )r:  r   r   s    re   r\  z_assert_async.<locals>.inner_fn  sB    ,,.**+=4+;+;+=e+DcJ /..s   *AAr]  )
r  r   r{   r   r?   r   rn   r   ri   r   )r   r   r\  assertion_ops   ``  re   _assert_asyncr    sm    LLND%**%DK ## nnDMMO$	L rg   c                    t        | |      S r   r  r   s     re   lower_assert_asyncr        s##rg   c                    t        | |      S r   r  r   s     re   lower_assert_functional_asyncr    r  rg   )r  c               v  	
 t        j                               dk(  sJ d       t        j                               dk(  sJ d       | j                         |k(  sJ d| d| j                                 t        | j                               k  s!J dt        | j                                       t        j                  | j                         	j                         
j                         	
fd}t        j                  | j                         || j                               S )	Nr.   r  r  Expecting input to have dtype , but got dtype: r  c                   |    f} |       } 
|      } |      }	j                   t        j                  k7  r$t        j                  |t        j                        }j                   t        j                  k7  r$t        j                  |t        j                        }t        j
                  t        j                  |t        j                        |      |z  }t        j                  |      }|S r   )r   r{   r  rQ   r   sub)rU  r  r  r  r  r$  r  r  r  r  r  r  r  s         re   r\  z=quantized_decomposed_dequantize_per_channel.<locals>.inner_fn  s    4ylS!k*'4
<<5==(LL6E-j%--@Jggcll5%--8*EMll3	*
rg   r]  	r   r   r   r{   r  r_  r?   r   rn   )r  r  r  r  r  r  r   r  r\  r  r  r  s    ```   ` @@@re   +quantized_decomposed_dequantize_per_channelr    s-    v !Q&=(==&{##%&!+G-GG+??% 
(/@AR@ST% #enn&'' 
!#enn&6"7!89' MM	$$&L&&(M$002  !~~	 rg   c                   | j                         t        j                  k(  rt        | t        j                        } | j                         t        j                  k(  sJ d| j                                 | j                         fd}t        j                  | j                         t        j                  |t        |      t        |            | j                               S )Nr  c                L    
|       }t        d|z  |t        j                        \  }}t        j                  ||z        |z   }t        t        j                        \  }}t        j
                  t        j                  ||      |      }t        j                  |	      S )N      ?r   )r  r{   r  rQ   r5  r  r  r   )rU  r  r  r  r  r$  r  r  r  r   r  r  r  s            re   r\  zBquantized_decomposed_quantize_per_tensor_default.<locals>.inner_fn  s    S! 1%K5==!
	: ii	)*Z7&y)5==Q
d++ckk#t4d;||GU++rg   r  r  r]  )r   r{   rG  r   r  r_  r?   r   rn   r  r#  r5  r   r   )r  r  r  r  r  r   r\  r  s      ``` @re   0quantized_decomposed_quantize_per_tensor_defaultr    s     ENN*.??- 
FuGXFYZ- $$&L, !""E%LS_
 ~~ rg   c               l   | j                         |k(  sJ d| d| j                                 t        j                  | j                         fd}t	        j
                  | j                         t        j                  |t        |      t        |            | j                               S )Nr  r  c                     |       }t        ||t        j                        \  }}t        j                  t        j
                  |t        j                        |      |z  }t        j
                  |      }|S r?  )r  r{   r  rQ   r  r   )rU  r  r  r  r$  r  r  s        re   r\  zDquantized_decomposed_dequantize_per_tensor_default.<locals>.inner_fnI  s]    S!-eZu}}Uzggcll5%--8*EMll3	*
rg   r  r]  )r   r{   r  r_  r?   r   rn   r  r#  r5  r   r   )	r  r  r  r  r  r   r  r\  r  s	         ` @re   2quantized_decomposed_dequantize_per_tensor_defaultr  3  s     ??% 
(/@AR@ST% MM	$$&L !""E%LS_
 ~~ rg   c           	     T  	
 | j                         t        j                  k(  rt        | t        j                        } | j                         t        j                  k(  sJ d| j                                 t        j                               dk(  s9t        j                               dk(  rj                         d   dk(  sJ d       t        j                               dk(  s9t        j                               dk(  rj                         d   dk(  sJ d       | j                         j                         	j                         
| j                         	
f	d}t        j                  | j                         || j                               S )Nr  r   r.   expect scale as scalar tensor"expect zero_point as scalar tensorc                l  	  
|       } t        j                               dk(  rdnd      } t        j                               dk(  rdnd      }j                  t        j                  k7  r$t        j                  |t        j                        }j                  t        j                  k7  r$t        j                  |t        j                        }rOj                  dk(  r@t        j                  |t        j                  |      |      }t        j                  |	      S 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{   r  rQ   r   r   rT  r  round_to_intr5  r  r  r  )rU  r  _scale_zero_pointr$  r  r  r  rv   r   r  r  r  r  scale_loaderr  zero_point_loaders           re   r\  zAquantized_decomposed_quantize_per_tensor_tensor.<locals>.inner_fnw  s5   S!c%..*:&;q&@dbI'ENN4D0E0JPRS;;%--'\\&%--8Fu}},,,{EMMBKfkkU*''%!7EC##C//iiv 667+E&y)5==Q
d++ckk#t4d;||GU++rg   r]  )r   r{   rG  r   r  r   r   r_  rn   r?   r   )r  r  r  r  r  r   r\  rv   r  r  r  s    ````` @@@@re   /quantized_decomposed_quantize_per_tensor_tensorr  Z  su    ENN*.??- 
FuGXFYZ- u~~ A%ENN"u~~'7':a'?'&'  z""$%*J!"a'J,?,?,A!,D,I,+,  $$&L$$&L"..0F, ,  !~~	 rg   c                 	
 t        j                               dk(  s9t        j                               dk(  rj                         d   dk(  sJ d       t        j                               dk(  s9t        j                               dk(  rj                         d   dk(  sJ d       | j                         |k(  sJ d| d| j                                 t        j                  | j                         j                         	j                         
	
fd}t        j                  | j                         || j                               S )	Nr   r.   r  r  r  r  c                V    |       } t        j                               dk(  rdnd      } 
t        j                               dk(  rdnd      }j                  t        j                  k7  r$t        j                  |t        j                        }	j                  t        j                  k7  r$t        j                  |t        j                        }t        j                  t        j                  |t        j                        |      |z  }t        j                  |      }|S )Nr.   r  r   )r   r   r   r{   r  rQ   r   r  )rU  r  r  r  r$  r  r  r  r  r  r  s        re   r\  zCquantized_decomposed_dequantize_per_tensor_tensor.<locals>.inner_fn  s    S!c%..*:&;q&@dbI'ENN4D0E0JPRS;;%--'\\&%--8Fu}},,,{EMMBKggcll5%--8+FOll3	*
rg   r]  r  )r  r  r  r  r  r   r  r\  r  r  r  s    ``   ` @@@re   1quantized_decomposed_dequantize_per_tensor_tensorr    sU    u~~ A%ENN"u~~'7':a'?'&'  z""$%*J!"a'J,?,?,A!,D,I,+,  ??% 
(/@AR@ST% MM	$$&L$$&L"..0
 
 !~~	 rg   c                     d   j                         j                  dk(  }|rt        d  D              rp D ]  }|j                           t        d  D              rt	        t
        j                  g  \   } t        t
        j                  j                         |      S t               dk(  rt         d         S t         d   |d      }t         dt        j                  i} D cg c]  }t        ||       c} ddd fd	t!        fd
 D              }dfdt"        j$                  rt'         |      S |r)t)        t*        j,                  j/                   |            S fdd}dddt               |k  s1t               t"        j0                  k  rt        fd D              rt        fdt2        j4                  j6                  D              }	t!        fd D              xr |	}
d fd} |       }t        fd D              xr t!        fd D               }|s|
s|r|st'         |      S t)        t*        j,                  j/                   |            S c c}w )z@Lower aten.cat, choosing between pointwise_cat and ConcatKernel.r   r   c              3  |   K   | ]4  }|j                         t        j                  t        j                  fv  6 y wr   )r   r{   int8uint8r   r  s     re   r   zcat.<locals>.<genexpr>  s-      DJ5ejj%++66Fs   :<c              3  T   K   | ]   }t        |j                               d k(   " yw)r   N)r   r   r  s     re   r   zcat.<locals>.<genexpr>  s"     >ves5>>#$)vs   &(r.   r   c                   t        | t              rJt        | j                  t        j                        r| j                  j                         S | j                  S t        | t        j                        r| j                  S | S r   )rk   rB   rm   r1   r5   r  r   r   s    re   unwrap_tensorzcat.<locals>.unwrap_tensor  sV    a#!&&"++.vv))++vva'66Mrg   c                    t        | t        j                        xr$ t        | j                  t        j                        S r   )rk   r1   r  rm   r@   )rw   s    re   is_reductionzcat.<locals>.is_reduction  s)    !R../TJqvvr||4TTrg   c                    t        | t        t        j                  f      r  |             S  |       xs> t        | t        j                        xr" t        fd| j                         D              S )Nc              3  h   K   | ])  } t         j                  j                  |             + y wr   )rR   r]   
get_buffer)r   readcan_fuse_reductions     re   r   z2cat.<locals>.can_fuse_reduction.<locals>.<genexpr>  s-      .D #177#5#5d#;<.s   /2)rk   rB   r1   r   r?   r  get_read_names)rw   r  r   r  s    re   r  zcat.<locals>.can_fuse_reduction  sh    a)R]]34%mA&677O !R\\*  ,,. 	
rg   c              3  .   K   | ]  } |        y wr   r   r   rw   r  s     re   r   zcat.<locals>.<genexpr>  s     B6a.q16   c                6   t        j                  |       r:t        j                  | d      \  }}t         j                  j	                  |       S t        | t        t         j                  f      r  |             S t        | t         j                        ryy)NF)freezeT)	r1   r  r  ConcatKernelcan_realize_into_without_copyrk   rB   r   r?   )r   r  _should_lower_cat_inputr  s      re   r  z#cat.<locals>.should_lower_cat_input  sx     ##A&11!EBJGQDDWMMMa)R]]34)-*:;;a&rg   c                H   t        | t        t        j                  f      r  |             S t        | t        j                        sy| j                         j                  }| j                         D ]*  }| t        j                  j                  |            z  }, |S r  )rk   rB   r1   r   r?   inner_fn_opcountnum_opsr  rR   r]   r  )r   countr  op_countr  s      re   r  zcat.<locals>.op_count  s    a)R]]34M!,-- !R\\*""$,,$$&DXagg00677E ' rg   r   r,   c                n    | t         j                  j                  t         j                  j                  fv S r   )r  catr  constant_pad_ndr`   s    re   additional_pointwise_opsz%cat.<locals>.additional_pointwise_ops#  s(    dhh&&(<(<(D(DEEErg   c              3  4   K   | ]  } |      k    y wr   r   )r   rw   MAX_SIMPLE_OP_COUNTr  s     re   r   zcat.<locals>.<genexpr>(  s     CFq22Fs   c              3  6   K   | ]  }t        |        y wr   rI   )r   user  s     re   r   zcat.<locals>.<genexpr>*  s!      
+ S":;+r   c              3  .   K   | ]  } |        y wr   r   r   r   r  s     re   r   zcat.<locals>.<genexpr>1  s     >v&s+vr	  c                    t         j                  yj                  d   } t        | t        t
        f      r| }n)t        | t        j                  j                        r| g}nyfdt        |      D ]T  \  }}t        |d      rt        |j                        dk  r+t        fd|j                  D              r y |      sT y y)NFr   c                    t        | t        t        j                  f      r  |             S t        | t        j                        S r   )rk   rB   r1   r   r?   )r   is_unrealized_pointwiser  s    re   r#  zKcat.<locals>.any_input_has_multi_consumers.<locals>.is_unrealized_pointwiseG  s7    a)R]]!;<2=3CDD!!R\\22rg   r_   r.   c              3  >   K   | ]  }|ust        |        y wr   r  )r   ur^   s     re   r   z=cat.<locals>.any_input_has_multi_consumers.<locals>.<genexpr>Q  s     WIq,AV'*Is   	T)rR   r^   rt   rk   ri   r   r{   fxNoder  r   r   r_   r  )fx_argsinput_nodesr   ir_inputr^   r#  rJ  r  s       @@re   any_input_has_multi_consumersz*cat.<locals>.any_input_has_multi_consumers;  s    >>L#"''*G'D%=1%GUXX]]3&i3
 "%[&!9XsG,CII!0C WCIIWW +84 ": rg   c              3  .   K   | ]  } |        y wr   r   r   s     re   r   zcat.<locals>.<genexpr>\  s      "
39C"3'6r	  c              3  .   K   | ]  } |        y wr   r   r  s     re   r   zcat.<locals>.<genexpr>^  s     <V(+Vr	  )r   zTensorBox | ir.StorageBoxr   	ir.IRNoder   r   r`   torch._ops.OpOverload)rn   r   r  r  require_channels_lastr  r  r  r  r   r  ri  r   r   rE  r   r  r/   force_pointwise_catr  rB   r1   r  r   max_pointwise_cat_inputsrR   r^   r_   )rJ  r   
cpu_devicer  r  r   r   fusable_reductionMAX_COMPLEX_POINTWISE_CATpointwise_usesfuse_pointwise_user+  has_multi_consumershorizontal_fuse_catr  r  r  r   r  r  r  s   `             @@@@@@@re   r  r    sX    %%',,5Jc DJ 
 EMMO >v>>-dhh@@IFA1 0 01&#>>
6{aVAY
q	3
*C	%D%L%LE /55fshsE"f5F
U

 B6BB !!VS)) //<==$ !"F 6{//	V77	7CFCC 
~~++
 
 >v>>Q> 		> <=! "
39"
 
 =<V<<< 	 ##6?P --R__++FC899[ 6s   6Jc                  
 | j                         t              t              t              t        k7  fd       t        j
                  j                  j                  t        j                  |d            }|r\t        j
                  j                  j                  t        j
                  j                  j                     |z            d      }n[t        j
                  j                  j                  t        j
                  j                  j                        |z
        d      }d
|r| df
nd|f
t              D cg c]  \  }}|fvs| }}}|j                  |       
fd}	t        t        j                   j#                  | ||	            S c c}}w )N)rU  rankc                     d  d S )Nz(diagonal dimensions cannot be identical z, r   )dim1dim2s   re   rC  zdiagonal.<locals>.<lambda>o  s     HbQUPVWrg   r   )r   r   c                    | d   }dgt        	      z  }d}t              D ]1  }|k(  r|d   z   ||<   |k(  r|d   z   ||<   %| |   ||<   |dz  }3 |t        	      dz
  k(  sJ |S )Nrh  r   r.   r,   )r   r  )
rU  diag_idxoriginal_idxcur_dimr  base_idxr?  r@  num_dimsoriginal_shapes
        re   	reindexerzdiagonal.<locals>.reindexer  s    r7sS00xADy"*Xa["8Qd"*Xa["8Q"%g,Q1 ! #n-1111rg   )r   r   r   r   rR   r]   r   evaluate_exprr   r^  evaluate_maxevaluate_minrj   ro   rB   r1   GenericViewr   )r  rx  r?  r@  offset_negative	diag_sizers   r  r  rH  rE  rF  rG  s     ``      @@@re   diagonalrO  g  s   ^^%N>"H84D84D	W gg&&44UXXfa5HIOGG$$11GG))t$v-t$ 
	 GG$$11GG))t$t$v- 
	 HGQ<v;$^4N441a$8MQ4EN	LL   R^^**5%CDD' Os    F:/F:c                0    t        t        | |||            S r   )r  rO  )r  rx  r?  r@  s       re   diagonal_copyrQ    s    %t455rg   c                P    t        |       }t        ||||      }t        ||       |S r   )r  rO  	mutate_to)r  srcrx  r?  r@  r0  ra   s          re   diagonal_scatterrU    s*    5\FffdD1FfcMrg   c                h   t        j                  |      }t        j                  | j                         |         }d }t        j                  j
                  j                  t        j                  |d            r||z   }n?t        j                  j
                  j                  t        j                  |d            r|}|t        |      rg| j                          | j                         }| j                         }| j                         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/        |	j1                                     \  }
}| j                          | j                         }| j                         }|
}t3        j4                  |
|| j                         j                  ||   | j                         |   d      }t        j                  j7                  |      |_        t        j                  j;                  |       ||= ||= t        | |||      S )Nr   r.   Frf  rd  )r   r-  r   rR   r]   r   r  r^  r\  r$   r  ry  rw  rx  r{  r  r  r%   ro  r^   r|  r   rF  iterr   r1   rz  rs  r   rt  )r   r   rU  r   actual_indexrP  
new_strider  slice_resultrd  unbacked_offset_symr  buffers                re   selectr]    sE   
,,s
C<<

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

SF ''))&1FKGGv&3a:/ABBrg   c           
     r   t        | |d      }|}t        |t        t        f      s`| j	                         |   }t
        j                  j                  j                  t        ||z   dz
  |            }|g|z  }||dz
  |z  z
  |d<   g }d}|D ](  }||z   }	|j                  t        | |||	d             |	}* |S )Nr   r.   rh  Frf  )ri  rk   ri   r   r   rR   r]   r   r  r(   ro   r  )
r   r  r   sizes_x_sizechunksr  r|  r   r}  s
             re   splitrb    s    
3
"CF edE]+c"!!++HVe^a5G,OP6!vzU22r
 FEdl 	fQUCu=>  Mrg   c                    t        | ||      S r   )rb  )r   r  r   s      re   split_with_sizesrd  		  s    E3rg   c                    t        | |d      }t        j                  j                  j	                  | j                         |         }t        |      D cg c]  }t        | ||       }}|S c c}w r  )ri  rR   r]   r   r  r   r  r]  )r   r   r`  rs   r  s        re   unbindrf  	  s`    
3
"CWW''

S(9:F).v7AfQQF7M 8s   A-c                @   | j                         }t        |      }t        ||      |dk(  rt        t	        | d      |d      S |   }t
        j                  j                  }|j                  ||       |j                  d       t        ||z
        dz   }|j                  |      dkD  r-| j                  |j                  t        ||z  |                   g |d  ||dz   d  |}	fd}
t        t        j                   j#                  | |	|
            S )Nr   F)r}  rg  r.   c                D    | d   |    z  z   }g | d  || dz   d S )Nrh  r.   r   )rU  dim_idxr   r~  s     re   rH  zunfold.<locals>.reindexer,	  s?    b'CHtO+8Tc8G8c#'B&788rg   )r   r   r   r  r>  rR   r]   r   	check_leqcheck_ltr(   r9  r:  r'   rB   r1   rL  r   )r   	dimensionr   r~  r  ndimdim_sizer   new_dim_sizeout_sizerH  r   s      `       @re   unfoldrq  	  s   JJLEu:D
4
+Cqyi1o4u==SzHwwHtX&aHtOT2Q6L&&x014	++GL44G,RS	
 EtDlDU379-=DtDH9 R^^**1h	BCCrg   c                    t        | |d      }t        | j                               }|j                  |t        j
                  j                         t        | |      S r  )ri  ri   r   insertr   r)  r*  r  )r   r   r  s      re   r>  r>  3	  sE    
3
"CQZZ\"IS%''++&9rg   c                    t        | |      }t        | t              sJ t        |t              sJ |j                  | _        | S r   )r>  rk   rB   rm   r#  s      re   
unsqueeze_ru  ;	  s>    
As
Ca###c9%%%XXAFHrg   c                   t         j                  j                  j                  j	                  t        j                  |            }t        | j                               }|dk  r|||z   z  }d|cxk  r	||z   k  sJ  J |S r  )	rR   r]   r   ro  rI  r   sympifyr   r   )r   r   rx  rm  s       re   ri  ri  D	  sr    
''


$
$
2
25==3E
FCqzz|D
Qwtf}#dVm#####Jrg   rh  c                   t        | |d      }t        j                  j                  j	                  | j                         |         dz  }t        | |d|d      }t        | |||dz  d      }t        |t        |            S )Nr   r,   Frf  )	ri  rR   r]   r   r  r   r  rz  sigmoid)r   r   new_lenr   r/  s        re   glur{  M	  ss    
3
"Cgg((c):;q@Gq#q'/Aq#w!59Aq'!*rg   c                N     |rt         j                           fd}d|_        |S )Nc                 x    d }t        j                  |t        j                  j                  g| i |      S )Nc                Z    t        | t        j                        r| j                         S | S r   rk   r1   r9   wrap_for_loweringr   s    re   wrap_tensorsz7fallback_handler.<locals>.handler.<locals>.wrap_tensors]	  #    ,6q")),D1&&(K!Krg   )rp  tree_mapr1   FallbackKernelr   )rt   r   r  kernels      re   handlerz!fallback_handler.<locals>.handler\	  s<    	L "++226KDKFK
 	
rg   T)r  r   _is_fallback_handler)r  r  r  s   `  re   r  r  X	  s&    f
 $(G Nrg   c                 .    t        j                  d       y )NzjTorchinductor does not support code generation for complex operators. Performance may be worse than eager.)warningswarnr   rg   re   _warn_complex_not_supportedr  j	  s    MMtrg   c                   | j                         rt                y| j                  ry| j                  ry| j                  t
        j                  k(  r|syt        |j                  t
        j                  j                        xrr |j                  t        j                  j                  t        j                  j                  t        j                  j                  t        j                   j                  fv xsE t        |j                  t
        j                  j                        xr t#        |j                         S y)z0Do not support reading or writing to this tensorTF)r  r  is_meta	is_sparser   r{   float8_e8m0fnurk   ra   r|   r}   r  r  r  r  r  
_scaled_mmrJ   )rw   rc   s     re   unsupported_input_tensorr  s	  s    ||~#%yy{{ww%&&&
 t{{EJJ$9$9: 		  

""''	Y 4;;

(=(=>W74;;CW

 
	
 rg   c                *   t         j                  j                  t        j                  j
                  j                  j                  f}||j                  |v r| j                         ryt        | |      ry| j                  xr t        j                  S )z2Do not support writing tensor but can read from itFT)r  r  r   r{   rQ   r  r  r  ra   r  r  is_cpur/   disable_cpp_codegen)rw   rc   supported_complex_viewss      re   unsupported_output_tensorr  	  so     					,,44 DKK+BBq||~4(882222rg   c                R     j                   t        j                  j                  u ry j                  dk(  ry j                   t        j
                  j                  u ry fd}t        j                   j                  i  j                  D ]  } ||d      s y  | d      S )NFplaceholderc                D   t        | t        j                  j                        syd| j                  vryt        j                  | j                  d         D ]F  }t        |t        j                  j                        s(|rt        |      s7 yt        |      sF y y)NFr$  T)rk   r{   r&  r'  r|  rp  tree_leaves_subclasses
FakeTensorr  r  )inp_out_node	is_outputr|  rc   s      re   check_skip_conditionzCfallback_node_due_to_unsupported_type.<locals>.check_skip_condition	  s    ,6)))&&|'8'8'?@DdE$5$5$@$@A,T48+D$7 A rg   )r  T)
ra   r  view_as_complexr  r`   lift_fresh_copyrp  arg_tree_leavesrt   r   )rc   allow_cpu_inputsr  r   s   `   re   %fallback_node_due_to_unsupported_typer  	  s    {{d**222ww- {{d**222* %%tyy@DKK@u5 A  55rg   c                @   | |       nt         }| |vs|s
J d|         |rt        t        j                  d            rt	        | g      rt
        j                  r&| t        j                  j                  j                  v si|sgt        j                  j
                  j                  r4dt        j                  j
                  _        t        j                  d       t        d|  d      fd}t!        | t        j"                  j$                        r*| j'                         D ]  }t)        | |      } ||        y t!        | t        j"                  j*                  t        j"                  j,                  f      r	 ||        y t/        d|  d	t1        |              )
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                n    t        |        t        |         t        | d       t        |             S NrP  )r   r   r$  r  )op_overloadlayout_constraints    re   register_fallbackz(make_fallback.<locals>.register_fallback	  s;    !+.(!+/@AG $G[)
 	
rg   zUnsupported fallback z with type )r3   r   osgetenvr4   r/   fallback_randomr{   _decompdecompositions_for_rngextra_random_decomps_dynamosuppress_errorslogwarningr   rk   r|   r   r   r   r}   HigherOrderOperatorRuntimeErrorr   )	r`   r  r  override_decompget_decomp_fncheck_decompsr  olr  s	    `       re   make_fallbackr  	  s\    (5'@MOnM]"o 
4RD95 	4!t$ ""emm::OOO ==//38EMM  0KKH RD !f f
 	

 "ejj112,,.B!"b/Kk* ! 
B..

0N0NO	P"22$k$r(LMMrg   c                T    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   r@  r{   r  )r?  numelr  s      re   philox_rand_offsetr  
  s-     E	 %u{{++rg   c                B  	
 t        j                  || t         j                  j                  |             j	                         
|j                         |j                         		
fd}t        j                  ||t        |             }t        |       }||fS )Nc                   t        j                   g       t        j                        }t        j                   g       t        j                        }t        j                  t        j
                   |       t        j                        |      }t        j                  ||      }t        j                  |      S r   )rQ   r   r{   r  r   rR  rand)	r:  seed_index_exproffset_index_exprrand_index_exprr  r   offset_loader
random_posseed_loaders	        re   r\  zphilox_rand.<locals>.inner_fn
  s     ,,{2DLLr):EKKH''NN:e,ekk:<M
 
 ||FE**rg   r]  )
r1   r  FlexibleLayoutcontiguous_stridesmake_indexerr_  r?   r   ri   r  )r   seedrx  r  rv   r   r\  random_values_nodeoffset_noder  r  r  s        `   @@@re   philox_randr  
  s     
,,T2	
 ln  ""$K&&(M+ #))Dz	 %T*K{**rg   c           	         t         j                  r\t        j                  t        j
                  t        j                  j                  t        j                  j                  | ||            S t        d      )Nz&should be handled in replace_random.py)r/   r  rp  r  rB   r   r1   r  r  native_dropoutr  r   )r   ptrains      re   r  r  6
  sY    $$T%8%8%@%@!QN
 	

 EFFrg   c                x   t         j                  s-| j                         t        j                  d      k(  sJ d       | j                          t        |      dk(  st        |d   t              rt        j                  j                  nt        j                  j                  }t        j                  || g|  | S )Nr   Tthis should be handled in decomps unless config.fallback_random or the device is CPUr   )r/   r  rn   r{   rv   r  r   rk   r5  r  
bernoulli_Tensorr1   InplaceBernoulliFallback)r   rt   r  s      re   r  r  A
  s    !!Q\\^u||E7J%J ^J IIK t9>ZQ7 	__## 
 Q66Hrg   c                    t         j                  s-| j                         t        j                  d      k(  sJ d       t        t        |       g| S )Nr   r  )r/   r  rn   r{   rv   r  r  )r   rt   s     re   bernoulli_pr  P
  sG    !!Q\\^u||E7J%J ^J eAh&&&rg   c                    t         r   r   )r  s    re   _foobarr  Y
  s    
rg   c                .    t         j                  d       y )Nz1using triton random, expect difference from eager)r  info)salts    re   _warn_triton_randomr  ^
  s    HH@Arg   c                 J    t        t        j                  j                         y r   )r  rR   r]   creation_timer   rg   re   warn_triton_randomr  c
  s    --.rg   )r  c                     |j                  d      t        | i |S t        j                  r|j	                  dd        t        | i |S t        d      N	generatorz-should have been handled in replace_random.py)r}  fallback_rand_generatorr/   r  r  fallback_rand_defaultr   rt   r   s     re   r  r  x
  sT    zz+*&777			

;%$d5f55
H
IIrg   c                     |j                  d      t        | i |S t        j                  r|j	                  dd        t        | i |S t        d      r  )r}  fallback_randn_generatorr/   r  r  fallback_randn_defaultr   r  s     re   randnr  
  sT    zz+*'888			

;%%t6v66
H
IIrg   c                l    t        j                  |      }t         j                  j                  | |      S r   )r1   get_stride_orderExternKernelrequire_stride_order)input_tensorr  stride_orders      re   inductor_force_stride_orderr  
  s)    &&v.L??//lKKrg   c                    t        d      )Nz.should be handled in fuse_seed_creation_pass()r  )rv   s    re   inductor_seedr  
  s    
I
JJrg   c                z    t                t        j                  t        j                  | t        |                  S r   )r  rB   r   r1   RandomSeedsrF   )r  rv   s     re   inductor_seedsr  
  s)    BNN5-2GHIIrg   c                |      fd}t        j                   j                          j                         |g       S )Nc                L    t        j                  j                               S r   )rQ   	load_seedget_name)r  r:  seedss    re   r\  z&inductor_lookup_seed.<locals>.inner_fn
  s    }}U^^-u55rg   r]  )r?   r   rn   r   )r   r:  r\  s   `` re   inductor_lookup_seedr  
  s9    6 !oo	 rg   c                V   t        | t        j                        st        j                  |       } | j                  dk(  rf| j                  }|t        j
                  j                         }t        j
                  j                  |      }|j                  |j                  z  }|S d}|}|S )Nrg  i   )
rk   r{   rv   r   r:  rg  current_deviceget_device_propertiesmulti_processor_countmax_threads_per_multi_processor)rv   rU  propthreads_per_round_CPU_GRAIN_SIZEs        re   get_threads_per_roundr
  
  s    fell+f%{{fll;**++-Czz//4&&)M)MM 	   +rg   )rx  align_dtypec                 
 t         j                  rJ dv sJ g | } t        j                  }|j	                         }t        j                  ||| t
        j                  j                  |       |      j                         
|j                         t         j                  r/|j                  dk(  r t        |      dd} ||      
fd}n
fd}t        j                  |||g |       }	|	j!                          |	S )	N)r  r  rx  rg  c                J    | t         j                  t         j                  fv ryy)Nr   r   )r{   rH  rG  )dts    re   _vec_from_dtypez(inductor_random.<locals>._vec_from_dtype
  s    emmU^^44rg   c           
          dg      } dg      }t        j                  ||t        j                   |       t        j                        t                    S )Nr   r.   )vec)rQ   
rand_eagerrR  r{   r  r   )r:  rng_seedbase_offsetr  r  r  r  s      re   r\  z!inductor_random.<locals>.inner_fn
  sS    "A3'H%qc*K>>!z%0%++>H rg   c                     t        t               g       t        j                   |       t        j                              S r   )r   rQ   rR  r{   r  )r:  moder  r  s    re   r\  z!inductor_random.<locals>.inner_fn
  s7    %73%Bz%0%++> rg   r]  )r  torch.dtyper   r   )r/   r  r{   r  rI  r1   r  r  r  r  r_  align_random_eagerr   r
  r?   r   r  )r   r  r  rx  r  r   rv   r  r\  r  r  r  r  r  s     `       @@@@re   inductor_randomr  
  s     %%%%$$$$T7DMME%%'FtR..AA$GPVln  ""$K  V[[F%:1&9	
 k*		 			 w	F NNMrg   r  c               f   	 t         j                  rJ g |}t        j                  }|j	                         }t        j                  |||t
        j                  j                  |      |      j                         |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   )rQ   	randint64rR  r{   r  r  )r:  highlowr  r  s    re   r\  z"inductor_randint.<locals>.inner_fn  sV    }}ONN:e,ekk:NN3,NN4-	
 	
rg   r]  )r/   r  r{   r  rI  r1   r  r  r  r  r_  r?   r   )
r  r  r   r  rx  r   rv   r\  r  r  s
   ``      @@re   inductor_randintr   
  s     %%%%T7DKKE%%'FtR..AA$GPVln  ""$K
 w	 rg   c                    | j                         }| j                         }t        d t        ||      D              |d   z   }| j	                         |d   ||d   fS )Nc              3  2   K   | ]  \  }}|d z
  |z    ywrH  r   )r   r  sts      re   r   z%_boundaries_helper.<locals>.<genexpr>  s     A/@eaa!er\/@s   rh  )r   ry  sumr  r  )tbr   r  
max_offsets       re   _boundaries_helperr'    s^    
 ;;=D]]_FAs4/@AAF2JNJ
Rr
	 rg   c                H    | j                         | j                         d   fS r=  )r  ry  r%  s    re   _sorter_helperr*  (  s    ;;="--/"---rg   	out_int32rightsidesorterc               J   
 d } |       r ||      r
9 |      s1 t        t        j                  j                  d       |||      S ||dk(  rd|rt        j
                  nt        j                  
|j                          j                          j                          t         j                               dk(  r

 fd}n	
 fd	}|j                         }t        j                  |
||j                  
      }	|	j                          |	S )Nc                ^    t         j                  j                  | t        j                        S r   )rR   r]   r  r  	BUCKETIZEr)  s    re   rC  zsearchsorted.<locals>.<lambda>6  s    AGG$7$7
N$$%rg   Fr  r+  r-  Tr.   c           
          |       }t        j                  |t              dd n
t              d       S d      S )Nr   r/  sorter_indicesrQ   	bucketizer'  r*  )rU  r$  index_dtyper-  sorted_sequencer/  values_loaders     re   r\  zsearchsorted.<locals>.inner_fn[  sV    $C=="?3%~t>&3I'-~t  <= rg   c                             }d fd}t        j                  |t               |      d n
t              d       S  |            S )Nc                    | j                         }t        j                  t        j                  t
        j                  d t        |d d d d       D                    S )Nc              3  ,   K   | ]  \  }}||z    y wr   r   )r   r  rs   s      re   r   zNsearchsorted.<locals>.inner_fn.<locals>.get_flattened_index.<locals>.<genexpr>r  s     &U9TAq1u9T   rh  )ry  rQ   rR  r  r  operatorr   r  )r%  stridesrU  r8  s     re   get_flattened_indexz;searchsorted.<locals>.inner_fn.<locals>.get_flattened_indexn  sU    --/~~$$ &UWSb\3sPR89T&U  	 rg   r4  )r%  rB   r6  )rU  r$  rA  r8  r-  r9  r/  r:  s   `  re   r\  zsearchsorted.<locals>.inner_fni  sh    $C =="?3#O4%~t>&3I'-~t  <Ov;V rg   r]  )r  r  searchsortedr  r{   r  r  r_  r  r   r   rn   r?   r   r?  )r9  selfr,  r-  r.  r/  validate_bucketizer\  rv   r  r8  r:  s   `  ` `    @@re   rB  rB  ,  s    /!$''9&'AT 1 1 8 8eT
 	
 DGO!*%++K$$&M 
?##%&!+
	 
		 	0 __Fzz	F NNMrg   r,  r-  c                  t        j                               dk(  sJ t        j                  j	                  | t
        j                        r.t        j                  j	                  t
        j                        s/ t        t        j                  j                  d      | |      S j                          | j                         }| j                         |rt        j                  nt        j                   fd}t#        j$                  ||| j                               }|j                          |S )Nr.   Fr  rE  c                \     |       }t        j                  |t              d      }|S r  )rQ   r7  r'  )r:  r$  r
  
boundariesr8  r  r-  s      re   r\  zbucketize.<locals>.inner_fn  s6    5!--z*
 rg   r]  )r   r   rR   r]   r  r  r2  r  r  r7  r  r  rn   r_  r{   r  r  r?   r   )	r  rH  r,  r-  rv   r\  r  r8  r  s	    ` `   @@re   r7  r7    s     z""$%*** 	
E>#;#;<GG
N,D,DEQ 5 55Q:%
 	
 F$$&L!*%++K
 ~~	F NNMrg   c                    t        | t        j                        xr+ t        | t        j                  t        j                  f       S r   )rk   r1   r9   NonTensorObjOpaqueMultiOutputr   s    re   _is_tensor_irnoderL    s:    a# J	BOOR112- ) rg   c                z    t        j                  t        t        j                  j
                  ||f      \  }}||fS r   )rp  tree_map_onlyrL  r1   r  require_stride1r  rt   r   s      re   require_denserQ    s7    ''2??::T6NLD& <rg   c                z    t        j                  t        t        j                  j
                  ||f      \  }}||fS r   )rp  rN  rL  r1   r  require_contiguousrP  s      re   rS  rS    s9    ''
**	vLD&
 <rg   c                z    t        j                  t        t        j                  j
                  ||f      \  }}||fS r   )rp  rN  rL  r1   r  r   rP  s      re   r   r     s;     ''
22	vLD&
 <rg   c                z    t        j                  t        t        j                  j
                  ||f      \  }}||fS r   )rp  rN  rL  r1   r  r2  rP  s      re   r2  r2    s9    ''
--	vLD&
 <rg   c           	        || S t        |t              st        |      r| S t        | t        j                        r.t        j
                  j                  | |j                               S t        | t              r!| D ci c]  }|t        | |   ||          c}S t        | t        t        f      r" t        |       d t        | |      D              S | S c c}w )Nc              3  :   K   | ]  \  }}t        ||        y wr   constrain_to_fake_tensor)r   r   f_as      re   r   z+constrain_to_fake_tensor.<locals>.<genexpr>  s!      
=O!S$Q,=OrI  )rk   r   r   r1   r9   r  require_exact_stridesr  r  rY  r   ri   r   r  )r   fake_argkeys      re   rY  rY    s    
(,-1J
#ryy!44S(//:KLL#tRUVRU3-c#hFFRUVV	C%	'tCy 
=@h=O
 
 	
 J Ws   >Cc           
         t        d t        | |      D              } |j                         D ci c]  \  }}|t        |||          }}}| |fS c c}}w )Nc              3  :   K   | ]  \  }}t        ||        y wr   rX  )r   r   r\  s      re   r   z,constrain_to_fake_tensors.<locals>.<genexpr>  s#      1MC 	!h/1rI  )r   r  r   rY  )rt   r   	fake_argsfake_kwargsr   r  s         re   r   r     sc      y1 D JPXAa)![^<<FX< Ys   Ac           
         fdt        fdt        || j                        D              }|j                         D ci c]  \  }}| || j                  |          }}}||fS c c}}w )Nc           	        t        |       r|j                  j                  d      }t        |t        j
                        s| S t        j                  |j                         t        j                  j                  j                        }t        j                  j                  | |      S t        | t              r| D ci c]  }| | |   ||          c}S | S c c}w Nr$  )rL  r|  r}  rk   r{   r  r1   r  r  rR   r]   r   ro  r  r  r  )r   fx_argfake_valr  r]  apply_constraints        re   rg  z1constrain_to_fx_strides.<locals>.apply_constraint  s    S!{{u-Hh5
..!177#3#3#=#=L ??77\JJc4 LOPCSC)#c(F3K@@CPP
 Qs   >Cc              3  6   K   | ]  \  }} ||        y wr   r   )r   r   re  rg  s      re   r   z*constrain_to_fx_strides.<locals>.<genexpr>  s"      9P+#vf%9Pr   )r   r  rt   r   r   )fx_nodert   r   r   r  rg  s        @re   r   r     st      9<T7<<9P D EKLLNSNDAqa!!W^^A%677NFS< Ts   !A*c                     fd fdt        fdt        t        | j                              D              }|j	                         D ci c]  \  }}| d| j
                  |          }}}||fS c c}}w )zGApply stride constraints to SDPA inputs, ensuring dense last dimension.c           
        t        |      s|S |j                  d   }|j                         D cg c]4  }t        |t        j
                        r|j                  j                  n|6 }}t        j                  j                  j                  }t        j                  |j                         |      }|r<|d   dk7  r4t        t        t!        t#        |j%                                                 }j&                  t(        j*                  j,                  k(  r| dv rt#        |      dk(  sJ d}d }t.        j0                  rd|j3                         }	t5        |      |	|rt7        |      nd f}|t        j                  j8                  v rt        j                  j8                  |   S  | ||||      }
||
t        j                  j8                  |<   |
S c c}w )Nr$  rh  r   )r   r   r   )r   r.   r,   r   )rL  r|  r  rk   r{   SymIntrc   exprrR   r]   r   ro  r1   r  ri   r(  r  r   r   ra   r  0_scaled_dot_product_efficient_attention_backwardr  r/   cache_sdpa_constraintmaybe_get_nameidr   sdpa_constraint_cache)rU  r   re  meta_valr  meta_stride_exprro  r  	cache_keyarg_namer  _apply_constraint_innerri  s              re   rg  z)sdpa_constraint.<locals>.apply_constraint,  s    %J;;u%GOGX
GX!:a6AFFKKA=GX 	 
 GG$$..	**8??+<iHL,1s3<<>/B)C DEL NNDDLLMv|$)))
 (L 	''))+H3'3l#I
 AGG999ww44Y??(h 0,
  7=AGG)))4Y
s   9Gc                ^   |j                   s,|j                  s t        j                  j	                  ||      S dj
                  t        j                  j                  j                  j                  k(  xr | dk(  }t        |t              sJ t        |j                               dvr|S t        j                  |      }|r3t        j                   t        j                  j#                  |      |      S t        |t$              rE|j'                         5|r3t        j                   t        j                  j#                  |      |      S |rt)        |j                               }g }|j'                         }	t+        t        |j                               dz
        D ]q  }
t,        j.                  j0                  j3                  ||
   d      s1|	3t,        j.                  j0                  j3                  |	|
   d      sa|j5                  |
       s dgt        |      z  }d|d<   d}t+        t        |      dz
  dd      D ]  }
||
dz      dk7  r|||
dz      z  }|
|v r@t,        j.                  j0                  j3                  t7        ||
dz            d      rd||
<   ]t,        j.                  j0                  j3                  t7        |      d      st9        |      z  }|||
<    t        j                  j;                  ||      S |r3t        j                   t        j                  j#                  |      |      S t        |t$              rE|j'                         5|r3t        j                   t        j                  j#                  |      |      S fd}t        |j<                  t        j>                        rQ ||      sI ||jA                               r3t        j                   t        j                  j#                  |      |      S t        j                  j	                  ||      S )	Nr   r   r   r   r.   r   rh  r,   c                    t         j                  j                  j                  t	        j
                  t        | j                         d         d            S )Nrh  r   )rR   r]   r   r  r   r  r*   r   )r   	ALIGNMENTs    re   
is_alignedzDsdpa_constraint.<locals>._apply_constraint_inner.<locals>.is_aligned  s@    77##22QZZ\"-y91= rg   )!is_cudais_xpur1   r  r  ra   r{   rQ   r  '_scaled_dot_product_efficient_attentionr  rk   rB   r   r   is_aligned_realized_tensortry_match_insignificant_stridesrealize_inputr9   maybe_get_strideri   r  rR   r]   r   rl  ro   r*   rE   r[  rm   r5   r  )rU  r   rs  rt  r  effn_attn_fwd_biasis_aligned_tensorrp  expanded_dimsmaybe_striders   out_stridesr  r|  r{  ri  s                 @re   rw  z0sdpa_constraint.<locals>._apply_constraint_inner_  s     HOO??77\JJ 	 NNyy~~EEMMN q 	 #y)))s||~f,J99#yI55--c24D 
 sF#$$&2!55--c24D  CLLN+HM //1L3s||~.2377##;;<LQ<OQRS ,((@@aRST!((+ 4 $X.KKOF3x=1,b"5q1u%*#hq1uo5F %ww''??KA.	:A *+A ww''??	*A %VY7)CF!'A% 6( ??88kJJ55--c24D 
 sF#$$&2!55--c24D 	
 chh,c?coo/0==55c:<L  33CFFrg   c              3  >   K   | ]  \  }\  }} |||        y wr   r   )r   rU  r   re  rg  s       re   r   z"sdpa_constraint.<locals>.<genexpr>  s+      "DC#v 	c6*"Ds   rh  )r   rj   r  rt   r   r   )ri  rt   r   r   r  rw  rg  s   `    @@re   sdpa_constraintr  )  s    1fdGL  "+Cgll,C"D D IOW1a!"a):;;FW< Xs   "A<)r  c                   t        |t        j                        s*t        || j	                         | j                               }|}| j                         |j                         k7  rt        || j                               }| j	                         |j	                         k7  rt        || j	                               }| j                         |j                         k7  r%t        || j                               }t        |      S t        |      S NrF  )rk   r1   r9   r@  r   rn   r  r   r   r-  r  )rC  rT  r  r   rq   s        re   r  r    s    c299%S 09JKACNN,,a*+~~3==?*Q()}}#,,.(Q(Sz8Org   )memory_formatc          	         t        j                  | j                         | j                         | j	                         t        | j                                     S Nr]  )r?   r   rn   r   r_  ri   r   )r   r  s     re   r  r    s@     ||~kkmAJJL!	 rg   c                   g }t        | t              rt        | j                  t        j                        rw| j                  } t        | t        j                        rF|j                  | j                                | j                  } t        | t        j                        rFt        |       } t        |       } |r;| j                  } |d d d   D ]  }t        j                  | |      }  t        |       } | S )Nrh  r  )rk   rB   rm   r1   r   ro   rw  r  )r   reinterpret_view_layoutsr  s      re   clone_preserve_reinterpret_viewr    s    !!YJqvvr7I7I$JFFB../$++ALLN;A B../ aLaAFF.tt4F""&9A 5aLHrg   r  c               \    fd}t        j                  t        |      || g      S )Nc                D    t        j                  | d   z  z         S )Nr   r   rQ   rR  )r:  r   r|  r~  s    re   r   ziota.<locals>.fn  s!    ~~dU1Xo5UCCrg   r]  )r?   r   rF   )lengthr|  r~  r   rv   requires_gradr   s    ```   re   iotar    s1    D V$x	 rg   )r   rv   r  
pin_memoryr  c               ^    |J t        || z
  |      }t        || |||||      S d|      S )Nr   r|  r~  r   rv   r  )rE   r  )	r|  r}  r~  r   rv   r  r  r  r  s	            re   arange_start_stepr    sV     S5[$'F+v# 
 27# rg   c           	        t        || j                               }| j                         t        | d      t        j
                  j                  j                  t        j                  d            r| j                            z   njt        j
                  j                  j                  t        j                  d            rn, t        t        j                  j                        | |      S t        j
                  j                  j!                  d       t        j
                  j                  j#                  | j                                   t%        t'        |      | j                               }|j                         fd}t)        j*                  | j-                         | j                         |t/        | j                                     S )Nr   c           	         t        j                  t        j                  t        j                  |    t        j
                        t        j                  t        j
                               |        |             S r   )rQ   r  eqrR  r{   r  )rU  r   r:  
src_loaderrN  s    re   r\  z select_scatter.<locals>.inner_fn  sX    yyFFs3x5uekk2 sOSM
 	
rg   r]  )r   r   r_  ri  rR   r]   r   r  r   r^  r   r\  r  r  select_scatterr  rj  rk  r-  r>  r?   r   rn   ri   )r   rT  r   r:  r\  r  rN  s     `` @@re   r  r    sT   
3
&C}}H
3
"Cww&&uxxq'9:

S))	
			(	(%);	< = 3 3 ; ;<QS%PPGGq%(GGeQZZ\#%67
3$ajjl
3C"J
 ||~kkmAJJL!	 rg   c           	     :   	
 t        | j                               } j                         
t         d       j	                            t
        j                  j                         \  t         j	                               }t        z
  dz
  z         |<   t        ||      }|j                         		 
fd}t        j                   j                          j                         |t         j	                                     S )Nr   r.   c                   	dk(  rk(  r
dk(  r |       S t        j                  |    t        j                        }t	        |       t        |    	z
  
      <   g }	dk7  rZ|j                  t        j                  |t        j                  t        j                  	      t        j                                     k7  rZ|j                  t        j                  |t        j                  t        j                        t        j                                     
dk7  rz|j                  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   re   rC  z1slice_scatter.<locals>.inner_fn.<locals>.<lambda>c  s
    Jw'rg   r  )rQ   rR  r{   r  ri   r(   ro   r  r   r-  r  r  r+   rS  r  r  r  r  r   r  )rU  r  r  src_valr  r   rn  r}  r  r|  r~  r   rN  s       @re   r\  zslice_scatter.<locals>.inner_fn>  s   A:#/daic?"..S5;;7s)C5 0$7A:KKNN5<<#6D (?KKNN5<<#4ekkB 19KKNN'C5(8!TBEKK LLEKK0	 t$/**' #A

 yySM
 	
rg   r]  )r   r   r_  ri  r   r1   ru  normalize_start_endri   r(   r-  r?   r   rn   )r   rT  r   r|  r}  r~  src_sizer\  rn  r  rN  s   ` ````  @@@re   slice_scatterr  /  s    
3
&C}}H
3
"Czz|C H 11!S%EJE3AJJL!HS5[D1H5t<HSM
h
C"J,
 ,
\ ||~kkmAJJL!	 rg   c                j    t        | t        t        f      rt        |       dkD  rt	        | d         S | S r  )rk   ri   r   r   _unwrapr   s    re   r  r  t  s-    !dE]#A
qt}Hrg   r   rv   r  r  c                   t        |d t        j                  fv d|        t        | d       t        t	               t
              rxs t        j                  nxs t        j                         g }t         t        j                        r fd}nt         t        t
        f      r fd}nt               dk(  s't         d   t        t
        f      rBt               dk  r4|j                  t        j                  t                             fd}n5t        j                  j!                  t        j"                   |            S t%        j&                  t)        |      ||	      S )
Nlayout=r  c                0    t        j                        S r   r  r:  rm   r   s    re   r\  ztensor.<locals>.inner_fn  s    >>$..rg   c                0    t        j                        S r   rQ   rS  r  s    re   r\  ztensor.<locals>.inner_fn  s    <<e,,rg   r   r   c                      fdt              dk(  rt        j                  d      S  dt                    S )Nc           	     j   | |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   )rQ   rS  r  r  rR  r{   r  )r|  r}  midbinary_searchrm   r   r:  s      re   r  z/tensor.<locals>.inner_fn.<locals>.binary_search  s    s{"{;!#<<UU;;U{q(50yyFFuQx=S%++6 "%-!#s+ rg   r   )r   rQ   rS  )r:  r  rm   r   s   `@re   r\  ztensor.<locals>.inner_fn  s7     4yA~||Au-- CI..rg   rF  r]  )r   r{   stridedrk   r  r   r  get_default_dtyper   r   r5  r   ro   IntegerrR   r]   add_tensor_constantr@  r?   r   rF   )rm   r   rv   r  r  r^  r\  s   ``     re   r@  r@  z  s   v$..'&0BC:~|,'$-%$2002!F$$	/ 
D5#,	'	- 
Ta:d1gs|<TaemmCI./	/( ww**LLU6:
 	
 V$	 rg   c                z    t        | t              r|t        | |      } |t        | |      } | S t	        | ||      S r  )rk   rB   r   r  r@  )rm   r   rv   s      re   	as_tensorr    sD    $	"D%(DT6*D$eF33rg   c                8    t        | t        j                        S r?  r  )rm   s    re   long_tensorr    s    $ekk**rg   c                   t        t        j                  j                  j                  t        j                  j
                  j                  d         }|J t        |      dk(  sJ |       t        t        |j                                     \  }}t        j                  |||       }t        j                  j                  |      |_        t        j                  j                  |       t        j                  j
                  j                  d   }t!        |t"        j$                  t"        j&                  t"        j(                  f      r|j*                  j,                  S t/        j0                  |      S )Nrd  r.   r$  )r%   rR   r]   r   ro  r^   r|  r   rF  rW  r   r1   DynamicScalarrs  r   rt  rk   r{   rl  SymFloatSymBoolrc   rm  r   rw  )rm   rd  binding_symr  r\  r$  s         re   _local_scalar_denser    s    2	""AGG$8$8$=$=>Q$R ((( !Q&9(99&   %6%<%<%> ?@Kk7D9F''))&1FKGGv& ''


#
#E
*C#ennemmDExx}}}}S!!rg   c                     y r   r   )rm   r   s     re   _assert_scalarr    s     rg   )rv   r  c                    y r   r   )r   r   r  r   rv   r  s         re   _assert_tensor_metadatar    s     rg   c                   | t        | t        t        f      st        d      rj                  t        t        t        f      rfd}nTt        t
        j                        rfd}n3t        j                               dk(  sJ j                         fd}t        j                  ||t        |            S )Nr   c                0    t        j                        S r   r  r:  r   r   s    re   r\  z_full.<locals>.inner_fn  s    <<u--rg   c                0    t        j                        S r   r  r  s    re   r\  z_full.<locals>.inner_fn  s    >>%//rg   r   c                     g       S r   r   )r:  value_loaders    re   r\  z_full.<locals>.inner_fn  s    ##rg   r]  )rk   r   r5  r   r   r   r   r   r   r_  r?   r   ri   )
fill_valuerv   r   r   r\  r   r  s     `  @@re   _fullr    s    Ej3,/GE74K%#u&	. 
E5;;	'	0 5>>#$)))((*	$ Dz	 rg   c                8     t        t        |            | fi |S r   create_tensor_liketensor_constructor)r   r  r   s      re   r)  r)    s     =0<=aJ6JJrg   c                "     d d d d dd d fd
}|S )NF)namesr   rv   r  r  r  c                   t        | d u d       t        |d t        j                  fv d|        t        | d       t        |      }|xs t        j                         }t        |      dk(  r6t        |d   t        t        t        j                  f      rt        |d         }|D ]  }t        |t        j                        sJ  |D cg c]  }t        j                  |       }}t        |||      S c c}w )Nnamed tensorsr  r  r.   r   )r   r{   r  rF   r  r   rk   ri   r   Sizerl  r   r-  r  )	r  r   rv   r  r  r  r   r  r  s	           re   r  z!tensor_constructor.<locals>.inner%  s     	5D=/26dEMM22gfX4FGz><0v&2002t9>ja4

2KLa>D A!!U\\222 )-.AQ.Z55 /s   C8r   )r  r  s   ` re   r  r  #  s#     6. Lrg   )r  r   r  rv   r  r  c                    t        | d u d       t        |      }t        |      dk(  r6t        |d   t        t
        t        j                  f      rt        |d         }t        |d ||||      S )Nr  r.   r   r   r  rv   r  )	r   rF   r   rk   ri   r   r{   r  empty_strided)r  r   r  rv   r  r  r   s          re   rO  rO  ?  sf     u}o.6"F
4yA~*T!WtUEJJ.GHT!W~d%v* rg   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   rv   r  r  r  c                  t        | d       t        |d t        j                  fv d|        || j                         }nt	        |      }|xs | j                         }t        | j                               } |||||      S )Nr  r  r  )r   r{   r  r   r   rn   ri   r   )r   r   rv   r  r  r  r   creation_fns          re   _constant_likez*create_tensor_like.<locals>._constant_likeW  s     	z><06dEMM22gfX4FG=KKME 'E)1<<>AJJL!fV

 	
rg   r   )r  r  s   ` re   r  r  R  s     duTX
 rg   c                *    t        t        |             S r   r  )r  s    re   r	  r	  i  s    0<==rg   c                     d d d d d fd
}|S )Nr  c               n   t        |t        t        f      sJ t        | d       t        |d t        j
                  fv d|        t        |      xs | j                         }|xs | j                         }|D cg c]  }t        j                  |       }}t        t        |      ||      S c c}w )Nr  r  )rk   ri   r   r   r{   r  r   r   rn   r   r  r  rF   )r   r   r   r  rv   r  r  r  s          re   _new_constantz#new_constant.<locals>._new_constants  s     $u...z><06dEMM22gfX4FGU#4q{{})1<<>*./$Qa $/Zv!6tDD 0s   =B2r   )r  r  s   ` re   new_constantr  r  s    tDT
E rg   r  c          	     ~    || j                         }|| j                         }t        |d ||t        |      |      S Nr  r   rn   r  rF   )r   r   r   r  rv   r  s         re   	new_emptyr    sF    }~V$ rg   c                  t        | t        t        f      sJ t        |t        t        t        d       f      sJ t	        |d t
        j                  fv d|        t        |      xs t        j                         }|xs t        j                  d      j                  }t        |      }t        d|||       }|j                          |j                  j                  }t        j                   |j                  dgt#        |       z        |_        t        |t$        j&                        sJ | D cg c]  }t)        j*                  |       } }|r#|D cg c]  }t)        j*                  |       c}nt$        j,                  j/                  |       }t%        j0                  ||| ||xs d      |_        |S c c}w c c}w )Nr  r  r   )r  rv   r   r   )r^  F)rv   r   r   r  	is_pinned)rk   ri   r   r   r   r{   r  r   r  r@  rv   rF   r  r  rm   dataclassesreplacer   r1   r  r   r-  r  r  r  r  )	r   r  r   r  rv   r  	pointwiser\  r  s	            re   r  r    s|    dT5M***ftUDJ7888v$..'&0BC<5#:#:#<E/u||C(//F6"F6TJI^^  F%%fkk1#D	/JFKfb//000%)*TELLOTD*  #))&Qa&)11$7 
 NN%FM  +)s   3F:F?c          	     ~    || j                         }|| j                         }t        ||||t        |      |      S r  r  )r   r   r  r   r  rv   r  s          re   new_empty_stridedr    sH     }~V$ rg   c                    t         j                  j                  j                  |      }t	        t        t        |            |j                        }t        j                  j                  | |      S )N)r]  )rR   r]   r   r   sortedr  r   __getitem__r1   r  r  )r   r  r  s      re   copy_stridedr    sN    WW55f=F%F,&2D2DEL??//<@@rg   c                V    |j                  d      J d        t        |      | fi |S )Nr   z(dtype should be handled by decomposition)r}  r  )r   r  r   s      re   fullr    s5    ::g*V,VV*)j)$9&99rg   c                   t        | t              sJ |j                         dk(  rt        | |j	                               S | j	                         t              dk(  }t        | |      |rt        | dg      } dg| j                         |j                         fd}t        j                  | j                         | j                         ||j	                               S )Nr   r.   c                    t        |       } t        j                   |                }t        |       dk(  r|g} n|| <    |       S r  )ri   rQ   indirect_indexingr   )rU  
gather_idxr   index_loaderr   rN  s     re   r   zgather.<locals>.fn  sL    3i**<+<d3iH
s8q=,C!CH}rg   r]  )rk   rB   	get_numelr  r   r   ri  r-  r_  r?   r   rn   r   )	r   r   r:  sparse_gradrx  r   r  r   rN  s	    `    @@@re   gatherr    s     a###AENN,--::<DY!^F
3
'C1qcNs}}H$$&L ||~kkm~~	 rg   c                0  	
 |r- t        t        j                  j                        | ||||      S |rJ t	        | t
              sJ t	        |t
              sJ dt        |j                               v sJ | j                         	|j                         t        |j                               | j                         
g |j                         
dd  	
fd}t        j                  | j                         | j                         |      S )Nr   r.   c                    t        |       t              k(  sJ |  d         | d        }t        j                  |d         gg | d  z   } |      S )Nz != r   )r   rQ   r  )rU  	var_index
weight_idxindices_loaderindices_ndimrP  weight_loaderweight_sizes      re   r   zembedding.<locals>.fn	  s{    3x3x=(@SEhZ*@@("3}#56	++I{1~FG K
K
 

 Z((rg   r]  )r  r  	embeddingr  rk   rB   strr   r_  r   r   r?   r   rn   )weightr
  padding_idxscale_grad_by_freqsparser   r  r  rP  r  r  s         @@@@@re   r  r    s   7 6 67G[*<f
 	
 :fi(((gy)))C))+,,,,&&(M((*Nw'')*L//#K6!!#6k!"o6H) )   " 	 rg   c           
        t        d | D              s)J d| D cg c]  }||j                          c}        t        d | D              rt        d      t	        |       D cg c]  \  }}t        |t              s| }}}t        |      dkD  sJ d       d gt        |       z  }t        |t        |D cg c]  }| |   	 c}       D ](  \  }}|j                         |k7  rt        d      |||<   * ||fS c c}w c c}}w c c}w )Nc              3     K   | ]T  }|P|j                         t        j                  t        j                  t        j                  t        j
                  fv  V y wr   )r   r{   r  r  r   r  r   rs   s     re   r   z.check_and_broadcast_indices.<locals>.<genexpr>  sA      A= 	
%++u{{EJJLLs   AAz)indices must be int64, byte or bool. Got c              3     K   | ]7  }||j                         t        j                  t        j                  fv  9 y wr   )r   r{   r   r  r  s     re   r   z.check_and_broadcast_indices.<locals>.<genexpr>!  s0      <Cqq}%**ekk22Gs   ?5?zFallback for bool indicesr   z"requires at least 1 non-None indexz.Fallback when indices is on a different device)r  r   r  r   rj   rk   rB   r   r  r   rn   )r
  rv   rs   r   
valid_idxsnew_indicess         re   check_and_broadcast_indicesr    s6      
 4G4eGqWXWdQ[[]G4e3fg   <C  ""=>> )' 2O 21jI6N! 2JOz?QD DD&3w<'KJ 1
3S
1GAJ
3S TU1 <<>V#%&VWWA V 
""# 5f P 4Ts   C=C='D D:Dc	           
     T   dt        j                        D ]  \  }	}
|
|	z
  dk7  sd t              D cg c]  \  }}|	| |    c}}g | t              t              z   d  d   }rz   nd | z   |d  z   f
d}|fS c c}}w )NFr.   Tr   c                  
 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  }|   *|t        |       k  sJ |j                  | |          |dz  }<|   }|J 
|   }|j                  t        j                   || |||z          |	              g || |d  }|S  |      S )Nr   rh  r.   r   wrap_neg)r   r  ro   rQ   r  )rU  r=  	new_indexfirst_tensor_indexstart_offsetnext_idxrs   loaderr   r   indexed_sizer
  indices_loadersnon_consecutive_tensorsoutput_sizetensor_indicestensor_sizer  rN  s            re   r   z*index_output_size_and_inner_fn.<locals>.fnW  sI   3x3{++++?#s<'8888;	+A.3q9K~b)A-.AL D qz!#c(***  X/A(+)))#A  ))s<,2EFG#!)	 /&

^
	 %,yE(92EErg   )r   pairwiserj   r   )r`  r
  r!  r"  r  r  rN  r   r  previouscurrentrs   r$  r  r   r  r   s    ````````      @@re   index_output_size_and_inner_fnr&  3  s    $ $&//?'X"&*# @ ,5W+=M+=C6!9+=MKSKS&[)9C<O)O)Q"RSK'*!K/ ++,,-./ 	 F  FD ?_ Ns   
B$B$c                    t        | ||      \  }}}t        j                  | j                         | j	                         ||      S r  )index_impl_helperr?   r   rn   r   )r   r
  r   r   r\  r  s         re   
index_implr)  |  sB    0GUCK1||~kkm	 rg   c                   t        |t        t        f      sJ | j                         t	        || j                               \  }}t        |      dkD  sJ d       |D cg c]  }||j                         nd  }}t        ||d      j                               }| j                         }t        t        |            D cg c]  }||   	||    }	}|rd|	v rd|vrt        d      t        t        |            D cg c]  }||   	 }	}t        ||||||	d ||	      \  }
fd}|
|fS c c}w c c}w c c}w )Nr   z Must have at least one valid idxz0index is out of bounds for dimension with size 0r  c                       |             S r   r   )rU  index_inner_fnrN  s    re   r\  z#index_impl_helper.<locals>.inner_fn  s    s+,,rg   )rk   ri   r   r_  r  rn   r   r   r  
IndexErrorr&  )r   r
  r   r  r!  rs   r  r"  r`  r  r   r\  r,  rN  s               @@re   r(  r(    sZ   ge}---}}H9'1<<>RG^~"F$FF"KRS7a!-q}}TA7OS w~a01::<=KZZ\F',S\':U':!gaj>TF1I':LUl"q';KLL',S\':;':!F1I':L;"@
#K- .009 T V <s   !D2
D7D7D<c                    	 t        | |d      S # t        $ r? | j                           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     re   r:  r:    sT    
!WD11 
			M

 1 1uMw
 	

s    AAAc                    t        | |d      S )NFr/  )r)  r0  s     re   _unsafe_indexr2    s    a..rg   c                6    t        t        |       |||dd      S )NTFr   may_realizeindex_put_impl_r  r   r
  r   
accumulates       re   	index_putr:    s     a'6:Tu rg   c                6    t        t        |       |||dd      S )NFr4  r6  r8  s       re   _unsafe_index_putr<    s     a'6:U rg   c                    |j                         | j                         k7  rt        || j                               }|rt        | |      }t        | t	        |d   ||             S r  )rn   r  r   rS  r  )rC  r
  r   r9  s       re   index_put_as_masked_fillr>    sV    T__..%!23D% T5UD9::rg   c                   ddl m} t        t        j                  t
        j                  j                  j                  j                        }t
        j                  j                  }t        j                  sF|D ||      r<d}|j                  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 : 
 )utilsr@  r   r  
index_put_rR   r]   r^   ra   _overloadnamer/   graph_partitionr|  r}  disable_cudagraphs_reasonr1   IndexPutFallback)	rC  r
  r   r9  r@  r  ri  r   rA  s	            re   index_put_fallbackrI    s    C$//177+?+?+F+F+T+TUK
 gg""G""8A\!,,**=$??;?E*;-8C,/)T7FJGKrg   c                $    t        | |||dd      S )NTr4  r7  rC  r
  r   r9  s       re   rD  rD    s    gvz4 rg   c                $    t        | |||dd      S )NFTr4  rK  rL  s       re   _unsafe_index_put_rN    s    gvzD rg   c           
        |rLd t        j                  |       |j                         v r$t        fd|D              s|j	                          |j                         dk(  rt        |      dk(  r|d   j                         t        j                  t        j                  fv r_|d   }t        t        |j                               t        | j                                     D ]  }t        |d      } t        | |g||      S t        j                         rt!        | |||      S |D ]E  }||j                         t        j                  t        j                  fv s7t!        | |||      c S  | j                         }	t        |	      }
|rLt#        | j                               r3|
dk(  rt%        | dg      } t!        | |||      } |
dk(  rt%        | g       } | S t'        || j                               }	 t)        || j+                               \  }}|D cg c]  }||j/                         nd  }}t1        | t2              sJ | j	                          |
dk(  rt%        | dg      } t5        ||d      j                               }t        t        |            D cg c]  }|	|   	 }}t7        |	|||||d |      \  }}t9        ||      }| j+                         }|J t        j:                  || j                         |j/                         |||rdnd       }t        j<                  d t        j>                  |       |	      }t@        jB                  jE                  |      |_#        t@        jB                  jI                  |       |
dk(  rt%        | g       } | S # t,        $ r t!        | |||      cY S w xY wc c}w c c}w )
Nc                   t        | t              rt        | j                  t        j                        r| j                  j                         } t        | t        j                        xr t        | j                  t        j                        xrg t        | j                  dd       xrN | j                  j                  j                  t        j                  j                  j                  j                  u S y)Nri  F)rk   rB   rm   r1   r5   r  r   r  r   ri  ra   r{   rQ   r  randpermr  )indices    re   indice_slice_from_randpermz3index_put_impl_.<locals>.indice_slice_from_randperm  s     &),FKK1U002vr}}5 V"6;;@VY=V ++22eiinn6M6M6U6UU	 rg   c              3  .   K   | ]  } |        y wr   r   )r   rR  rS  s     re   r   z"index_put_impl_.<locals>.<genexpr>  s      H
=D6&v.Wr	  r.   r   rh  r/  
atomic_addrv   r   r\  r^  output_indexerscatter_moder   r  rm   )%r1   try_get_namer  r  r  r  r   r   r{   r   r  r  r   r>  r>  $are_deterministic_algorithms_enabledrI  rK   r  r   r  rn   r   r_  rk   rB   ri   r&  r-  Scatterr  MutationLayoutSHOULDREMOVErR   r]   rs  r   rt  )rC  r
  r   r9  r   r5  r  r  r:  r`  x_ndimr!  rs   r  r"  r  expected_vals_sizer\  rv   scatterr\  rS  s                        @re   r7  r7     s   	 ??4 F$9$9$;;C H
=DH
 E
 NN 	aLAAJ  "uzz5;;&??qzs4==?+S-ABAT2&D C'tffjII 113!$DD !2uzz5;;6O!O%dGVZHH  ]]_F[FB4>>CSTQ;qc?D!$DQ;b>Dfdnn./FE"=T__&#
 LSS7a!-q}}TA7OSdI&&&LLN {D1# w~a01::<=K',S\':;':!F1I':L;#A	$  F./F __Fjjnn##%!%/\TG ,,T2F
 ''))&1FKGGv&{D"~Kc  E!$DDE T <s   3N N7N<N43N4r  c                   	 t        | |dd      \  }}j                         | j                         		fd}t        j                  | j	                         | j                         ||      S )NFr  c                     j                   t        j                  k7  r+t        j                          t        j                        }n        }t        j
                  | fd      S )Nc                                     S r   r   )_unsafe_index_fnrU  self_loaders   re   rC  z8_unsafe_masked_index.<locals>.inner_fn.<locals>.<lambda>  s    K8H8M,Nrg   )r   r{   r   rQ   r   r  )rU  mask_valrd  fillr  mask_loaderre  s   ` re   r\  z&_unsafe_masked_index.<locals>.inner_fn  sK    ::#||K$4ejjAH"3'Hzz($NPTUUrg   r]  )r(  r_  r?   r   rn   r   )
rC  r  r
  rg  r^  r  r\  rd  rh  re  s
    ` `   @@@re   _unsafe_masked_indexri    sw    "3gUU#FA ""$K""$KV V  nn	 rg   c           	         t        ||d      }| j                         }t        t        |            D cg c]#  }||   rt	        ||   ||    ||   dz
        nd % }}t        | ||d      S c c}w )Nr   r.   T)r9  )r  r   r  r   rg  r<  )r   r  r
  r   masked_valuer?  rs   clamped_indicess           re   #_unsafe_masked_index_put_accumulaterm    s    vq)LJJLE s7|$$A 7>ajgaj58)U1X\2dJ$   Q$OOs   (A-c                V    t        j                  |t        j                  ||             S r   )rQ   r  r  r   minmaxs      re   rg  rg    s    ;;sCKKQ/00rg   c                P    t        |       }t        ||||      }t        ||       |S r   )r  r{  copy_)rC  rT  r   r  re  r0  output_views          re   as_strided_scatterru    s*    4[FVT6>BK	+sMrg   c                2    t        t        |       |||fi |S r   )scatter_r  )r   r   r:  rT  r   s        re   r`  r`    s    E!Hc5#888rg   r  include_selfc               :   t        |t              }t        | ||j                         t	        t
        j                  |r|j                         n
t        |            |r|j                         j                  nd|      rt        j                  | ||||||       |S y )Nznot implrx  )rk   rB   rP   r   r
   r{   r   r   rn   r1   ScatterFallback)r  rC  r   r:  rT  r  ry  src_is_tensors           re   scatter_fallbackr}    s     sI.MU[[]#--/S	J!.J 	%	
 rg   r  c                  |dv sJ |[t        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  rw  rR   r]   r^   ra   rE  r}  scatter_reduce_)rC  r   r:  rT  r  r  fallback_results          re   rw  rw    s    ....~dmmQWW-A-A-H-H-V-VW*sE3v
 &""	:	4eS&99rg   c                0    t        t        |       |||      S r   )scatter_add_r  r   r   r:  rT  s       re   scatter_addr    s    a#uc22rg   c                     t        | |||d      S )Nr$  )r  r  s       re   r  r    s    1c5#u55rg   c                4    t        t        |       ||||fi |S r   )r  r  )r   r   r:  rT  reduction_typer   s         re   scatter_reducer    s    58S%nOOOrg   )ry  c          
         |dv sJ t        t        j                  j                               dk(  r dt        j                  j                         v sJ d       t	        t
              rt               t        t        j                  j                   |||      }|r|S t	         t              sJ dt        |j                               v sJ t         j                               }|dk(  rt         dg       t	        t              r)t        j                               dk(  rt        dg      t	        |t              r)t        |j                               dk(  rt        |dg      }|j                         dk(  r S t                j!                          |j#                         t	        t              rj#                         nd  fd} fd	}	d
 }
 j%                         }|J |st'        j(                  | j                          fd|j                         |d       }t'        j*                  d t'        j,                         |      }t.        j0                  j3                  |      |_        t.        j0                  j7                  |       t'        j(                  | j                         |	|j                         | |
|            }t'        j*                  d t'        j,                         |      }t.        j0                  j3                  |      |_        t.        j0                  j7                  |       |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_rx  r   r   c                    j                         }t        |      }t        |       }t        j                   |       |dk(  rdn|   d      |<   |S )Nr   r.   F)r  )r   r   ri   rQ   r  )rU  r?  rm  indirect_idxr   r  rC  s       re   rW  z'scatter_reduce_.<locals>.output_indexer&  sU    5zCy11DAIq5:
S rg   c                `    r |       S t        j                  j                               S r   rQ   rS  r   )rU  rC  rT  r  s    re   r   zscatter_reduce_.<locals>.fn0  s*    c?" <<T^^%566rg   c                    | dk(  ry| J y )Nr$  rU  r   r~  s    re   backend_reduce_strz+scatter_reduce_.<locals>.backend_reduce_str8  s    U? >!>rg   c                L    t        j                  dj                               S r  r  )r:  rC  s    re   rC  z!scatter_reduce_.<locals>.<lambda>H  s    3<<4>>3C#Drg   rV  rY  )r   r  r  r   rk   r    r)  r}  r  rB   r  r   r   r  r  ri  r  r_  rn   r1   r\  r  r]  rR   r]   rs  r   rt  )rC  r   r:  rT  r  ry  r  rm  rW  r   r  rv   zero_outr\  r`  r  r  s   `` `           @@re   r  r    s   BBBBD  **,-2T))3355U UU	6 #vc"&  !O dI&&&C)****t}}DqyD1##y!c#,,.&9Q&>3n%#ENN,<(=(BUQC A
c
"CLLN$$&L&0i&@"dJ7 __F::.."D>>#)
 ""006

 gg--f5	""6*
 jjnn~~%'/G ,,T2F
 ''))&1FKGGv&qyD"~Krg   c                T   | j                          | j                         | j                          d  | j                         d   }D cg c]+  }t        j                  j
                  j                  |      - c}t        |      k(  sJ |}t        |      D cg c]
  \  }}||z   c}}t        |      D ]  \  }}	|		d|	z  |<    fdfd}
t        j                  | j                         | j                         |
g ||      S c c}w c c}}w )Nr  c                   t        j                  | t        j                        } r8t        j                  | t        j
                  dt        j                              } t        j                  | t        j
                  |t        j                              } t        j                  | t        j                        } t        j                  | |d      S )N      ?Fr/  )
rQ   rR  r{   r  r   rS  rz  r   r  r  )r   r  r   exacts      re   scale_fnz$upsample_nearestnd.<locals>.scale_fn  s     NN1emm,3<<U]];<AGGAs||E5==9:LLEKK($$QE::rg   c                    |  d  }| d   } 
g |t        |      D cg c]  \  }}} 	|||       c}}}      S c c}}}w r   )r  )rU  r   r/  rs   r  r   i_sizes
inv_scalesr  r  rN  s         re   r   zupsample_nearestnd.<locals>.fn  sc    H!HWaW#aW:UV:UJAq$8Aq$':UVW
 	
Vs   Ar]  )realize_hintr_  r   rR   r]   r   r  r   r  rj   r?   r   rn   r   )r   r   scales_xr  r  batchrs   o_sizesor  r   r  r  r  rN  s      ``      @@@@re   upsample_nearestndr  m  s    NN}}HjjlA23GJJL1"E6=>gqww))!,g>Gx=AG$'$9:$9DAq!a%$9:Jh'5%KJqM (	;
 
 ||~kkm!!!	 9 ?
 ;s   0D'D$c                "    t        | ||fd      S )Nr.   r  r  r   r   r  s      re   upsample_nearest1dr    s    avi1==rg   c                $    t        | ||fdd      S )Nr.   Tr  r  r  r  s      re   _upsample_nearest_exact1dr    s    avi1DIIrg   c                $    t        | |||fd      S )Nr,   r  r  r   r   scales_hscales_ws       re   upsample_nearest2dr    s     ax.BaHHrg   c                &    t        | |||fdd      S )Nr,   Tr  r  r  s       re   _upsample_nearest_exact2dr    s     ax.BatTTrg   c                &    t        | ||||fd      S )Nr   r  r  r   r   scales_dr  r  s        re   upsample_nearest3dr    s     ax8.LPQRRrg   c                (    t        | ||||fdd      S )Nr   Tr  r  r  s        re   _upsample_nearest_exact3dr    s#     	;8X6!4 rg   c                ,     t         fd|D              S )Nc              3  J   K   | ]  }t        j                  |        y wr   r  )r   r   r   s     re   r   z$_create_constants.<locals>.<genexpr>  s     6Aa's    #)r   )r   rt   s   ` re   r  r    s    6666rg   c                    | j                         | j                         fd}t        j                  | j	                         | j                         |      S )Nc                    t        |       } t        |       t              k(  sJ D ]  }|   dz
  | |   z
  | |<     |       S r  )ri   r   )rU  r   r  r  rN  s     re   r  zrev.<locals>.loader  sR    3i3x3u:%%%Cc
Q#c(2CH  }rg   r]  )r_  r   r?   r   rn   r   )r   r  r  r  rN  s    ` @@re   revr    sM     }}HJJLE ||~kkm	 rg   c                   d } |       ryt        |      dk7  st        | j                               dk7  ry| j                          t        | t        j
                        rt        | j                  t        j                        rt        | j                  j                  t        j                        s>t        j                  rNt        | j                  j                  t        j                        r | j                  j                  j                  sy| j                          t	        j                  |       \  }}|j                  }|d   dk7  ry|d   dk7  s|d   dk7  s|d   dk7  ry|d   }|dk(  ry|d   }|j                   d   }	||	|z   k  ry| 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                  } | yt        | j                        }t        |      dk(  rD|d   j                  t        j                  j                  t        j                  j                  fv ryy)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.
        Tr.   r   F)rR   r]   r^   r   r_   r   ra   r  mmr  addmm)r^   r_   s     re   _padding_can_be_fusedz6inplace_constant_pad_nd.<locals>._padding_can_be_fused  sk     ww++l(()u:?uQxGGOOJJ3
  
 rg   Nr   r,   r.   r   r   F)r   r|  r}  rg  inductorinplace_padding)r   r   r  rk   r1   rB   rm   r   r  r/   can_inplace_pad_graph_inputInputBufferr   freeze_layoutr  r  r   rR   r]   buffer_to_padded_sizer{  rx  r  fill_r   )r   paddingr  r  r  r  r@  npadstride0rowsizebufnamepadded_size	resized_xsliced_xs                 re   inplace_constant_pad_ndr    s   (  7|qC

-2
 IIK
 q",,'!&&"--0qvv{{B$5$5622qvv{{BNN; vv{{OO((+IAvmmGqzQqzQ'!*/WQZ1_1:DqyajGkk!nG4ffkkG;;q>6;;q>D#89K-8AGG!!'*		I iQg7T>QVWH	(JZ*+q0+rg   c           	     0   t        d |D              sy| j                         }t        |      }t        t	        |ddd   |ddd               }d}d}t        |      D ]*  \  }\  }	}
|	dk7  r y|
dkD  r| y|dz
  |z
  }|
}$|
dk  s* y |y| j                         }||j                  dk(  ryt        |      }|||<   | j                         } t        |      |      } t        |      |||      }t        d   d	xx   dz  cc<   t        | |g|      S )
a7  Decompose right-pad into cat([x, fill], dim) and delegate to cat lowering.

    The cat lowering already has heuristics for choosing between pointwise_cat
    (fusion) and ConcatKernel (memory planning / zero-copy).  By routing through
    cat() we reuse those heuristics rather than duplicating them here.
    c              3  <   K   | ]  }t        |t                y wr   )rk   r   r   r  s     re   r   z_pad_as_cat.<locals>.<genexpr>O  s     37az!S!7r  Nr,   r.   r   r   rF  r  pad_rewritten_as_cat)r  r   r   ri   r  rj   rn   r   r   r   r  r   r  )r   r  r  r  rm  	pad_pairspad_dim
pad_amountrs   leftr-  rv   	pad_shaper   fill_value_typed
pad_tensors                   re   _pad_as_catr  E  sF    3733JJLEu:DS1wqt!t}56I GJ%i0=D%1919"QhlGJQY 1  \\^FfkkU2 UI#IgKKME+}U+J75#$45vJ Z/0A50:((rg   c                   t        |      dz  dk(  sJ t        d |D              rt        |       S t        j                  rt        | |      }|r|S t        | |      }||S | j                         }t        t        t        t        |d d d   |dd d                           t        |      t              z
  g D ]?  \  }}j                  t        j                  j                  j                  |      |f       A t        |d        }g t        |d        D ]C  \  \  }}	}
j                  |
       |j                  t!        j"                  |
|z   |	z                E t        |      t        |      k(  sJ  t%        | j'                                     fdfd}| j)                         t+        j,                  | j/                         | j'                         ||      S )Nr,   r   c              3  &   K   | ]	  }|d k(    ywr   Nr   r  s     re   r   z"constant_pad_nd.<locals>.<genexpr>}  s     
#7a167r  r.   c                B    g }t         	d        D ]J  \  }\  }}}|dk7  r|j                  t        |d             |dk7  s0|j                  t        ||             L t	        j
                  t        j                  |      }t        j                  | 
fd      S )Nr   c                             S r   r   )r:  rN  s   re   rC  z/constant_pad_nd.<locals>.mask.<locals>.<lambda>  s	    rg   )	r  ro   range_mask_lowrange_mask_highr  r  rQ   r  r  )r:  r  rU  r  r  r  boundsr  
mask_sizesr  rN  s   `     re   r  zconstant_pad_nd.<locals>.mask  s    (+E!"Ivz(J$C#tfaxN323qyOC89	 )K
 $/zz$ 7DDrg   c                    t        | d        }t        | d        D ]  \  }\  }}|j                  ||z
          t        |      t        |       k(  sJ  |      S r   )ri   r  ro   r   )r:  r  rU  r  _highbounds_precompr  r  s        re   	offset_fnz"constant_pad_nd.<locals>.offset_fn  se    rO	!$U12Y!?C#uS3Y' "@9~U+++Irg   r]  )r   r  r  r/   r  r  r  r   ri   r(  r  ro   rR   r]   r   lookup_precomputed_sizer   r-  r   r   r_  r?   r   rn   )r   r  r  rq   r  lhr   r  r  r   r  r  r  r  r  r  rN  s     `         @@@@@@re   r  r  z  s   L1"""

#7
##Qx%a*=J a*
-C

JJLE(4GCaCL'!$Q$- @ABCFE
S[ A 68N1qww//GGJANO  uRay/KJ qr3dT$5<<s
T(9:; 4 {s5z)))-q{{}-j9JE E }}H||~kkm	 rg   c                    t        j                  t        j                  | t        j                        t        j                  t        j                  |      t        j                              S r   )rQ   r  rR  r{   r  r   r  )rs   r  s     re   r  r    s@    66q%++&u}}S)5;;7 rg   c                    t        j                  t        j                  | t        j                        t        j                  |t        j                              S r   )rQ   r  rR  r{   r  )rs   r  s     re   r  r    s7    66q%++&tU[[) rg   c                V    t        j                  t        | |      t        | |            S r   )rQ   r  r  r  )rs   r  r  s      re   
range_maskr    s'    88q#4  rg   c                      j                          d   j                         xs dgz   fd}|S )Nr   c                N   | d   |  d  t        j                  t        j                  t	              D cg c]  }t        |   |   
|   z   
|          ! c}      }	rt        j                  |fd      S t        j                  |fd      S c c}w )Nc                 4     t               g       S )Nr   )constant_boundary_condition)r   ihpad_fill_valueprefixr   s   re   rC  z;constant_boundary_condition.<locals>.load.<locals>.<lambda>  s"    O3A~3O"fNrNrg   c                      g        S r   r   )r  r  rN  s   re   rC  z;constant_boundary_condition.<locals>.load.<locals>.<lambda>  s    (>V>b>*Brg   )r  r  rQ   r  r  r  r  )r:  rs   r  r  r  r   r  r  r  r  	padding_hr   rN  s      @@re   rm  z)constant_boundary_condition.<locals>.load  s    uC45\HHLQRUJWJqZ1qtil2Yq\MBJW
  JJ 
	
 D"BJO
	
 Xs   $B"
)r   r_  )	r   r  r  r  r   rm  r  r  rN  s	   ````` @@@re   r  r    sE     	


cTUA}}H$A39I
 
* Krg   dilationc               (   |dgt        |      z  }t        | d||   z  z   ||   ||   dz
  z  z
  ||   dz
  z   ||         }|rt        | d||   z  z   ||   ||   dz
  z  z
  d||   dz
  z  z   ||         }t        j                  j                  j                  t        j                  |dz
  ||   z  | z
  ||   z
  d            r|dz  }t        j                  j                  j                  t        j                  ||            rd}||fS |}||fS )Nr.   r,   r   F)	r   r(   rR   r]   r   r  r   r\  r  )	r   rs   kernel_sizer  r  	ceil_moder   x_outx_alts	            re   pooling_sizer    sN   3W%	A
NXa[KNQ,>??6!9q=Qq	E
 '!*nqk[^a/01 6!9q=!" 1I
 77**HHeai6!9,q071:=qA
 QJE77**588E5+ABI ) E)rg   c               l    t        | |      } t        j                  t        j                  |       }|dkD  S )N   )rL   r  r  r?  rz  )r  n_dimwindow_sizes      re   %should_fallback_max_pool_with_indicesr  	  s/    {E2K""8<<=Krg   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        | j	                               |dz   |dz   fv sJ t        ||      }|||k(  sJ |||||fS )Nr   r.   r,   r	  )rL   rk   rB   r   r   r  )r   r  r  r  r   r	  r  use_fallbacks           re   max_pool_checksr    s    !|#+1}3;{E2K&%(F7E*GHe,Ha###{u$$$v;%w<5   x=E!!!qzz|EAI 66668ERL"...<??rg   c                  | j                          | j                  d   }| j                   d  }t        t              D 	cg c]  }	t	        ||	   |	||       c}	 \  }
}| j
                  }|t        j                  u rdn5|j                  rt        d      nt        j                  |      j                  }t        |      t        |
      z   }t              st        |      st        d D              rt        | |      n| j                         fd}t!        j"                  d| | j%                         |||||      }t!        j"                  d	| | j%                         t        j&                  ||||      }t)        |j*                  j*                  t               r|j-                          t)        |j*                  j*                  t               r|j-                          ||fS c c}	w )
Nr  F-infc              3  &   K   | ]	  }|d kD    ywrH  r   r  s     re   r   z)_max_pool_with_offsets.<locals>.<genexpr>K  s     ,EHqQUHr  r  c                    | d   }|  d  }t              D cg c]  }||   	|   z  ||   |   z  z   |   z
  ! }} 
g ||      S c c}w r   r  )rU  reduction_idxr  bhrs   r  r   r	  r  r  rN  s         re   fn_innerz(_max_pool_with_offsets.<locals>.fn_innerP  s    Wuf%\ 5\
! UVAY=#3hqk#ABWQZO! 	 
 &2''	
s   $Arq  r  
input_noderv   	dst_dtyper  r\  r^  reduction_rangesargmax)r  r?  r  r  r  r   r{   r   rw  r5  r  rp  ri   r  r  r_  r@   r   rn   r  rk   rm   r  )r   r  r  r  r   r  r	  r  dhwr  dhw_outr   	min_valuerP  r  r  offsetsrN  s     ``` `          @re   _max_pool_with_offsetsr#  ,  s    NNGGGeVE
''5&'
C
 5\	

 " A;X "	

GY GGE EJJ 	$66eFmEKK<N<R<R  E{T']*H
7|s9~,EH,E)E.q)G==?( ( ||~$	F ||~++$	G &++""I.',,##Y/7?s

s   G8c           
        t        |      }t        | |||||d      \  }}}}}t        j                  d      5  t	        | ||||||      \  }}	|t        |	t        j                        fcd d d        S # 1 sw Y   y xY w)NFr  r  unroll_reductions_thresholdr  )r   r  r/   r   r#  r   r{   r  )
r   r  r  r  r   r  r	  r  r  r"  s
             re   !_low_memory_max_pool_with_offsetsr'  w  s     E 1@	1-K(A 
"	50
 x44 
6	5	5s   1A55A>c                P   t              | j                         t        j                  t	        j
                  t        j                              fd}t        j                  | j                         t        j                  || j                               }|S )Nc                     	|       }t        j                  |
      }t        j                  |      } | |      }t        j                  t        j
                  | d        t        j                        S r   )rQ   r  r0   _flattened_index_to_ndrR  _flatten_indexr{   r  )rU  rx  offset_sympyr  idhwincrements_to_index
input_sizer  r	  offsets_loaderr
  s        re   offsets_to_indicesz4_pool_offsets_to_indices.<locals>.offsets_to_indices  sm    $,,V[A&==lKX"36~~))$
E670CDekk
 	
rg   r]  )r   r_  r   rw  r  r  r?  rz  r?   r   rn   r{   r  r   )	r"  r  r/  r.  r1  r
  r	  r0  r
  s	    ```  @@@re   _pool_offsets_to_indicesr2    s     E((*N--	 0 0{ KLK
 
 !!#kk#!	G Nrg   c                L    t        |      fd}t        | |||      S )Nc                    |  d  }t              D cg c]  }||   |   z  ||   |   z  z   |   z
  ! c}S c c}w r   r  )rU  r  r  rs   r   r	  r  r  s       re   r.  zD_low_memory_max_pool_offsets_to_indices.<locals>.increments_to_index  sd    %\ 5\
! UVAY=#3hqk#ABWQZO!
 	
 
s   $<)r   r2  )r"  r  r/  r  r  r   r.  r	  s      ``` @re   '_low_memory_max_pool_offsets_to_indicesr5    s-     E
 $j*= rg   c           	         t        | |||||      \  }}}}}t        | ||||||      \  }}	t        |	|| j                  | d  |||      }
||
fS )Nr  )r  r#  r5  r?  )r   r  r  r  r   r  r	  r  rq   r"  r
  s              re   _max_pool_with_indicesr7    s}     1@	;1-K(A *	;9ELC 6	G <rg   c           	     &    t        | |||||d      S Nr,   r  r7  r   r  r  r  r   r  s         re   max_pool2d_with_indicesr<          "	;9A rg   c           	     &    t        | |||||d      S Nr   r  r:  r;  s         re   max_pool3d_with_indicesr@    r=  rg   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        |j                               dv sJ | j	                          | j                         }t        |t              rt        |j                  j                  t              r|j                  j                  }	|	j                         }
|
J t        j                  d t        j                  |
|	j                         |	j                               |	      }|j                          |j                         }n|j                         }|d uxr |d   dk(  xs |d uxr |d   dk(  }t        d |D              rt!        | ||||      S |j                         ^ }}| j                         ^ }|j#                         | j#                         t%        |j                               }t'        fdt)        d   dz        D              t'        fd	t)        d   dz        D              z  }|d
kD  rt!        | ||||      S |j                         fd}t        j*                  | j                         | j                         ||      }|rt        j,                  j/                  |      S |S )Nr   r.   r,   ry  )rv   r   r   rY  c              3  &   K   | ]	  }|d k7    ywrH  r   r  s     re   r   z3max_pool2d_with_indices_backward.<locals>.<genexpr>:  s     
$8a168r  c              3     K   | ]=  }t        t        |d          t        d t        |d    z
  d                z
  d       ? ywr   r.   Nrq  r(   r   r  r  r  s     re   r   z3max_pool2d_with_indices_backward.<locals>.<genexpr>G  M      *A 	HQq	"SHQQ5GPQ,S%TTVWX*   AAc              3     K   | ]=  }t        t        |d          t        dt        |d    z
  d                z
  d        ? ywr.   r   NrE  r   wr  r  s     re   r   z3max_pool2d_with_indices_backward.<locals>.<genexpr>K  rG  rH  r  c                   | ^ }}}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 ]3  }
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                  ||      }|	;t        j                  ||t        j                  dt        j                              }	t        j                  t        j                  t        j                   ||      t        j                   ||            |      }t        j                  |t        j                  |	|      |	      }	$ 6 |	J |	S )Nr   r.   Fr/  rh  r  )rQ   rR  r{   r  r(   r  rS  r  r  r   r  r  r  r  r  r  r  )rU  r  r  rL  
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                      re   r   z,max_pool2d_with_indices_backward.<locals>.fnZ  s3   A^^AIM5;;?

N
N..QQ'&)3VAY?
 ..QQ'&)3VAY?
 x6!9595;;Gx6!9595;;G++gs||Au{{'CD++gs||Au{{'CDE3>>-#MNE3>>,#LM'C]+WWWcll3&DEWWWcll3&DE))Bs||Au{{7S(TU$R(# ))Bs||Au{{7S(TU$R(#
  .j9'
3	|Z8#"yyy#,,sEMM*J H 88FF2u-FF2u- D  #yyswwx/KXVHE , (H ###rg   r]  )rk   rB   r   r   r  r  rm   r?   rn   r1   r  r  r   decide_layoutry  r  )fallback_max_pool2d_with_indices_backwardr_  ri   rq  r  r   r  r2  )grad_outputr   r  r  r  r   r  r
  	gO_striderm   rv   x_bufferx_strideis_channels_last_batch_heightr  rP  r
  r   rq   r\  r]  r  r^  r_  r`  ra  rb  s     ```                @@@@@@@@re    max_pool2d_with_indices_backwardrl    s    !|a&1}q6a###{q   v;!w<1x=Aqzz|&&& ,,.I!YJqvv{{I$Fvv{{"!!!$$$$nn&]]_
 
 	 &&(%%' ,A!1A 3)A,!"3  
$8
$$8K(Iw
 	
  jjlVWe&1&:&:&<#Q|((*N))+KAJJL!H {1~)* M  {1~)* M
  -/KR8K(Iw
 	
 ##%L9 9v 

%%'##%	C 44S99
rg   c                6    | j                         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                 ,     g  z   z         S r   r   )h_start_indexr  iwr  w_start_indexrN  s   re   rC  z3pad_adaptive_loader.<locals>.load.<locals>.<lambda>  s$    HNvN}r'9N=2;MNOrg   )rQ   r  r  rR  r{   r  r  )r  
incrementsstart_indicesend_indicesh_end_indexw_end_indexr  rp  r  rq  rr  pad_valrN  s   `      @@@@re   rm  z!pad_adaptive_loader.<locals>.load  s    B'4$}#. [xxFF}r15;;?{EKK8 FF}r15;;?{EKK8	
 zzOO
 	
rg   r_  )r   rx  rm  rN  s    ` @re   pad_adaptive_loaderrz    s    }}H
, Krg   c                    t        j                  | ||      }t        j                  |||      }t        j                  | ||      }t        j                  |||      }	||||	fS )N)out_diminp_dim)r  r#  )
r  r  h_inw_inh_outw_outrp  rv  rr  rw  s
             re    compute_indices_adaptive_poolingr    sa    %%k5$OM##IudKK%%k5$OM##IudKK+}kAArg   c                d    |\  }}|\  }}	t        | |||||	      \  fd}
|
S )Nc                   | ^ }}} |      } |      } |      } |      }d }	t        j                  t        d         t        d               D ]$  \  }
} |||
|g||g||g      }|	|}	 ||	      }	& |	S r  )r   productr  )rU  r  r  r  bwrp  rv  rr  rw  r  r  rq  r$  h_end_index_fnh_start_index_fnkernel_maxes
pooling_fnw_end_index_fnw_start_index_fns                re   r   z _adaptive_pooling_fn.<locals>.fn  s    R(,$R((,$R(''l1o(>lSTo@VWFBR.k*	C ~#C0 X rg   r  )r  r  r  in_sizes	out_sizesr  r~  r  r  r  r   r  r  r  r  s     `  `     @@@@re   _adaptive_pooling_fnr    sP     JD$LE5 	)YdE5	
 . Irg   c                h   
 |\  }|\  }}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                  ||	      ||
      }
|	|}	| ||	      }	 |
S r  )	r   r  r  rQ   rR  r{   r  r  gt)rU  r  r  r  r  rp  rv  rr  rw  maxvalmaxindexr  rq  r$  r:  r  r  r  r  r  r  r  s                  re   r   z)_adaptive_pooling_fn_with_idx.<locals>.fn  s   R(,$R((,$R(''l1o(>lSTo@VWFBR.k*	C NN#t+m;b@%++E  99SVVC%8%J~#C0) X, rg   r  )r  r  r  r  r  r  r~  r  r  r   r  r  r  r  r  s     `  `    @@@@@re   _adaptive_pooling_fn_with_idxr    sQ     JD$LE5 	)YdE5	
! !F Irg   c                6     j                         t        j                  k(  rt        d      t	         t
              sJ t        |      dk(  sJ  j                           j                         ^ }}}t        j                  j                  j                  |      }t        j                  j                  j                  |      }|\  }}||k(  r||k(  rt               S |dk(  s|dk(  r2g |||}t        | j                          j                               S ||z  dk(  r,||z  dk(  r$t!        ||      t!        ||      g}t#         |      S t%        ||z   dz
  |      }	t%        ||z   dz
  |      }
t'        |      ||gz   } j                         }|	|
z  }|dkD  rt)         |      S d }d }t+        |||	|
g||g||gt,        j.                  	      t1        t3                      fd
}t5        j6                   j                         |||      }|S )Nz0'adaptive_avg_pool2d' not implemented for 'Long'r,   r   rF  r.   r  c                     t        | |z  |      S r   r(   r:  r|  r}  s      re   r  z)_adaptive_avg_pool2d.<locals>.start_indexP      733rg   c                2    t        | dz   |z  |z   dz
  |      S r  r  r  s      re   r  z'_adaptive_avg_pool2d.<locals>.end_indexS  "    g-7!;WEErg   r  r  r  r  r  r  c                ^    t        j                   | t                     |             S r   )rQ   truedivrz  )rU  fn_sumones_loaderr   s    re   r   z _adaptive_avg_pool2d.<locals>.fna  s-    {{3+A./[1I
 	
rg   r]  )r   r{   r  r  rk   rB   r   r  r   rR   r]   r   r  r  rO  rn   r(   
avg_pool2drE   ri   fallback_adaptive_avg_pool2dr  rQ   r   rz  	ones_liker?   r   )r   r   r  r~  r  r  r  o_sizer  h_kernel_maxw_kernel_maxrP  r   r
  r  r  r   rvr  r  s   `                 @@re   _adaptive_avg_pool2dr  *  s   {{}#MNNa###{q   NNUD$77%%d+D77%%d+DLE5 u}QxzUaZ'5'%''V1;;=HHe|qTE\Q.e,htU.CD![))D5L1,u5LD5L1,u5LE{eU^+HKKME-KR+A{;;4F ""L1%.77F &il3K

 
		||~	
B Irg   c                     j                         t        j                  k(  rt        d      t	         t
              sJ t        |      dk(  sJ  j                           j                         ^ }}}t        j                  j                  j                  |      }t        j                  j                  j                  |      }|\  }}|dk(  s|dk(  r\g |||}t        | j                          j                               t        |t        j                   j                               fS ||z  dk(  r||z  dk(  rt        t!        ||z   dz
  |      }t!        ||z   dz
  |      }	t#        |      ||gz   }
 j                         }||	z  }|dkD  rt%         |      S d }d }t'        ||||	g||g||gt(        j*                  	      t-        ||||	g||g||gt(        j*                  	       fd
} fd}t/        j0                   j                         |||
      }t/        j0                   j                         t        j                  ||
      }||fS )Nz,adaptive_max_pool2d not implemented for Longr,   r   rF  r.   r  c                     t        | |z  |      S r   r  r  s      re   r  z(adaptive_max_pool2d.<locals>.start_index  r  rg   c                2    t        | dz   |z  |z   dz
  |      S r  r  r  s      re   r  z&adaptive_max_pool2d.<locals>.end_index  r  rg   r  c           	     <     | t        t        d                  S Nr  rz  r5  )rU  inner_func_max_valr   s    re   inner_fn_max_valz-adaptive_max_pool2d.<locals>.inner_fn_max_val      !#':1eFm'LMMrg   c           	     <     | t        t        d                  S r  r  )rU  inner_func_max_idxr   s    re   inner_fn_max_idxz-adaptive_max_pool2d.<locals>.inner_fn_max_idx  r  rg   r]  )r   r{   r  r  rk   rB   r   r  r   rR   r]   r   r  rO  rn   
ValueErrorrE   ri   fallback_adaptive_max_pool2dr  rQ   r  r  r?   r   )r   r   r  r~  r  r  r  r  r  r  rP  r   r
  r  r  r  r  r  rir  r  s   `                  @@re   adaptive_max_pool2dr  u  sL   {{}#IJJa###{q   NNUD$77%%d+D77%%d+DLE5zUaZ'5'%''V1;;=H%%++allnK
 
 	
 e|qTE\Q.D5L1,u5LD5L1,u5LE{eU^+HKKME-KR+A{;;4F ."L1%.;; 7"L1%.;;NN 
		||~!	
B 
		||~kk!	
B r6Mrg   c                h               j                          fd}|S )Nc                `   j                         }t        |      dk(  r[t        |       dk(  r d| d   dz
  
z
  g      }nJt        |       dk\  r | d   | d   dz
  
z
  g      }n$ dddz
  
z
  g      }n g | dz
  
z
        }t        j                  |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   rQ   rR  r   r{   r  r  r   float64r  r  r7  r  r  r   rw  )r  rs   samples_shapesamplei_exprdiffout_sz_exprrZ  seq_ir  r   in_sz	kernel_szndimsout_szsamplessamples_loaders             re   rm  z)_fractional_pooling_offsets.<locals>.load  s     ((*}"6{a (F1Iuqy3(GHV! (F1Iuqy3(OP (Auqy3(?@ $$>f$>eai#o$>?F7#4#4#67~~ei/=nnVaZ=LLu}}-s||K/W
 		#&&a0!U;		6F?e34syy%7PPUEKK0vvfk*$$SYYtUD%A5==QVCWXXrg   ry  )r  r  r  r  r   r  rm  r  s   `````` @re   _fractional_pooling_offsetsr    s@    C[F#JE#I((*N!Y !YF Krg   c                "    t        | |||d      S r9  _fractional_max_poolr   r  r   random_sampless       re   fractional_max_pool2dr        ;^STUUrg   c                "    t        | |||d      S r?  r  r  s       re   fractional_max_pool3dr    r  rg   c                   | j                          | j                  d   | j                   d  }}t        j                  d      5  t	              D cg c]  }t        |||||       c}| j                         fd}fdt        |      t        |      z   }	| j                         }
t        j                  d| | j                         |
|
||	|      }t        j                  d| | j                         t        j                  |
||	|      }t        |t              sJ |       t        |j                   j                   t              r|j#                          t        |t              sJ |       t        |j                   j                   t              r|j#                          t%        ||| j                        }||fcd d d        S c c}w # 1 sw Y   y xY w)	Nr  r%  )r  r  r  r  r  r   c                6    | d   } g | | |            S r   r   )rU  r  r  r.  r	  rN  s      re   r  z&_fractional_max_pool.<locals>.fn_inner  s.    5&\FOfO':3'NOPPrg   c                ~    | d   }|  d  }t              D cg c]  } |   |||         ||   z    c}S c c}w r   r  )rU  r  r  bdhwr  dhw_index_fnr	  s        re   r.  z1_fractional_max_pool.<locals>.increments_to_index  sc    5&\Fvw<D u%A  QQ0=3CC%  s   :rq  r  r  )r  r?  r/   r   r  r  r_  ri   r   r@   r   rn   r{   r  rk   rB   rm   r  r2  )r   r  r   r  r	  r  inp_dhwr  r  rP  r   r  r"  r
  r  r.  rN  s       `         @@@re   r  r    s   NNWWWuf%qwwvw'77E	"	5 5\

 " (&"% "

 ==?	Q	 ;k!22!! <<>(	
 ""#<<>kk(	
 &),4f4,fkk&&	2NN'9-6w6-gll''3OO*[!''+>
 w{ 
6	5

 
6	5s   GG1EGGG%c                ~     j                           j                         ^ }}}t        j                  j                  j                  |      }t        j                  j                  j                  |      }|^ }}}	||z  dk(  r,||	z  dk(  r$t         t        ||      t        ||	      gd      S t        ||      }
t        ||	      }d fd}t        ||
|g||g||	gt        j                         fd}t        j                   j                          j                         |t!        |            }|S )	Nr   r.   )divisor_overridec                F    t        | |z  t        j                  |            S r   )r'   r   rw  r  s      re   r  z0upsample_nearest2d_backward.<locals>.start_indexV  s    uwg(>??rg   c                     | dz   ||      S r  r   )r:  r|  r}  r  s      re   r  z.upsample_nearest2d_backward.<locals>.end_indexY  s    EAI99rg   r  c                (     | t                    S r   )rz  )rU  r  r   s    re   r   z'upsample_nearest2d_backward.<locals>.fne  s    c.q122rg   r]  )r  r   rR   r]   r   r  r  r(   rE   r  rQ   r   r?   r   rn   r   ri   )r   r   r/  r  r  rj  inp_hinp_wout_hout_wr  r  r  r   r  r  r  s   `              @@re   upsample_nearest2d_backwardr  A  s1    NNJJLVUEGG&&u-EGG&&u-E 'VUEu}eemq0&(>?RS
 	
 5%(L5%(L@: ""L1%.77F3 
		||~kkmJ
B Irg   c           
     (    t        | ||||||d      S )Nr,   r  _avg_poolndr   r  r  r  r  count_include_padr  s          re   r  r  s  )     		 	rg   c           
     (    t        | ||||||d      S )Nr   r  r  r  s          re   
avg_pool3dr    r  rg   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        | j	                               dz   dz   fv sJ | j                          | j	                         d   }| j	                          d  t        t              D 	cg c]  }	t        |	   |	|       c}	 \  }
}t              st        |      rt        | d      d}n| j                         d}t        |      t        |
      z   }| j                         }t        | t        j                   d      }fd	}t#        j$                  t&        j(                        }|d
kD  r6t        d t              D              rt*        dz
     } || ||      S |d
k  rt-        j.                  d
      nt1        j2                         }| j5                         }|J |5  t7        j8                  d| |||||      }d d d        t;        j<                  d      r4t        |j<                  j<                  t6              r|j?                          |r|r|r|n|}tA        ||      }n=fd}tC        j8                  | j5                         |||      }tA        ||      }tE        ||      S c c}	w # 1 sw Y   xY w)Nr   r.   r,   r  r  TF)r   r   c                    | d   }|  d  }|}t              D cg c]  }||   |   z  ||   z   |   z
   }} 	g ||      S c c}w r   r  )
rU  r  r  r  r  rs   r   r  r  rN  s
         re   r  z_avg_poolnd.<locals>.fn_inner  sr    Usd#Z>CCjIjbefQi"Q%''!*4jI&2'' Js   A	r  c              3     K   | ]D  \  }}t         j                  j                  j                  t	        j
                  ||             F y wr   )rR   r]   r   statically_known_truer   Ne)r   r   r  s      re   r   z_avg_poolnd.<locals>.<genexpr>  s;       ,DAq 	
..uxx1~>,s   A
Ar%  r$  r  rm   c                   |  d  }g }t              D ]  }||   |   z  |   z
  }t        j                  |
|   z   	|   |   z         }s/t        j                  |d      }t        j                  |	|         }t	        j
                  ||z
  t        j                        }|j                  |        t        j                  t        j                  |      S r  )r  r   r`  ra  rQ   rR  r{   r  ro   r  r  rz  )rU  r  divide_factorsrs   hstarthendfactorr  r   r  r  r  r  s          re   fn_countz_avg_poolnd.<locals>.fn_count  s    cTUBN3ZA*WQZ7yy+a.!8!A$:KL("YYvq1F 99T1Q40Dvu{{C%%f-   ##CGG^<<rg   r]  )#rL   rk   rB   r   r   r  r  r  r  r  r  r_  ri   r   r   r   rE  r  r  r?  rz  fallbacks_avg_poolndr/   r   
contextlibnullcontextrn   r@   r   r   rm   r  div_primr?   r   )r   r  r  r  r  r  r  r   r  rs   r  
ceil_modeshad_paddingrP  r   output_dtyper  r
  fallbackcontextrv   r  divisorr  r  divide_factorr  rN  s    ``` ` `                  @@re   r  r    s    #){C0K&#&F7C(Ga###{s"""v;#w<3qzz|q#' 2222NNJJL3$E	

cTUA 3Z

 1q+vw	J

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


*&6"K"g&	= 	= "((<<>	
 "m,FE""I

t 
s   )L2LLc                   dk7  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        |j                               dv sJ | j	                          |j                         ^ }t        d|      \  }	}
t        d|      \  }}| j                         d   xs d   xs |
xs || j                         ^ }t        |j                               }|j                         }t        fdt        d   dz        D              t        fdt        d   dz        D              z  }|dkD  rt        | ||      S fd	fd
}t        j                  | j                         |||      }|S )Nr   divisor must be not zeror,   ry  r.   c              3     K   | ]=  }t        t        |d          t        d t        |d    z
  d                z
  d       ? ywrD  rE  rF  s     re   r   z&avg_pool2d_backward.<locals>.<genexpr>X  rG  rH  c              3     K   | ]=  }t        t        |d          t        dt        |d    z
  d                z
  d        ? ywrJ  rE  rK  s     re   r   z&avg_pool2d_backward.<locals>.<genexpr>\  rG  rH  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.   )
rQ   rS  r{   r  r  rz  r  r   rR  r  )rW  rX  stride_hstride_wpad_hpad_wkernel_hkernel_wr  wstartr  wendr  heightr  r  r  rb  s                re   !compute_pool_size_without_paddingz>avg_pool2d_backward.<locals>.compute_pool_size_without_paddingo  s   
 <<q	5;;7<<q	5;;7WQZ5WQZ5<<A<<<A<X.6X.6{{GGFH%GGCNN65;;7?
 {{GGFH%GGCNN5%++6>
 VS\\!U[[%ABVS\\!U[[%AB{{4!DE{{4u{{!CDf 5swwtV7LMrg   c                8   | ^ }}}|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                  ||            }|;t        j                  ||t        j                  dt        j                               }t        j                  |t        j                  ||      |      } . |J |S )Nr   r.   Fr/  r  )rQ   rR  r(   r{   r  r  rS  r  r  r   r  r  r  r  r  r  r  )rU  r  r  rL  rP  rQ  rR  rS  rT  rU  rV  rW  rX  r  partr  r  r  r  r\  r]  r  r  r  r_  r`  r  ra  s                   re   r   zavg_pool2d_backward.<locals>.fn  s$   A
N
N..QQ'&)3VAY?
 ..QQ'&)3VAY?
 x6!9595;;Gx6!9595;;G++gs||Au{{'CD++gs||Au{{'CDE3>>-#MNE3>>,#LM'C]+WWWcll3&DEWWWcll3&DE#/,E&k'N[^;E=b"EE{{#11 #$&s||Au{{7S(T!" !.&+  11 #$&s||Au{{7S(T!" !-&+& ). xxFF2u%FF2u% #"yytS\\#u}}5UVH"yyswwx/FQHS , (V ###rg   r]  )rk   rB   r   r   r  r  r_  ri   r   rq  r  fallback_avg_pool2d_backwardr?   r   rn   )re  r   r  r  r  r  r  r  r  _h_out
ceil_mode1_w_out
ceil_mode2rP  r   r
  r   r  r  r\  r]  r  r  r_  r`  ra  rb  s     ``` ``          @@@@@@@@@re   avg_pool2d_backwardr  /  s    #'71'<X>XX<a&k9---a###{q   v;!w<1qzz|&&&

Q%;FJ &eQVWiXFJ))+K!*F
FjFJK&1&:&:&<#Q|AJJL!HKKME {1~)* M  {1~)* M
  -/KR+	
 		
 8? ?B 
		%%'	
B Irg   c                    dk7  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        |j                               dv sJ | j	                          |j                         ^ } t        d|      \  }	}
t        d|      \  }}t         d|      \  }}| j                         t              xs
 |
xs |xs || j                         ^ }t        |j                               }|j                         }fdt        d      D        \  z  z  }|d	kD  rt        | ||      S  fd
fd}t        j                  | j                         |||      }|S )Nr   r  )r   r   r   r   )r   r   r.   r,   c              3  j   K   | ])  t        fd t           dz        D               + yw)c           
   3  t   K   | ]/  }t        |   z  t        d |   z
     z        z
  d       1 ywrD  )rq  )r   r  rs   r  r  s     re   r   z0avg_pool3d_backward.<locals>.<genexpr>.<genexpr>  sF      
. VAYQ[^);q	(I!JJAN.s   58r,   N)rq  r  )r   rs   r  r  s    @re   r   z&avg_pool3d_backward.<locals>.<genexpr>  s=      3
 A	 	 
;q>A-.
 	
 s   /3}   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  d   K   | ](  }t        j                  |t        j                         * y wr   rQ   rS  r{   r  r   r  s     re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  s      'UfQ(Df   .0c              3  d   K   | ](  }t        j                  |t        j                         * y wr   r  r  s     re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  s      MWs||Au{{;Wr!  c              3  d   K   | ](  }t        j                  |t        j                         * y wr   r  )r   r   s     re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>   s#      (
2=QCLLEKK(+r!  c              3  x   K   | ]2  \  }}}t        j                  t        j                  ||      |       4 y wr   )rQ   r  rz  )r   r  r  pads       re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>$  s7      "
	1c GGCGGAqM3's   8:c           
   3     K   | ]i  \  }}}}t        j                  t        j                  ||      t        j                  t        j                  |t        j
                        |             k y wr   )rQ   r  r   rR  r{   r  )r   r|  r   r   r%  s        re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>*  sX      

'"q#s KKq!3773>>#u{{+KS#Q's   A/A1c              3     K   | ]<  }t        j                  |t        j                  d t        j                               > ywr  rQ   r  rS  r{   r  )r   r|  s     re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>5  s2      "
1 KKs||Au{{;<1   AAc              3     K   | ]?  \  }}t        j                  |t        j                  |t        j                               A y wr   rQ   r  rR  r{   r  )r   r}  r   s      re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>9  s6      
KS KKS^^C=>K   AA)r  rQ   rz  r  )pdrW  rX  stride_dr  r  pad_dr	  r
  kernel_dr  r  dstartr  r  dendr  r  r  depthr  r  r  r  rb  s                      re   r  z>avg_pool3d_backward.<locals>.compute_pool_size_without_padding  sH   'Uf'U$(HMWMue(
2=(
$(H"
 Rx8<ueU>S"


 '*(8X.&u%	'	

dD"
 &&1"

tT 2UFE4JK
dD GGCGGD&)3774+@A3774QWCX
 rg   c                .   | ^ }}}}d t        |||g      D        \  }}}d t        |||g!      D        \  }}}d t        |||g!      D        \  }}	}
d |||fD        \  }}}d t        ||	|
g g      D        \  }}	}
d }t              D ]{  }t              D ]i  }t        "      D ]W  }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                  ||
            }|;t        j                  ||t        j                  dt        j                              }-t        j                  |t        j                  ||      |      }Z l ~ |J |S )Nc              3  ,   K   | ]  \  }}||z     y wr   r   )r   r  r%  s      re   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>D  s     A)@vq#1s7)@r>  c              3     K   | ]<  \  }}}t        j                  t        ||z
  |z   |      t        j                         > y wr   rQ   rR  r(   r{   r  )r   r  r   r  s       re   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>F  s<      %
>1a NN8AEAIq15;;?>r)  c              3     K   | ]8  \  }}t        j                  t        ||      d z   t        j                         : ywrH  r7  )r   r  r  s      re   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>K  s6      
.1 NN8Aq>A-u{{;.s   >A c              3     K   | ]<  }t        j                  |t        j                  d t        j                               > ywr  r(  )r   pstarts     re   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>P  s2      %
5 KKQ <=5r)  c              3     K   | ]?  \  }}t        j                  |t        j                  |t        j                               A y wr   r+  )r   pend
pooled_dims      re   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>T  s;      
% j KKcnnZEF%r,  c              3     K   | ]?  \  }}t        j                  |t        j                  |t        j                               A y wr   )rQ   r   rS  r{   r  )r   r:  p_s      re   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>`  s:      "+JFB R(EF+r,  r   r.   r,   Fr/  r  )r  r  rQ   r  r  r  r  rS  r{   r  r  r  r  r  r   )#rU  r  r  r  rL  pdstartrP  rQ  pdendrR  rS  rT  pd_rU  rV  r-  rW  rX  r  r  r  r  r  d_window_sizer  r\  r]  r  r  r  pooled_depthr_  r`  r  ra  s#                        re   r   zavg_pool3d_backward.<locals>.fnB  s   AqAaAY)@A1a%
1ay+v>%
!'

Q1Iv.
ue
%
"GW5%
!'
$'u%m\'R%
ue 'C]+ /C"*-$gw7#sC+"JBB (3 0*+ +AQ ?+a. P A"b" M;;#!' # 5 5$'KK(*CGGE3<<5;;;W,X%& %1*/!" !$ 5 5$'KK(*CGGE3<<5;;;W,X%& %2*/!"  !$ 5 5$'KK(*CGGE3<<5;;;W,X%& %1*/!"!4 7D< 88E!2CFF2u4EFr5)D  '#&99 $S%--(H$ $'99T3778T3JH#Um 0 , (r ###rg   r]  )rk   rB   r   r   r  r  r_  r  ri   r   r  fallback_avg_pool3d_backwardr?   r   rn   )!re  r   r  r  r  r  r  r  rj  _d_outceil_mode_dr  ceil_mode_hr  ceil_mode_wr  rP  r   r
  r   r  r  rC  r3  r\  r]  r  r  rD  r_  r`  ra  rb  s!     ``` ``             @@@@@@@@@@@@re   avg_pool3d_backwardrJ    s    #'71'<X>XX<k9---a###{q   v;!w<1qzz|&&&$%JJL!VUFE&q+vw	FK ';FK 'q+vw	FK ))+Kg,K+KKK4?4H4H4J1Qm\AJJL!HKKME3
 q3/M=-  -/-?KS+	
 		
# #JU U Un 
		%%'	
B Irg   c                   | j                         }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  rFn t	        |      dk(  r	||   dk(  r_J  t	        t        |            t	        |      k(  sJ d       |S )Nr   )r   r  rh  zinvalid axis: r.   zreduction axis not unique)r   rk   r   r  r   r   ri   r&   )r   r  r   rs   s       re   _validate_reduction_axisrM    s    ::<D$vSY
4yA~T{//H>$1HH/	:D3t97Q;GCIs4y14GDG'c$i'CINtAw!|LL  z$ CI-J/JJ-Krg   )r  c               J   |t        | |      } | j                         t        t           t	        | |            }g }g g g t        t                    D ]Q  }||v r&j                  |       j                  |          -j                  |       |j                  |          S d|dv rt              dkD  rt        |       rt        | j                  t              rdnt        | j                  t        j                        sRt        | j                  t        j                        rat        | j                  j                  t        j                        r3| j!                         }|j#                         xs |j%                          fd}	r0t'              }
D ]  }t(        j*                  j,                  |
|<   ! n|}
| j/                         t1        | j3                         |xs | j5                         | j5                         |	|
      S )NFr  argminr.   Tc                |   t        |      t              k(  sJ 
r+t        |       t              k(  sJ D cg c]  }| |   	 } }t        |       t              k(  sJ d gt        |       t        |      z   z  }t        j                  t        |       t        |            D ]
  \  }}|||<     	|      }ru|D cg c]  }t	        j
                  |       }}|d   }t        dt        |            D ]  }||   z  ||   z   } |t        j                  |t        j                        fS |S c c}w c c}w r  )r   r   r   r  r   r-  r  rQ   rR  r{   r  )r:  reduction_indexrs   r  rU  varr   rindex
linear_idxinner_loaderkeepdimskept_idxreduced_idxreduced_sizesshould_compute_logical_indexr   s            re   r  z%_make_reduction_inner.<locals>.loader  sB   ?#s;'7777u:T***'/0x!U1XxE05zS]***Fc%j3+??@	!% #k?"C
HC !IcN
 Y' (/>?!ell1oF?  J1c&k*'-*::VAYF
 + 3>>*ekkBCC+ 1 @s   D4D9)rv   r  r  r\  r^  r  )r   r   r&   r   rM  r  r   ro   r:   rk   rm   r>   r1   r   r   Bufferrw  is_transposedis_contiguousri   r   r)  r*  r_  r  rn   r   )r   r  rW  r   rK  r  
kept_sizesrs   r  r  rP  rV  rX  rY  rZ  r[  r   s     `        @@@@@@re   _make_reduction_innerr`    s    Q::<Dc?3At<=DJHMK3t99q!  a)OOAd1g&  $) .."aLaffk*+/( 2 23qvvr}}-*QVV[["))2T\\^F$$&Df.B.B.D*D ) 6 :A''++HQK  ==?L||~'81;;=++-& rg   c                     dd d fd}|S )Nr   c               V   dv rF| j                         t        j                  k(  r%t        |       rt	        | t        j
                        } t        | |||      }t        j                  d| d|}t        |j                  j                  t              r|j                          |S )NrO  )r  rW  r   rK  r  )r  r  r   )r   r{   r   r:   r   r  r`  r@   r   rk   rm   r  )r   r  rW  r   r   r  rK  r  s         re   r  zmake_reduction.<locals>.inner  s    
 22+!EKK(A&"7)
 !!XAXQWXKK
 NNrg   r   r   )r  rK  r  s   `` re   make_reductionrc    s    T 4 Lrg   c                   |t        | |      } t        | |      }t        | j                         | j	                         f| j                         f| j                         |      S )N)rv   dtypes	inner_fnsr   r  )r   ri  r  rn   r   r_  r   )r   r  r   s      re   _make_scan_innerrg  $  sZ    QD!D||~==?$ZZ\ rg   r   c                   |t        | |      } | j                         t        | |      }| j                         }|t        j
                  t        j                  fv rt        | t        j                        } t        | ||      }t        fd|D              }t        j                  || j                         | j                               }t        j                  |t        |j                                     }t        t!        ||      |      S )Nc              3  (   K   | ]	  }|     y wr   r   r   rs   r   s     re   r   zmean.<locals>.<genexpr>=       04a$q'4r   r9  )r   r   rM  r   r{   rH  rG  r5  sum_rO   r1   r8   rn   r7   r   ri   div)r   r  keepdimr   r  
sum_resultdenomr   s          @re   r  r  2  s    Q::<D#At,D;;=Lu~~66Q$aw'J0400Ee1;;=XEeT**=*=*?%@AEC
E*L99rg   c                2  
 |d}| j                         
t        | |      }t        | |d      }|r|j                          t	        t        | |            }t        |||      }t        
fd|D              }|rt        j                  ||z
  d      }t        j                  || j                         | j                               }t        j                  |t!        |j                                     }t#        ||      }	|s|	fS |r|nt%        ||      }|	|fS )Nr.   T)rn  c              3  (   K   | ]	  }|     y wr   r   rj  s     re   r   z var_mean_sum_.<locals>.<genexpr>P  rk  r   r   r9  )r   rM  r  r  squarer  rl  rO   r   ra  r1   r8   r   rn   r7   r   ri   rm  r  )r   r  
correctionrn  return_meanx_meandiffsro  rp  x_varr   s             @re   var_mean_sum_ry  C  s    
::<D#At,D!T4(F3q&>"EeT7+J0400E		%*,a0e1;;=XEeT**=*=*?%@AE
E"ExVGFD$9F&=rg   c                l   t        | |      }t        | ||d d       }|d   }t        |d         }| j                         }|r|j                  dk(  st
        j                  }nt
        j                  j                  }t        |t        j                        xr t        |      |k  xr t        |      dk7  S )Nr  rW  r   rK  r^  r  r   r.   )rM  r`  rO   rn   r   r/   r&  cppuse_two_step_variance_thresholdrk   r   r  r   )r   r  rn  r   r^  reduction_numelrv   	thresholds           re   use_two_step_variancer  ]  s     $At,D"	wd$F HF#F+=$>?O\\^Fv{{e+66	 JJ>>	?EMM2 	' I-	'&!Q&rg   c                  dt        | ||d d       }|j                  d      }|j                  d       |j                  d       t        j                  j                  d|fd| j                         d|\  }}}	|j                          | j                         | j                         t        | |      }t        fd|D              d	 fd
}
 t        |
      |      }|r|j                          ||fS |fS )Nr.   r{  r\  r  r  welford_reduce)rf  r  r   c              3  (   K   | ]	  }|     y wr   r   rj  s     re   r   z$var_mean_welford_.<locals>.<genexpr>  s     1Dq47Dr   c                    t        | t        j                        rD| j                  s8t	        j
                  t	        j                  | t        j                        |      S t	        j                  | |      S r   )
rk   r   r   	is_numberrQ   r   rR  r{   r  rS  r;  s     re   get_constant_or_index_exprz5var_mean_welford_.<locals>.get_constant_or_index_expr  sG    a$Q[[<<q%++ >FF||Au%%rg   c                           }       }t        j                  d      }| t        j                  |||z
        z  S r  )rQ   rS  r  )rm   cNzerort  r   r  rnumels       re   r  z#var_mean_welford_.<locals>.scale_fn  sE    &z59&vu5||Au%ckk$A...rg   r   )r`  r  r1   WelfordReductionr   r   r  r   rM  rO   r  )r   r  rt  rn  ru  r   r  r  m2r  r  rS  r   r  r  r   s     `         @@@@re   var_mean_welford_r  u  s    
"	wd$F ZZ
#F
JJ{
JJ{%%,, )'kkm 	KD"a JJLKKME::<D#At,D1D11F&
/ #.
"2
&CDy6Mrg   c               ,   | j                         t              }t        | |d      } t        | ||||      }t        j
                  j                  st        | ||      rt        di |n
t        di |}t        fd|D              }|s|d   S |S )NFr  )r   r  rt  rn  ru  )r  rn  c              3  :   K   | ]  }t        |d         yw)Fr  N)r   )r   r   r  s     re   r   z#var_mean_helper_.<locals>.<genexpr>  s     Fv!8Ayu55vs   r   r   )r   r   r   r  r/   mtiadisable_welford_reductionr  ry  r  r   )	r   r  rt  rn  ru  r   r   r0  r  s	           @re   var_mean_helper_r    s    I))4MM.A
F KK11$QT7C 	
 ((  FvFFF'6!93V3rg   )rt  rn  c               "    t        | |||d      S )NFr  rt  rn  ru  r  r   r  rt  rn  s       re   var_r    s    	W% rg   c               "    t        | |||d      S )NTr  r  r  s       re   var_meanr    s    	W$ rg   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_recursiverQ   r  rS  rz  )r   r1  r   r  s       re   r  r    s    1uS^^A.E::Av||Au%%Av1a1fe,FWWVV$F	A!|#Mrg   c                .    t        j                  | |      S r   )rQ   r  r   r/  s     re   
pow_nativer    s    771a=rg   )r  c                z    t        t              r%j                         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              xr dcxk  xr dk  nc xs	 |xr dk\  }|rZ j                          fd}t        j                   j                          j                         | j                               S t         t              r: dk(  rt!        d      S  d	k(  r$t#        j                               rt%              S |rDt         t              rt'               S t        t              rt)               S t+               S t-               S )
Nr  r.   c              3  r   K   | ]/  }t        |t        j                        s|j                          1 y wr   )rk   r1   rB   r   r  s     re   r   zpow.<locals>.<genexpr>  s$     N1*Q2Ms   77i    r   c                F    t         |       j                               S r   )r  r   )rU  r   r/  r  s    re   r   zpow.<locals>.fn  s     a??rg   r]  r,   )rk   r5  r   r  r   sqrtr  rF  r   r_  r?   r   rn   r   r   r    r)  r   exp2fallback_pow_scalarfallback_pow_tensor_scalarfallback_pow_tensor_tensorr  )r   r/  r   is_integer_powembed_exponentr   r  s   ``    @re   r  r    sv   !U1c!f~	Au	!s(Aw	As	QQx NANNE%e,N  3' a"32AF  	@ <<>++-::<	
 	
 !V6Q?"6nQ[[]37Na &q!,,6"-a33-a33arg   c                   t        | t              r| j                  }n| }t        |t              r|j                  }t        |t        j                        st        j                  | j                         | j                         |j                         | j                               }t        |t        t        f      sJ |j                  }t        |t        j                        sJ t        |t        j                        rg|j                         sW|j                         sGt        |j                  t        j                        s#|j!                          |j                  |_        | S t        j"                  j%                  |||       | S )Nr]  unsafe_alias)rk   rB   rm   r1   r   r?   r   rn   r   r_  r   r5   r;   is_input_bufferis_module_buffer	NopKernelr  r]  realize_into)changedr$  r  changed_datarc   s        re   rS  rS    s*   '9%||#y!hhc2==)%%'##%__&##%	
 $: 6777ii#r}}---,.$$&((*l''6 	HH!!..\ /  Nrg   c                .    t        | t        | |            S r   )rS  r)  )r   r  s     re   r  r  8  s    Q	!Z011rg   c                    | |u r| S t        || j                               }t        || j                               }t	        || j                               }t        | |      S r   r  rn   r   r   r-  r   rS  )r  rT  r  s      re   rs  rs  =  sR    
cz

C)
*C
3
(C
clln
%CS#rg   c                .    t        j                  | |      S r   )rQ   floordivr  s     re   r  r  H      <<1rg   c                .    t        j                  | |      S r   )rQ   truncdivr  s     re   r  r  M  r  rg   c                .    t        j                  | |      S r   )rQ   div_rnr  s     re   _div_rnr  R  s    ::arg   c                >   t        |       xr t        |      }t        |       xr t        |      }|dk(  r,|rJ d       |rt        | |      S t        t	        | |            S |dk(  r,|rJ d       |rt        | |      S t        t        | |            S t        | |      S )Nr2  z5floordiv operands can not be boolean at the same timer7  z5truncdiv operands can not be boolean at the same time)r   r   r  r2  r  r  r7  rm  )r   r/  rounding_modeboth_integerboth_booleans        re   div_moder  W  s    "1%</!*<L"1%</!*<L X!XX
 ".x1~G5A3GGX!XX!-x1~C5Q3CCq!9rg   c                    t        |       xr t        |      }|rt        | |      S t        t        j                  j
                        } t        |      | |      S r   )r   logical_andr=   r  rz  r  r  )r   r/  	both_boolr   s       re   rz  rz  k  sN    "9q'9I1a  **+!~b!!Q''rg   c                f   t        | t        j                        rt        | j                        S t        | t        j
                        rt        | j                               S t        | t        j                        r| S t        | t        j                        syt        j                  j                  j                  | j                               }t        j                  |      5  t!        j"                  t        j$                  dd      5   | j&                  | j)                          }ddd       ddd       t        t        j                  j*                  j,                        sJ t        |j.                  t        j                        r|j.                  S y# 1 sw Y   rxY w# 1 sw Y   vxY w)z:Try convert an arbitrary IR node into an ir.Constant valueNallow_indexingT)rk   r1   r;   get_constant_valuerm   r5   r  r   Loopsr{   	_inductorops_handlerExtractConstantsHandlerrn   rR   set_ops_handlerr   objectr  r\  inner_fn_argsvirtualizedOpsValuer   )r   r  rq   s      re   r  r  u  s!   
 !R]]#!!&&))!R[[!!!--/22!R[[! a"oo))AA!,,.QG	'"R&&(8$?ajj!//+, 	@ 	#
 c5??66??@@@#))R[[)yy 	@? 	#"s$   &&F'F*F'F$	 F''F0c                l   t        d | |fD              }|rt        | |      S t        |      x}q| j                         j                  dk7  rT|j
                  dk(  r*t        j                  t        d      |j
                        }nd|j
                  z  }t        | |      S d } t        |      | |      S )Nc              3  L   K   | ]  }t        |      xs t        |        y wr   )r   r   r  s     re   r   zdiv_prim.<locals>.<genexpr>  s#     O1oa(>OA,>>s   "$r   r   infr  c                 &    t        j                  |  S r   )rQ   r  r  s    re   r   zdiv_prim.<locals>.fn  s    {{D!!rg   )r  r  r  rn   r   r   mathcopysignr5  rz  r  )r   r/  is_integralr   r  r   s         re   r  r    s    OAOOK1~ &a((5!,,.:M:MQV:V ==AuU|W]]CJw}},J1j!!" >"a##rg   c                Z    t        | |ft        j                        \  } }t        | |      S r  )rN  r   INT_TO_FLOATr  r  s     re   rm  rm    s0     	
A$C$P$PDAq Aq>rg   c                h    t        |       xs t        |       }|rd }nd } t        |      | |      S )Nc                .    t        j                  | |      S r   )rQ   modr  s     re   r   zfmod.<locals>.fn  s    771a= rg   c                .    t        j                  | |      S r   )rQ   fmodr  s     re   r   zfmod.<locals>.fn  s    88Aq>!rg   )r   r   r  )r   r/  r  r   s       re   r  r    s9    !!$:(:K	!
	" >"a##rg   c                   t        | j                               st        | j                               r|t        j                  }t        d|      } || |||      S )Nr$  r  r   r   r   r   r{   r  rc  r   r  rW  r   r   s        re   rl  rl    sK     	'+;AKKM+J
-	U	;Baxu--rg   c                   t        | j                               st        | j                               r|t        j                  }t        | j                               dk(  r(|dv sJ |xs | j                         }t        | |d      S d }t        | ||      }t        j                  j                  d	i |d|i\  }|t        | ||      S |S )
Nr   r   rh  Tr  c                @    | \  }|\  }t        j                  ||      fS r   )rQ   r   a_tupleb_tupler   r/  s       re   
combine_fnzcumsum.<locals>.combine_fn  #    1rg   r  r   r  r   r   r   )r   r   r   r{   r  r   r   r   rg  r1   Scanr   fallback_cumsumr   r  r   r  r   r  s         re   cumsumr    s     	'+;AKKM+J
-
1::<Aw&5t,, 
 ad%8F??J?IV~qd%88Mrg   c                   t        | j                               st        | j                               r|t        j                  }t        | j                               dk(  r(|dv sJ |xs | j                         }t        | |d      S d }t        | ||      }t        j                  j                  d	i |d|i\  }|t        | ||      S |S )
Nr   r  Tr  c                @    | \  }|\  }t        j                  ||      fS r   )rQ   rz  r  s       re   r  zcumprod.<locals>.combine_fn   r  rg   r  r  r  r   )r   r   r   r{   r  r   r   r   rg  r1   r  r   fallback_cumprodr  s         re   cumprodr    s     	'+;AKKM+J
-
1::<Aw&5t,, 
 ad%8F??J?IV~t599Mrg   c                
   d }| j                         }t        | j                               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                  ||      }||k7  t        j                  |       z  }t        j                  |t        j
                  t        j                  ||z
              |z   |      fS r   )rQ   r  r  r'  r  log1pexp)r  r  r   r/  min_vmax_vr  s          re   log_add_exp_helperz(logcumsumexp.<locals>.log_add_exp_helper  s{    Aq!Aq!CIIe$4#45		$		#''%%-*@ AE I1MOOrg   r   r  r  r  r  r   )	r   r   r   r  rg  r1   r  r   fallback_logcumsumexp)r   r   r  r   r   r  s         re   logcumsumexpr    s    P KKME
1::<Ag~~Qxac7FGG4FGIV~$QC00Mrg   c                   t        | j                               dk(  r,dv sJ t        |       t        | t        j
                        fS | j                         }t        j                  d|d      }t        | |      }|t        j
                  f|d<   | j                         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  re  c                R    t        j                  |    t        j                        S r   rQ   rR  r{   r  rU  r  s    re   rC  zcummax.<locals>.<lambda>1      CNN3t9ekk:rg   rf  r  r  r   )r   r   r  r  r{   r  r   r1   get_reduction_combine_fnrg  r_  r  r   fallback_cummaxr   r  r   r  r   r   r
  s    `     re   cummaxr	  "      
1::<AwQxAU[[999KKME,,5J ad%8Fu{{+F8	:F; ggnnEvE*EOFG~qd++7?rg   c                   t        | j                               dk(  r,dv sJ t        |       t        | t        j
                        fS | j                         }t        j                  d|d      }t        | |      }|t        j
                  f|d<   | j                         fd	f|d
<   t        j                  j                  di |d|i\  }}|t        |       S ||fS )Nr   r  r   rP  Fr   r  re  c                R    t        j                  |    t        j                        S r   r  r  s    re   rC  zcummin.<locals>.<lambda>H  r  rg   rf  r  r  r   )r   r   r  r  r{   r  r   r1   r  rg  r_  r  r   fallback_cumminr  s    `     re   cumminr  9  r
  rg   c                   t        | j                               st        | j                               r|t        j                  }t        d|      } || |||      S )Nr  r  r   r  r  s        re   r  r  P  sK     	'+;AKKM+J
-	e	<Baxu--rg   c                ^    t        | t        j                        }  t        d      | ||      S )Nr  r  rW  )r   r{   r   rc  r   r   rn  s      re   
reduce_anyr  [  s(    EJJA >% w??rg   c                Z    |t        | ||      t        | ||      fS t        | d |      S Nr  )reduce_amaxreduce_argmaxr  s      re   
reduce_maxr  a  :    
g6!#8
 	

 qtg66rg   c                Z    |t        | ||      t        | ||      fS t        | d |      S r  )reduce_aminreduce_argminr  s      re   
reduce_minr  l  r  rg   xor_sumrq  rp  r  r  rP  
logical_or)r  r  rv  stabler   
descendingc          	        |d}| j                         }| j                         }t        t        |      |      }t        |      dk(  r't	        |       t        d|t        j                  |      fS t        |      r||   nd}t        j                  j                  rt        j                  }nt        j                  }t        j                  j                  j!                  |t        j"                  |      j$                        st'        | |||      S t)        |dd||d      }dgt        |      z  }	t        |      r||	|<   t+        ||	      }t-        ||      }t.        j0                  j3                  || j4                  |j4                  f| j7                         |j7                         f||||      \  }
}|
t'        | |||      S |J |
t9        |t        j                        fS )NFr   r.   r   r  )rv   re  rf  r   r  r!  r"  )r   rn   r   r   r  r  r{   r  r/   tritondecompose_sort_opsr  int16rR   r]   r   statically_known_ltr  rq  sort_fallbackr  r  r-  r1   Sortr   r   r_  r   )r   r!  r   r"  r?  rv   rn  	idx_dtyper
  
view_shaper   s              re   sort_stabler,    s   ~JJLE\\^F
3u:s
+C
5zQQxq&%++u=== ZuSzQH }}''KK	KK	77//%++i:P:T:TUQv3:NN6QVG s5z!J
5z"
37J'GWe$Gggnn'==?G$7$7$9: % OFG ~Qv3:NN8GU[[111rg   c                     t        | d||      S )NFr   )r,  )r   r   r"  s      re   sortr.    s    qCJGGrg   c                @   t         j                  j                  st        |       S | j	                         }t        j                  t        j                  |t        j                  d            }t        | |g      }t        |d      \  }}|dz
  dz  }t        |d|      S )Nr.   r   r  r,   )r/   r$  r%  median_fallbackr   r  r  r?  rz  r   r  r  r,  r]  )rC  r   r  flatsorted_valsr  r   s          re   median_defaultr3    s    ==++t$$==?DX\\4q1ABEugD 1-NK	qA+q!$$rg   c                   t         j                  j                  st        | ||      S | j	                         }t        |      }|dk(  r5t        |       t        d| j                         t        j                  |      fS t        ||      }t        | d|      \  }}||   }|dz
  dz  }t        |||      }	t        |||      }
|rt        |	|      }	t        |
|      }
|	|
fS )Nr   Tr!  r   r.   r,   )r/   r$  r%  median_dim_fallbackr   r   r  r  rn   r{   r  r   r,  r]  r>  )rC  r   rn  r?  rm  r2  sorted_idxsr  r   r   r
  s              re   
median_dimr8    s    ==++"4g66MMOEu:DqyT{E!T__%6UKKK
4
%C*4#FKc
A	
Q1AKa(F[#q)G63'GS)7?rg   c                p    t         j                  j                  st        | ||      S | j	                         }t        |      }| j                         }|dk(  r't        |       t        d|t        j                  |      fS t        ||      }t        | d|      \  }}||   }t        |ddt        j                  |d      }	t        j                  d      g|z  }
||
|<   t!        |	|
      }	t#        |	|      }	|	j%                         fd}t'        j(                  t+        |      t        j                  ||      }t-        |||      }|j%                          |j%                         |	j%                          fd	}t'        j(                  t+        |      t        j.                  ||      }|j%                         |	j%                         fd
}t'        j(                  t+        |      t        j                  ||      }t1        ||      \  }}|	j%                         |j%                         fd}t'        j(                  t+        |      t        j                  ||      }t3        ||d      }t-        |||      }t-        |||      }|st5        ||      }t5        ||      }||fS )z9Lower aten.mode via sort-based decomposition or fallback.r   Tr5  r.   Fr  c           	         t        j                  t        j                   |       t        j                  dt        j
                              t        j                  dt        j
                              S Nr.   r   )rQ   r  r  rS  r{   r  )rU  positions_loader0s    re   prev_pos_fnz!mode_default.<locals>.prev_pos_fn  sI    {{GG%c*CLLEKK,HILLEKK(
 	
rg   r]  c           
         t        j                  t        j                   |        |             t        j                   |       t        j                  dt
        j                                    S r  )rQ   or_ner  rS  r{   r  )rU  positions_loadershifted_loadersorted_loaders    re   is_boundary_fnz$mode_default.<locals>.is_boundary_fn  sO    wwFF=%~c':;FF#C(#,,q%++*FG
 	
rg   c                    t        j                   |        |       t        j                  dt        j                              S r=  )rQ   r  rS  r{   r  )rU  is_boundary_loaderpositions_loader2s    re   boundary_pos_fnz%mode_default.<locals>.boundary_pos_fn.  s6    yys#c"LLU[[)
 	
rg   c                    t        j                  t        j                   |        |             t        j                  dt        j
                              S r  )rQ   r   r  rS  r{   r  )rU  last_boundary_loaderpositions_loader3s    re   
run_len_fnz mode_default.<locals>.run_len_fnC  s@    wwGG%c*,@,EFLLEKK(
 	
rg   r  )r/   r$  r%  mode_fallbackr   r   rn   r  r  r{   r  r   r,  r  r   r  r  r-  r_  r?   r   rF   r  r   r	  r  r  )!rC  r   rn  r?  rm  rv   r2  r7  r  	positionspos_view_shaper=  prev_positionsshifted_valsrD  is_boundaryrH  boundary_poslast_boundaryr  rL  run_lenmax_pos	mode_vals	mode_idxsrF  rJ  rA  r<  rG  rK  rB  rC  s!                            @@@@@@@@re   mode_defaultrY    s    ==++T300MMOEu:D__FqyT{E!VU[[%@@@
4
%C*4#FKc
A 	%++fEI mmA&'$.NN3Y/Iy%(I "--/
 %%V$kk	N +sN;L++-M!--/N ,,.
 ""V$jj	K %002!--/
 ##V$kk 	L lC0M1 "--/(446
 V$kk	G G#=G{C1I{C1IIs+	Is+	irg   c                z   t         j                  j                  st        | ||||      S | j	                         }t        |      }|dk(  r5t        |       t        d| j                         t        j                  |      fS t        ||      }t        | d||      \  }}t        ||d|      }	t        ||d|      }
|	|
fS )Nr   Tr   )r/   r$  r%  topk_fallbackr   r   r  r  rn   r{   r  r   r,  r  )rC  r   r   largestr  r?  rm  r2  r7  r   r
  s              re   topkr]  \  s    ==++T1c7F;;MMOEu:DqyT{E!T__%6UKKK
4
%C*Tsw K Ka+F[#q!,G7?rg   c                   t         j                  j                  st        | |||      S | j	                         }t        |      }|dk(  r5t        |       t        d| j                         t        j                  |      fS t        ||      }t        | d|      \  }}t        |||dz
        }t        |||dz
        }	|rt        ||      }t        |	|      }	||	fS )Nr   Tr5  r.   )r/   r$  r%  kthvalue_fallbackr   r   r  r  rn   r{   r  r   r,  r]  r>  )
rC  r   r   rn  r?  rm  r2  r7  r   r
  s
             re   kthvaluer`  m  s    ==++ q#w77MMOEu:DqyT{E!T__%6UKKK
4
%C*4#FKKa!e,F[#q1u-G63'GS)7?rg   c                <    t        | |t        j                  |      S )Nr   r   r  )r  r   r  )r`   r   r  s      re   register_pointwise_numericrc    s!    
;HH'	 rg   c                b    t        | j                         t        | t        j                        S r  )rN   r  r  r   r  r  s    re    register_pointwise_numeric_ldf64re    s'    '4
;HH rg   )r  )r   c                 	
 t        | ||t        j                        | j                         |j                         |j                         	| j	                         }j
                  xr1 t        j                  j                   xr |duxr |j                  dv 
	
fd}t        j                  | j	                         || j                               S )az  
    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.
    rP  Nrg  xpuc                    |       } |       } 	|       }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  )
rQ   rT  mul_rnrz  rk   r   r   rR  rS  r   )rU  self_valt1_valt2_valt1_times_t2
value_exprr   re  	t1_loader	t2_loaderuse_fmar   s         re   r\  zaddcmul.<locals>.inner_fn  s    s#33A:'7766844  **VV4K''&&1K eU[[)u5JeU3J77:{H== 778SWWZ%EFFrg   r]  )r   r   rE  r_  rn   rw  r{   rx  ry  r   r?   r   r   rC  tensor1tensor2r   rv   r\  r   re  rp  rq  rr  s      `  @@@@@re   addcmulrv    s      ;CC	E ""$K##%I##%I __F 	+!!!	+$	+ KK?*	 G G<  }}	 rg   c                 	
 t        | ||t        j                        | j                         |j                         |j                         	| j	                         }j
                  xr1 t        j                  j                   xr |duxr |j                  dv 
	
fd}t        j                  | j	                         || 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)

    For value=1: self + tensor1 / tensor2 (no FMA needed, just add the division)
    For value!=1: fma(value, div_rn(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.

    We use div_rn (round-to-nearest division) to force proper rounding, preventing
    Triton from fusing operations in ways that change the rounding behavior.
    rP  Nrg  c                    |       } |       } 	|       }
rt        j                  ||      }nt        j                  ||      }dk(  rt        j                  ||      S t	        t
        j                        rt        j                        }nt        j                        }
rt        j                  |||      S t        j                  |t        j                  ||            S r  )rQ   r  r  r   rk   r   r   rR  rS  rT  rz  )rU  rk  rl  rm  	t1_div_t2ro  r   re  rp  rq  rr  r   s         re   r\  zaddcdiv.<locals>.inner_fn  s    s#33 

662IFF3IA:778Y// eU[[)u5JeU3J77:y(;; 778SWWZ%CDDrg   r]  )r   r   r  r_  rn   rw  r{   rx  ry  r   r?   r   r   rs  s      `  @@@@@re   addcdivrz    s    " ;HH	E ""$K##%I##%I __F 	+!!!	+$	+ KK?*	 E E:  }}	 rg   r   r  logical_not)rv  )r   r  rK  identity)r  pointwise_overrides_datac              #  F  K   t         |   t        | j                  d       }|y fd}t        |t        j
                  j                        r9|j                         D ]%  }t        ||      }|j                   ||      f ' y |j                   ||      f y w)Nc                4    j                   t        |       S y r   )r$  r  )r`   rm   s    re   make_triton_fallbackz6_get_pointwise_overrides.<locals>.make_triton_fallback  s    ;;#B'' rg   )	r}  r   r   rk   r{   r|   r   r   r   )nsr   r`   r  olnamer  rm   s         @re   _get_pointwise_overridesr    s     #D)D	TYY	%B	z( "ejj112llnFV$Bd..0DR0HHH % $**,@,DDDs   BB!rb  c                d    | t         |<   t        j                  |        fd}t        | |       y )Nc                      | i |}g }t        | d   |      D ]"  \  }}|j                  t        ||d             $ |S )Nr   Tr  )r  ro   rS  )rt   r   resultsmut_resultsr   r  outplace_ops         re   r   z$register_foreach_inplace.<locals>.fn  sO    t.v.tAw0KCyf4HI 1 rg   )rZ   r  r   r  )aten_opoutplace_aten_opr  r   s     ` re   register_foreach_inplacer    s.    07,-G$ w+rg   c                2    t        | d       fd       }|S )NrP  c                 l     | i |}t        || d   j                               }t        | d   |      S r  )r   r   rS  )rt   r   r  r  s      re   r   zregister_inplace.<locals>.fn  s<    d-f-&$q'"3"3"56a&))rg   )r$  )r  r  r   s    ` re   register_inplacer    s#    wD9* :*
 Irg   c                     y r   r   ro  s      re   sym_constrain_ranger  7      rg   c                    t         j                  j                  j                  d   }t	        |t
        j                        r|j                  j                  S t        |      S rd  
rR   r]   r^   r|  rk   r{   rl  rc   rm  r   r   r   r$  s      re   r  r  <  A    
''


#
#E
*C#u||$xx}}3xrg   c                    t         j                  j                  j                  d   }t	        |t
        j                        r|j                  j                  S t        |      S rd  r  r  s      re   
sym_strider  E  r  rg   c                "    | j                         S r   )r  r   s    re   	sym_numelr  N  s    ;;=rg   c                 ~    t        |       dk(  rt        | d   t        t        f      r| d   } t	        j
                  |  S r;  )r   rk   ri   r   r   Addr  s    re   sym_sumr  W  s8     4yA~*T!WtUm<Aw99drg   c                    t        d      )NzHelpful for debuggingr   )rC  rt   r   s      re   foobarr  `  s    
5
66rg   c                8    | j                          t        |       S r   )r  r  r   s    re   _realizer  e  s    IIK8Org   c                R    | j                          t        j                  | |       | S r   )r  r1   ResizeStorageBytes)variablerP  s     re   resize_storage_bytes_r  k  s#    (H-Org   c                    | j                          |j                          t        j                  t        j                  | |            S r   )r  rB   r   r1   SetSourceTensorKernel)rC  source_tensors     re   set__source_tensorr  r  s5    LLNB44T=IJJrg   rs  c                    | |u r| S t        || j                               }t        || j                               }t	        || j                               }t        | |      S r   r  )r  rT  s     re   
fsdp_copy_r  {  sR    #:JS^^-.sCMMO,S#,,.)c""rg   c               8   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 | j                         | j                         }| j                         }t	        j                         rit        j                  j                   j"                  rEt%        |      rt'        d      n0t)        |      r t	        j*                  |      j,                  ndndt.        j0                  j2                  j5                  d      rt7        |||      S | j9                         }|d uxr t;        d	 |D              }|rt=        | g      }nRt        | j>                  t@        jB                        r| j>                  jE                         | _        tG        | gd
g      }|jI                         t@        jJ                  jM                  ||      }tA        jN                  ||||      jQ                         fd}	tS        jT                  |||	t        |            }
|
S )Nzunsupported memory format: r   r   nanTr  r   rF  c              3  p   K   | ].  }t         j                  j                  j                  |d        0 ywr  )rR   r]   r   rl  r   s     re   r   zresize.<locals>.<genexpr>  s+      2@G100A6s   46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   re   rC  z*resize.<locals>.inner_fn.<locals>.<lambda>  s    ZL(Arg   )rQ   rR  r{   r  r  r  )	rU  flat_index_exprlimitr  r  r  	old_numelout_indexeruninitialized_vals	       @re   r\  zresize.<locals>.inner_fn  sZ     %
..U[[Ay%++6vvou-zz$ ACTUUrg   r]  )+rk   rB   ri   r   r{   contiguous_formatpreserve_formatr  channels_lastr   channels_last_3dr  r   rI  r[  rC  deterministicfill_uninitialized_memoryr   r5  r   r  rq  rR   r]   r   rl  r  r  r  r  rm   r1   r5   r  r{  r_  r   stride_ordered_for_memory_formatr  r  r?   r   )r   r   r  r   rv   r@  has_overlappingx_flat
out_strider\  rq   r  r  r  r  s              @@@@re   resizer    s0   a###dT5M***//---8HII+++4yA~~...4yA~~IKKME""$F 	224KK%%??%  %ee$ %E 2 6 6 $  ww//	1=D+5HH  "GT) c 2@G2 /O a)% affbkk*VV'')AFA	{QC0$$&K""CCD-XJ..jANNPKV 

UXd4jC Jrg   )auto_functionalizedc                    ddl m} |j                  |      }t        j                  | ||i ||       |j                         D ci c]  \  }}t        |t              s|| c}}S c c}}w )Nr   )kernel_side_table)
kernel_idxgridtma_descriptor_metadatakernel_args)*torch._higher_order_ops.triton_kernel_wrapr  get_constant_argsr1   UserDefinedTritonKernelr   rk   rB   )	r  constant_args_idxr  r  r   r  constant_argsr]  r$  s	            re   triton_kernel_wrap_r    sn     M%778IJM 7/v//	 &,\\^R^cz#y7QCH^RRRs   A*!A*c                ^   t        d | g|D              rTd}t        j                  j                  j                  j                  dd       x}r| d| }|t        j                  _        t        j                  j                  | |||      }t        t        t        j                  |            S )Nc              3  V   K   | ]!  }t        |t              xr t        |       # y wr   rd  r  s     re   r   zcond.<locals>.<genexpr>  s&     
M;La:a 1Yq\1;Lrf  z"control flow operator: torch.cond.rA  rB  )r  rR   r]   r^   r|  r}  rG  r1   Conditionalr   ri   maprB   )predtrue_fnfalse_fnoperandsr   rA  r  s          re   r   r     s     
MD;L8;L
MM2''..3377tLL;LE*;-8C,/)^^""4(HEFI$$f-..rg   c                   t         j                  sit        d ||z   D              rTd}t        j                  j
                  j                  j                  dd       x}r| d| }|t        j                  _        t        j                  j                  | ||||      }t        |t              sJ t        t        t        j                  j                   |            S )Nc              3  V   K   | ]!  }t        |t              xr t        |       # y wr   rd  r  s     re   r   zwhile_loop.<locals>.<genexpr>  s+      *3A 	1f.)A,.3rf  z(control flow operator: torch.while_loop.rA  rB  )r/   rF  r  rR   r]   r^   r|  r}  rG  r1   	WhileLoopr   rk   r   ri   r  _maybe_wrap_as_tensor_box)cond_fnbody_fncarried_inputsadditional_inputsstack_outputr   rA  r  s           re   
while_loopr    s     !!c *"33* ' 9''..3377tLL;LE*;-8C,/)\\  .*;\F fh'''BLL::FCDDrg   )r  c                    t        j                  j                  | g| }t        t	        t
        j                  |            S r   )r1   InvokeSubgraphr   ri   r  rB   )subgraph_fn
identifierr  r  s       re   invoke_subgraphr     s5    %%k=H=FI$$f-..rg   c                l   t         }t        | j                  j                        D ]]  \  }}|j                  dk(  r?|t
        j                  j                  vsJ ||   t
        j                  j                  |<   U|j                  dk(  r]t
        j                  j                  |      \  }}t        j                  j                  j                  t
        j                  |||      }|t
        j                  j                  vsJ t
        j                  j                  }	 |t
        j                  _        t
        j                  j                  |      t
        j                  j                  |<   |t
        j                  _        ` |t         u rt        d      |S # |t
        j                  _        w xY w)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

    r  r0  zNo output node found in graph)_MISSINGrj   r]   nodesr`   rR   envfetch_args_kwargs_from_envr{   r&  Interpreterr0  r^   run_noder  )r  rt   r0  rs   rc   output_argsr   saved_current_nodes           re   process_subgraph_nodesr     s;    F\//556477m#qww{{*** $QAGGKKWW "#''"D"DT"JKXX))00$VTFqww{{***!"!5!5:'+$$%GG$4$4T$:D!'9$! 7$ :;;M (:$s    AFF3)control_depsc                   g }| D ]B  }t        |t              s|j                          |j                  |j	                                D t
        j                  j                  j                  }d}t        |      |z   t        |      k(  sJ t        t
        j                  j                        }t        |j                  j                  j                  d            t        |      k(  sJ t        |j                  t        |            }| sJ t
        j                  j                  |d D ]E  }	|D ]>  }
|	j                  }|J t
        j                  j                   |   j#                  |
       @ G |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)rk   r9   r  ro   r  rR   r]   r^   rt   r   
operationsr  
find_nodesr  ri   operation_nameadditional_buffer_depsr   )additional_depsr  rt   	dep_namesdeporiginal_args
arg_offsetoperation_lenr0  r`   dep_nameop_names               re   control_deps_op_loweringr  ;   s=    I#v&(  GG((--MJt9z!S%7777**+M{''--88M8JKsSWyXXX $K$<$<d4jIF? gg  0!H''G&&&GG**7377A " 1 Mrg   )schemec                  d }t         j                  j                  j                  j	                  dd       }|J t        | j                  j                  j                        D ]  \  }}|j                  dk(  r!||   t         j                  j                  |<   7|j                  dk(  rt         j                  j                  |      \  }}t        j                  ||j                               D ]  }	|	j                          |j                  r7t         j                  j                   j#                  |	j%                                t         j                  j&                  j#                  |	j%                                 t(        j*                  j,                  j/                  t         j                  |||      }Xt         j                  j1                  |      t         j                  j                  |<    |S )Nquant_optionsr  r0  )rR   r]   r^   r|  r}  rj   r  r  r`   r  r  r   r   r   r  codegen_low_precisionlow_precision_codegen_opsr   r  invoke_quant_opsr{   r&  r  r0  r  )
r  r 	  r  r0  r	  rs   rc   rt   r   r  s
             re   invoke_quant_tracerr	  j   sa   FGG((--11/4HM$$$[55;;AAB477m# (AGGKKWW 77==dCLD&__T6==?;		 66GG5599!:N:N:PQ((,,Q-A-A-CD < XX))00$fMF ! 0 0 6AGGKK% C( Mrg   c                  
 ddl m}m} t        |      dkD  rt	        d      t        j                  ||      D cg c](  } ||j                         |j                               * }} || |      

fd}t        |d   dd       }t        d |D              |d	<   t        d
 |D              |d<   t        j                  j                  d|dd|}	|	d   t	        d      |	S c c}w )Nr.   )InputDescriptorlower_pointwise_subgraphr   zSUnable to generate code for associative_scan op, because there are lifted argumentsrF  c                d     g t        j                  |       t        j                  |       S r   )rp  r  )lhsrhslowered_combine_fns     re   wrapped_combine_fnz,associative_scan.<locals>.wrapped_combine_fn   s6    ! 
$
$
 	
rg   r  c              3  <   K   | ]  }|j                           y wr   r   r  s     re   r   z#associative_scan.<locals>.<genexpr>   s     7BqQ[[]Br  re  c              3  <   K   | ]  }|j                           y wr   ry  r  s     re   r   z#associative_scan.<locals>.<genexpr>   s     <Ar  rf  F)r  can_fallback_to_atenz/Unable to generate code for associative_scan opr   )r  r	  r		  r   r  r   r   r   rn   rg  r   r1   r  r   )r  xsr  r	  r		  r   subgraph_inputsr	  r   r  r	  s             @re   associative_scanr	     s    M
!a
 	
 R((A 	akkmALLNC(   2*oN
 be!48F7B77F8<<<F;WW^^ %" F
 ayLMMM-s   -C'c                     y r   r   )tokenss    re   _sink_tokensr	     r  rg   c                      y r   r   r   rg   re   _make_tokenr	     r  rg   c                L   ddl m}m}  ||      }||t        j                  j
                  j                  u rddlm}m	} |j                         }	|	r|	j                  j                  t        j                  j
                  j                        }
|
rMt        |
|      sJ |
j                  |d         }|r)t        |      dk(  sJ d       t!        t#        |            }t        t$        j&                  j(                        }|t*        v r+t+        |   |i |}t-        j.                  t0        d |       n:d }t-        j2                  |t5        j6                  j8                  |g|i |      }t        t$        j&                  j(                  |d       dkD  sJ d	| d
       |rt$        j&                  j:                  j=                  |      }t$        j&                  j(                  |d D ]W  }d |_        |s|jA                         }t$        j&                  jB                  |   jE                  |jA                                Y t$        j&                  j:                  |<   	 d }t-        j2                  |||f      \  }} ||||      }t        |jR                        dk(  r| |fS t        |jR                        dk(  r| |fS | g|S # tF        $ rQ}tI        |      }tJ        jM                  d||       t        |tN        tP        f      r| g|cY d}~S | |fcY d}~S d}~ww xY w)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                "    | j                         S r   )r  r  s    re   rC  zwith_effects.<locals>.<lambda>   s
    !))+rg   c                Z    t        | t        j                        r| j                         S | S r   r  r   s    re   r  z"with_effects.<locals>.wrap_tensors   r  rg   zCNo operation nodes were generated when lowering effectful operator .c                      yr   r   r   rg   re   rC  zwith_effects.<locals>.<lambda>   s    drg   c                <   t        | t        j                        r| j                         S t        | t              rK	 | j
                  }t        |d      r0t        |j
                  d      r|j
                  j                         S | S | S # t        t        f$ r Y | S w xY w)Nrm   get_example)
rk   r1   TorchBindObject	get_valuerB   rm   r   r%	  AttributeErrorr   )r   r  s     re   convert_ir_to_valuez)with_effects.<locals>.convert_ir_to_value   s    !R//0{{}$Ay)ffGw/Gm5  '||7799 H	 '(;< s   AB BBz5Failed to get schema for %s: %s. Assuming list output)*torch._higher_order_ops.effectsr	  r	  r{   rQ   higher_orderr  torch._guardsr	  r	  try_gethop_dispatch_set_cache	get_cacherk   get_effectsr   rF  rW  rR   r]   r  rW   rp  rN  rB   r  r1   r  r   effectful_opsr}  has_side_effectsr  additional_star_depsr   r  r  r  r  r   ri   returns)tokenr`   rt   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                           re   with_effectsrA	     s    I b/KrUYY%;%;%K%KKE$,,.$/$F$F$P$P		&&66%! %!"79LMMM/;;DGDw<1,D.DD,"&tG}"5K **+M 
Y2//Y(=vF	L "++222GGG
 qww!!-.12Q6 
MbTQRS6 WW2266{Cgg((8F&2F#! //+,,W599:L:U:U:WX 9  	
k* #	$ &,__$&
"] Rm< 6>>av	V^^		!v  #F	CR	
 fudm,#F##6?"#s*   #(K	 		L#=LL#LL#L#)register_comm_loweringsregister_symm_mem_loweringsc                N   t        | |ddd      }|d   }t        j                  j                  j	                  t        |            }t        j                  j                  di |d|d\  }}|dk(  rXt        j                  j                  j                  |t        j                        r t        j                  d| d|d	|\  }}||fS 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_hintzTOnline softmax is disabled on the fly since Inductor decides to split the reduction.)rW  r   )r`  rR   r]   r   simplifyrO   r1   r@   
num_splitsstatically_known_geqr/   r&  r<   r   r  debugr  rW   r  r  r  rl  )r   r   r   r  r  hint	num_split
max_tensor
sum_tensorr  r  xsums               re   prepare_softmax_onlinerQ	  #!  s'   
 #	d$dF 01WW&&}5E'FGFll-- 
.OD) A~!''**??22 "8!>!> "
Qt"
?E"

J :%% 			b	
 1cD1!#a,/Ct,Tzrg   c                     t         j                  j                         xr! t         j                  j                         dk\  S )z.Check if we're on SM100+ hardware (Blackwell).)r   r   )r{   rg  is_availableget_device_capabilityr   rg   re   _is_sm100_or_laterrU	  Q!  s,    ::""$V)I)I)Kw)VVrg   c                   t               st        d      | j                         }|t        j                  t        j
                  t        j                  fvrt        d|       |t        j                  k7  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)rU	  r   r   r{   r  rH  rG  r  r   r  r#  rQ   inline_asm_elementwiseuint16r  r  )r   r   r   r  s       re   cvt_e8m0_rceil_loweringr]	  V!  s     !T
 	
 MMOEU]]EMM5>>BBOPUwW
 	

 sEMM*			""7ll
B  ^B$FFEKK((rg   )rY	  rZ	  c           	     <   	 t        | }t        d |D              |D cg c]  }|j                          c}	 	fd}t        j                  j                  |d   j                         |t        |d   j                                     S c c}w )Nc              3  <   K   | ]  }|j                           y wr   r	  re  s     re   r   z/lower_inline_asm_elementwise.<locals>.<genexpr>!  s     ;FSFr  c           
          t         fdD              }t        j                  |	d}t        j                  t        j
                  fv r$t        j                  |t        j                        }|S )Nc              3  .   K   | ]  } |        y wr   r   )r   r  rU  s     re   r   zAlower_inline_asm_elementwise.<locals>.inner_fn.<locals>.<genexpr>!  s     7wVVC[wr	  )rW	  rX	  r   rY	  rZ	  input_dtypes)r   rQ   r[	  r{   rH  rG  r   r  )
rU  valsr  asm_strrX	  r   rb	  rY	  rt  rZ	  s
   `  re   r\  z.lower_inline_asm_elementwise.<locals>.inner_fn!  sh    7w77++#%
 U]]ENN33\\&%--8Frg   r   r]  )	r   r   r_  r1   r?   r   rn   ri   r   )
rd	  rX	  r   rY	  rZ	  rJ  r   r\  rb	  rt  s
   `````   @@re   lower_inline_asm_elementwisere	  |!  s     'F;F;;L,23FSs F3G $ <<ay##%F1I&&()	   ) 4s   B)r  )quantized_lowerings)mkldnn_lowerings)jagged_loweringsc              #  X  K   t        | t        j                  j                        sJ d       t        j                  |       }	  t        |       t        |              d |r
|t        | <   yt        j                  |        y# |r
|t        | <   w t        j                  |        w xY w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)	rk   r{   r|   r}   rW   r}  r$  r  r  )r`   old_handlers     re   force_fallbackrk	  !  s      b%**//0 50 --#K".r23'IbMMM" 'IbMMM"s   AB*B ""B*#B''B*r/  )rp   zIterable[Any]r   z4defaultdict[tuple[Any, bool], list[tuple[int, Any]]])r   Callable[..., Any]r   zCallable[..., Any] | None)r   ztorch._C.Tagr   z%Callable[..., tuple[Any, Any]] | None)r   r   r   r  r   None)r   zuCollection[torch._ops.OpOverload | torch._ops.OpOverloadPacket] | torch._ops.OpOverload | torch._ops.OpOverloadPacketr   zlist[Any] | None)r   z3torch._ops.OpOverloadPacket | torch._ops.OpOverloadr   zCallable[..., tuple[Any, Any]]r   rm	  )r   zint | torch.dtyper   r  )r   r	   r   z'TypeGuard[TensorBox | sympy.Expr | int])r   r	   r   zTypeGuard[TensorBox | bool])rt   r	   r   r   r   r   r   r  )r`   z9Any | torch._ops.OpOverloadPacket | torch._ops.OpOverloadr   r  r   r   )r   rB   rv   torch.devicer   rB   )rt   	list[Any]r   zdict[str, Any]r  r   r   &ELEMENTWISE_TYPE_PROMOTION_KIND | Noner  r   r   z tuple[list[Any], dict[str, Any]])r   r1  r  rl	  r   rl	  )
r  rl	  r  r   r   rp	  r  r   r   rV   )r   rp	  r   z.Callable[[Callable[_P, _T]], Callable[_P, _T]]rc  )NNNFFN)FrZ  )FT)r   rB   r   r  r  r   rl  r   )r   rB   r   r  )r   rB   rv   rn	  )F)r   rB   r  rB   r   )r   rB   r  zSequence[sympy.Expr]r   rB   r  )r  rB   r  rB   r  rB   r  r   r  r   r  r   r   r  r   rB   )r  rB   r  rB   r  rB   r  r   r  r   r  r   r   r  r  torch.dtype | Noner   rB   )r  rB   r  r5  r  r   r  r   r  r   r   r  r   rB   )r  rB   r  r5  r  r   r  r   r  r   r   r  r  rq	  r   rB   )r  rB   r  rB   r  rB   r  r   r  r   r   r  r   rB   )r  rB   r  rB   r  rB   r  r   r  r   r   r  r  rq	  r   rB   )r   r   r.   )rx  r   r?  r   r@  r   rL  )T)rw   ztorch.Tensor)rc   ztorch.fx.Node)NTFN)rv   rn	  )
r   	list[int]r  rB   r  r  rx  r   r  r  )
r  r   r  r   r   rr	  r  rB   rx  r   )r%  rB   r   z.tuple[str, sympy.Expr, sympy.Expr, sympy.Expr])r%  rB   r   ztuple[str, sympy.Expr])r9  rB   rC  rB   r,  r   r-  r   r.  
str | Noner/  TensorBox | Noner   rB   )r  rB   rH  rB   r,  r   r-  r   )r.   )r   r   r:  r   )r   NNr.   )NNN)rh  FF)r   r   )r  r1  r   r   r  rs	  ry  r   )r   r   r  rs	  )r   r   ry  r   )r,   F)r  ztuple[float | None, ...]r  r   r  r   )r  float | None)r  ru	  r  ru	  )r  ru	  r  ru	  r  ru	  )r   rB   r  zSequence[int]r  r5  r   rt	  )rs   
sympy.Exprr  zsympy.Expr | int)rs   rv	  r  rv	  )rs   rv	  r  rv	  r  rv	  )Nr  N)
r"  rB   r  Sequence[int | torch.SymInt]r/  rw	  r.  zlCallable[[Sequence[int | torch.SymInt], Sequence[int | torch.SymInt]], torch._inductor.virtualized.OpsValue]r   rB   )Nr   r.   F)r  )NNNN)r   r   FTN)r  rS   r   )r   r.  r   zir.Constant | None)rh  F)rh  TT)r`   ztorch._ops.OpOverloadPacket)r   z-list[ir.TensorBox | ir.ShapeAsConstantBuffer])r  ir.Subgraphr  r  )r  ztorch.fx.GraphModulert   ro	  )r  rx	  )r  rx	  r  ztuple[torch.Tensor]r0  (B  
__future__r   r  r  r  r   loggingr  r?  r  rm  r  collectionsr   collections.abcr   r   r   r   typingr	   r
   r   r   r   typing_extensionsr   unittest.mockr   r   r{   $torch.ao.quantization.fx._decomposedtorch.fxtorch.utils._pytreerC  _pytreerp  torch._dynamo.utilsr   (torch._higher_order_ops.associative_scanr   r  r   "torch._library.fake_class_registryr   torch._library.opaque_objectr   torch._library.utilsr   torch._prims_commonr   r   r   r   r   r   r   r   r   r   r    torch.fx.experimental.sym_noder!   r"   rh  r#   r$   r%   torch.utils._ordered_setr&   torch.utils._sympy.functionsr'   r(   r)   r*   r+   _dynamo.utilsr-    r/   r0   r1   r2   decompositionr3   r4   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   r  rQ   rR   r  rS   rT   rU   FALLBACK_ALLOW_LIST	getLoggerr  r  rW   __annotations__rX   rY   r|   r}   r  r  tr_c10dr  r   _higher_order_opsr  rb   r  rZ   quantized_decomposedrf   rx   r   r~   r   r   r   r{  r  r  r  bmmconvolutionconvolution_backwardr<  r@  rl  r  r  r  _int_mmr  r  r&  r  r  rH  r  r  	complex32	complex64r   rG  r   r   r   r   r   r   r   r   r
  r  r!  rE  r$  r2  rN  r`  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!  r%  r'  r-  r/  r2  r5  r  r7  r-  rB  rD  rS  _unsafe_viewreshaperV  slicern  r  r  r  quantize_per_channelr  r  r   r  _functional_assert_asyncr  dequantize_per_channelr  quantize_per_tensorr  dequantize_per_tensorr  r@  r  r  r  rO  rQ  rU  r]  rb  rd  rf  rq  r>  ru  ri  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	rand_like
randn_likerandint_likestreamsrecord_event
wait_eventsynchronize_eventsynchronize_deviceforce_stride_orderr  r  r  r   r  lookup_seedr  r
  randomr  rand_eager_offsetrand_eager_offsetsr   r'  r*  rB  r  r7  	NO_OPMATHrL  rQ  rS  r   r2  rY  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_normrh  _is_compiledembedding_dense_backwardr  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_backward	nanmedianrQ  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run_dtensor_rng_opmasked_scattermasked_scatter_backwardr  angle_efficientzerotensor(_sparse_coo_tensor_with_dims_and_tensors	to_sparse
_to_sparser   r  r  rn  #_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  arange
start_stepr  r  r  r  scalar_tensorr  
LongTensorr  r  r  r  r  r)  r  rO  r  r	  r  r  
zeros_liker  r  r  r  r  r  r  r  r  r&  r)  r(  r:  r2  r:  r<  r>  rI  rD  rN  r7  ri  fallback__unsafe_masked_indexrm  ,fallback__unsafe_masked_index_put_accumulaterg  ru  r`  r}  rw  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r#  r'  r2  r5  r7  rd  rz  r  r  r  r  r  r  r  r  r  r  r  r  r  
avg_pool1dr  r  r  rJ  rE  rM  r`  rc  rg  r  ry  r  r  r  rS  r  r  r  r  r  Tensor_Tensorr  Scalarr  Tensor_Scalarr  rS  r  rs  r  r  r  rm  r  rz  r  r  true_divider  r$  rl  r  r  r  r  r  r  r	  r  r  r  r  r  r  rq  r  rp  r  r  r  r  r  r  r  r  rP  r  r   r.  r!  r(  r,  r]  r[  r`  r_  medianr0  r   r6  r  rM  r3  r8  rY  rc  re  rsqrtr  r  expm1relury  r  rs  r  rv  rz  _foreach_addcmul_foreach_addcmul_scalar_foreach_addcdiv_foreach_addcdiv_scalarcossinabsbitwise_andbitwise_left_shiftbitwise_not
bitwise_orbitwise_right_shiftbitwise_xorlgammaerfspecial_erfr  tantanhr  r{  r  logical_xorr  r  	clamp_min	clamp_maxnegr  	remaindersignsignbit	_neg_viewler  r  r  r  r@  coshsinhacosacoshasinasinhatan2atanatanhr  erfcerfinvhypotlog10log2	nextaftercodegen.commonr  r}  r  r   r`   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_clone_foreach_copyforeach_copyr  _foreach_add__foreach_mul__foreach_div__foreach_copy__foreach_addcmul__foreach_addcdiv_r  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__r  r  r   r  r  r   methodfuncr  r  _inductor_testr  r  r  r  set_source_Tensorr  fsdpr  r  *torch._higher_order_ops.auto_functionalizer  r  r+	  r   r  while_loop_stack_outputr#  r  r  r  r  .torch._inductor.fx_passes.control_dependenciesr  r  invoke_quantr	  r	  r	  r	  rA	  comm_loweringrB	  rC	  rQ	  rU	  cvt_e8m0_rceilr]	  r[	  re	  r  rf	  register_quantized_opsregister_woq_mm_opsrg	  register_onednn_fusion_opsrh	  register_jagged_opscontextmanagerrk	  r   rg   re   <module>r
     sH=   "        	 
  # D D ? ? '    +  $ $ ( H U ? 8 :    M 
 0  - 8 8 =    $      * T]t_ !   g!@B	= BBD? DTV Q V-Juzz,,-/	yy~~
))

		9
5::#8#89; /j../
))*
 8j!6!679 NP K Pyy55 90B	B*BF
	" 	"3;3.3 
3   !!$$$$--&&, {{zz{{{{{{}}}}}}

."# "'CC8C C 	C( ANQ	(O
OO O @	O
  O &Od"/A42!2 2 @	2
  2 F2n ,44 4*#4?D:  $fR5pB UYEE$E,0EMQE2 5**77TR) S)X 5--4HN IN* @E 5. 499??=& >& ;@e K 5##>F ?F 
7?? $,^ $7DD 4::4H I>>   	
 4::DI J0 4))UPTU V. DJJT\\499emmTU V 4&doo&s+ 4<<T:B ;B, 4$$$?" @" DMM?# $ 4::C C 4::C C 499! ! 4::! ! 4::%%&% '% 4::! ! 4;;D9> :>4 5))tD
 E
 4>>t<# =# 4;;4  4n 4$$$?499$74<<T:1 ; 8 @1 4<<T:> ;> 4::481#++AT eE 9eEP 4??=J >J8 4##> ? 4''TB C
BJ '<<RVW/// / 	/
 / / / / X/d$ 4%%))*$ +$ 400445$ 6$ //T %)/// / 	/
 / / / "/ //d ,,44$!!! ! 	!
 ! ! !!H ..66D %)!!! ! 	!
 ! ! "! !!H ,,33/// / 	/
 / / //d ..554 %),,, , 	,
 , , ", ,,^ 488b: b:J 4==d;5E <5Ep 4%%4@6 A6 4((dC D 4;;D9=C :=C@ 4::48 94 4((dC  D  4;;D9 : 4;;D9D :D8 4>>t< = 4??= > 488 $  D
3%6T 	6Nr	, 599%%11tL#+ M#+L 4&&DAG BG 4??= > 4>>##>' ?' 4<<  ! QB B/
 )):):; *499+>+>? )$***<*<= +DJJ,@,@A  dll  dnnd 3 doot 4 d 6 eii,,44 5 eii**22 3 eii1199 : eii22:: ; 499J J 4::J J >44$OL PL
 >&&DAK BK >''TBJ CJ
 >--4H	 I	( >((dC $}}5
5
5 5
 5 5 D5p n.. / n// 0 >))tDLM	(09FI E8 . 4$$++F
 #^^
^ 	^
 ^ ^ ^ ^ G^B NN/N/X/X 777 	7
 77t 0 d))113J KaJ d'' ( d&& ' d== > dll ' d&&U 3 d!!#5 6 d,,5 9 d"" /99%%E 	::U
 d!! " d"" # doo  d!!)) * d++33 4 djj  dnn## $ d))11 2 d,,44 5 dkk  d$$5 1 d . d'')@ A doo} - d&&(: ; doo} - d&&(: ; d!!#5 6 d..0B C d** + d== > d== > d22 3 d== > d33 4 d00- @ d00 1 d// 0 d// 0 d11 2 d11 2 d-- . d-- . d-- . d..0B C d00 1 d++ , d""$6 7 dnn  dmm  dll  doo  d  d-- . d   ! d'' ( d## $ dnn  d&& ' d"" # d$$ % dnn  d"" # d## $ d** + d  dnn  djj  d'' ( d// 0 d   d## $ d%% & d## $ d!! " djj  dmm  dll"" # dhhU + d((- 8 ell$$;; < ell$$77 8 ell$$AA B ell$$77 8
 d!! " d** + d""$6 7 djj  d'' ( d;; < dnn  doo  djj  0088	
 99AA	
 ,,44	
 ,,66	 55==	
 ,,44	
 55==	
 44<<	
 ==EE	
 99AA	
 BBJJ	
 d++33_ E d++55 6 d,,44o F d//77 I d0088/ J d   d$$++T B d22::<N O
 499$7 8 4::"  * 4"#+d**+E2 5:: ( 4;;))tD 

  E. 4&&DA  B F 4%%4@A AAH ELL$"4"4e6I6IJKtDU 5 L5p 5??#4 $4 5##$+ %+ 4++,)" -)"X 4&&' ( 4//0%)59$ 1<K8 EKK,- 
 .$.> 0t/0B50IJ
1!45	 21 56
  4>>" $T$4  # 4%%&Td '@ 4))*"4 +" 5%%--.A /A EJJ		*+: ,:
 4;;D9! :!H 4>>t< =B#F FR"1J 4::48
 9
 4%%4@/ A/ 4>>t< = 4))* +;. 4??= > >44$O Psl !1%%5!  0@,,44%0 ,
 4,,$G H, 4;;QUV	P W	P 1 1 4**E F 4<<T:9 ;9 & 
  D 4==d;AE : <:" 4##>3 ?3 4$$$?6 @6 4&&DAP BP 4''TBPT q Cqp , ', 	,
 ,^ 4**223> 4> 41199:J ;J 4**223LPI*I=II 4I 41199:LPU*U=IU ;U 4**223 "!!S S 	S
 S 4S 41199: "!!	 	 		
 	 ;	7 599$$% &*\\(\6;\\~2)2)(2)6;2)2)j 4''TB8 C8v :>@ MQ < JN@:HV 5::PTU 5 V5D- -	 > 	11t$< 4//TJ 
 K
 4//TJ 
 K
 -=))11- ) 488dSN TNb8B&R2j  0%%5  
 4,,-B .BJ  0$$%  
 4++,N -Nb)X 4--.V /V 4--.V /VAH 433;;<BF. =.b 4??=  >* 4??=  >, T__,,%HT__,,%HT__,,%H A#H  0$$%  
 4++F b GbJ  0$$%  
 4++F D GDN& HLOd < 499:t : : 40(V4. DHHeii()T5  * 4==!u  "   .HH  'txxER -HH 
 488t,- --`"J 4::2 2 4::48 9       488t, -& DHH:.( /(@ EII;$/$ 0$. 	txx'7DD

 DIIuzz*d;$ <$  DHHeii(). . *. #4;;#6#67#DLL$8$89 ():):)B)BC "4;;#6#67"4;;#6#67 4;;  . 4<<  !. 4$$% &* 4;;D9 :, 4;;D9 :, 499. . . 488@ @
 48867 77 48867 77 ! %--  	!: ;*		*>%+@A*		*>%+@A.!$++.85;;? /!$++.85;;? HH ,	 !!1!1uM 499##>!re +2 ?+2\ 499$$$?H @H !!2!2N$T]]%:%:PUV "4;;#6#6ER&t{{ER  !2!2N 4;;&&DA% B% 4;;??= >& 499$$$?j  @j Z 499$$$? @  4==((dC D$ 	#4::.&txx0!$)),"4::.$))$
*4<<
8'		2	DKK	(t4 4<<40-. F 1FR 4<<40-. F 1FR 5  't'  5  't' 
 ! *   *" !1!12'(?(?@  -  0
()A)AB  !1!12 4;; ' * *I*V*V 4:: & 488 $ 499 %   * **	 !**	  OO**	
 !**	 T\\
*
T\\
* ! $.. !' * ! $.. !' *""'8
 4>> "$))L 499  4<<uzz B ! $.. !# & 477%** = 477%** = 477%** =uzzB 477%** = 477%** = 499 % 499 % 499 % 4:: & 499 % 4:: & 4:: & 499 % 4:: & 4== ) 499 % 4;; ' 4:: & 4:: & 499 % 4>> * D DE$ %D4Ld50 	 3+		
5 5Mt50 	 3+		
5 %, .CT  0ct  4,,33Sd K-d.?.?.D.DcJ  4,,33S 9/0A0A0H0H#N  4,,113 7 4,,33S 9 4,,44c : 4,,44c : 4,,33S 9 4,,113 7 4,,<<c B-d.?.?.D.DcJ  4,,33S 9/0A0A0H0H#N  4--t 4 4.. 6 40055w ? 40077 A 40055w ? 40077 A 42277 A 422997 C 42277 A 422997 C 433Z @ 4--t 4 4.. 6)$*<*<dC, T..335E t00779K T..335E t00779K T..335E t00779K !3!3!;!;\ !!  
 !!   C   ""K 0 ))+= > ""K 0 !!: . **,? @ ""K 0 C   !!3 ' && 1 ""K 0 ""K 0 !!: . ""K 0 C   T "  (   $,,  , " $// "#5 6  $++ z * " $// "#6 7  $,,  ,  - !!4?? 3 t{{ + !!4?? 3  - 4++, - 4==$$% & 4??&&' ( 4>>" # (M'')LFD1(01$7 * 5==! " 4<< 7 !7 599++334 5
 599%%;;< = 599>>&&445K 6K 599>>7#uyy~~++334# 5# 599>>(()%) = *=@ K J ! " 12S 3S( 599))..DI/2/ J/ 599))44$OE PE& 	II22)JT24
 599))99tT/ U/
 8 H H G <T:+ ;+\ 5**77TRDH  S8 &DA  4G  B F 599??//778 9 599??..667 8 599))66DQc  Rc L P O     >88dS* T*ZW
 >00dK") L")J 	22 37QF     ! ! +  * * , '  ' ' )   ,  + + -   %  $ $ &  rg   