
    i	                   y   U d dl mZ d dlZd dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
Z
d dlZd dlmZ d dlmZmZmZmZ d dlmZmZmZmZmZmZmZ d dlmZ d dlmZ d dlZd dl Z d dl!Z d dl"Z d dl#m$c m%Z& d dl'm(Z( d d	l)m*Z* d d
l+m,Z, d dl-m.Z. d dl/m0Z0 d dl1m2Z2m3Z3m4Z4m5Z5m6Z6m7Z7m8Z8m9Z9m:Z:m;Z;m<Z< d dl=m>Z>m?Z? d dl@mAZAmBZBmCZC d dlDmEZE d dlFmGZGmHZHmIZImJZJmKZK ddlLmMZM ddlNmOZOmPZPmQZQmRZR ddlSmTZTmUZU ddlQmVZVmWZWmXZXmYZYmZZZm[Z[m\Z\m]Z]m^Z^m_Z_m`Z`maZambZbmcZcmdZdmeZe ddl$mfZfmgZgmhZhmiZimjZjmkZkmlZlmmZmmnZnmoZompZpmqZq ddlrmsZsmtZt erddlumvZv  ed      Zw ed      Zx eEddg      Zy ej                  e{      Z|i Z}de~d <   i Zd!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dvd%Z	 	 	 	 dwd&Zdxd'Z	 	 	 	 dyd(Zdzd)Z	 	 	 	 d{d*Z	 	 	 	 	 	 d|d+Z eej(                  e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                  g       e jD                  e jF                  e jH                  e jJ                  e jL                  e jN                  e jP                  e jR                  e jT                  e jV                  e jT                  e jX                  e jZ                  d,Zd}d-Zd~d.Zdd/Zd0d1	 	 	 	 	 	 	 dd2Zd3 Z	 	 	 	 	 	 dd4Zdd5Z	 	 	 	 	 	 	 	 	 	 	 	 dd6Z	 	 	 	 	 	 dd7Z	 	 	 	 	 	 	 	 	 	 dd8Zd0e7jr                  d0e}f	 	 	 dd9Zd: Zdd;Z	 	 	 	 	 dd<Zdd=Zd> Zddd?Z ee j                  j                  d@      dA        Z eej                  d@      ddB       Zd0dCddDZ eej                  j                  d@      ddE       Zd0d0dFddGZ eej                  d@      dddH       ZddIe7jr                  d0ddd0dfdJZ endKe7j                  dL        eej                  dIdM      ddN       ZdO Z eΫ        	 ddPZ eej                  d0dM      dQ        Z eej                  d0dM      dR        Z eej                  ej                  ej                  ej                  ej                  g      dS        Z eedT      r  eej                        e׫        eej                  d@      ddU       Z eej                  d@      ddV       Z eej                  g      ddW       Z eej                        dX        Z eej                        dY        Z eej                        dZ        Z eej                        d[        Z eej                  j                        d\        Z eej                        d]        Z eej                  d@      d^        Z eej                  d@      d_        Z eej                  d@      d`        Z eej                        da        Z eej                  d@       eej                  d@       eej                  d@      ddb                     Z eej                  d@      dc        Z eej                  d@      ddd       Z eej(                  d@      dde       Z eej                  d@      ddf       Z eej*                  d@      ddg       ZddhZ eej                  d@      	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddi       Zdj Z eej                  j                        dk        Z eej                  j                        dl        Z eej                  d@      ddm	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddn       Z eej                  j                  d@      	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddo       Z eej                  j                  d@      ddm	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddp       Z eej                  j                  d@      	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddq       Z eej                  j                  d@      ddm	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddr       Z eej                        dds       Z eej                   d@      dddt       Z  eej                  d@      dddu       Z eej                  d@      dddv       Z eej                  d@      dw        Z eej                  d@      ddx       Z eej
                  d@      ddy       Z eej                  d@      ddz       Z eej                  d@      d{        Z eej                  d@      d|        Z eej                  d@      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(                  j*                  d@      d        Z eej,                  d@      d        Z eej.                  d@      d        Z eej0                  j2                  d@      d        Z eej6                        d        Z ej8                  d      d        Zd Z eej>                  j                        Z  eej>                  jB                        Z" eejF                  j                        Z$ eejF                  jB                        Z% eejL                          ee j                  jN                  jP                  j                          ee j                  jN                  jR                  j                          eej>                        d        Z eejF                        d        Z# eePjT                  d@      d        Z+ eePjX                  d@      dd       Z- eePj\                  d@      d        Z/ eePj`                  d@      d        Z1 eePjd                  d@      d ddd       Z3 eePjL                  d@      d d	 	 	 	 	 	 	 	 	 dd       Z4ddZ5ddZ6 eejn                  jp                  d@      d0d0ddd	 	 	 	 	 	 	 	 	 	 	 	 	 dd       Z7 eejr                  jp                  e7jt                  @      d0d0d	 	 	 	 	 	 	 dd       Z9d Z;d Z<d Z=d Z>d Z?d Z@d ZAd ZB eej                          eej                          eej                          eej                  d0        eej                  j                  d0        eej                  e<        eej                  d0        eej                  d0       e j                  j                         r eej                  d0       e j                  j                         r eej                  d0        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                  d0        eej                  e;        eej4                  eA        eej                  e;        eej                  e<        eej                  e;        eej                  e<        eej                  e<        eej                  e<        eej                          eej                          eej                          eej                          eej                          eej                          eej                  e;        eej                          eej                          eej                          eej                          eej                          eej                          eej                          eej                          eej                  e<        eej                          eej                          eej                  e<        eej                          eej                  j                          eej                          eej                          eej                          eej                          eej                          eej                          eej                          eej                          eej                           eej                          eej                          eej                          eej                          eej
                          eej                          eej                          eej                          eej                          eej                          eej                          eej                          eej                          eej                          eej                          eej                           eej"                  j$                          eej&                          eej(                          eej*                          eej,                          eej.                          eej0                          eej2                          eej4                  j                          eej6                  j                  d0        eej8                  e;        ee j:                  j<                  j>                          ee j:                  j<                  j@                          ee j:                  j<                  jB                          eejD                          eejF                          eejH                  e<        eejJ                          eejL                          eejN                          eejP                          eejR                          eejT                  jV                          eejX                  j                  eBd0        eejZ                  j                  eBd0        eej\                  j                  eBd0        eej\                  j^                  d0        eej`                  j                  eBd0        eejb                  j                  eBd0        eejd                  j                  eBd0        eejf                  j                  eBd0        eejh                  j                  eBd0        eejj                  j                  eBd0        eejl                  j                  eBd0        eejn                  j                  eB        eejn                  j^                          eejp                  j                  eB        eejr                  j                  eB        eejt                  j                  eB        eejv                          eejx                  jp                  dI        eejz                  j                  e<        eej|                  d@      dd       Z eej~                        ddd       Zd Z eed      r  eej                        e        eej                        d        Z eej                  d@      dd       Z eej                  d@      dd       Zd Z ee j                  ej                  g      dddd0d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@      ddǄ       Z eej                  d@      ddȄ       ZdɄ Z	 ddʄZd˄ Zdd̄Z eej                  d@      d̈́        Z eej                  d@      d΄        Z eej                  d@      ddτ       Z eej                        ddЄ       Zdф Zd҄ Z eej                  d@      ddӄ       Z eePj                  d@      ddԄ       ZddՄZ eej                  j                  d0֫      Z eej                  j                  d0֫      Z eej                  d@      dׄ        Z eej                  d@      d؄        Zedل        Z eej                  d@      ddڄ       Z eej                  d@      ddۄ       ZddIdܜ	 	 	 	 	 	 	 dd݄Z eej                  d@      ddޜdd߄       Z eej                  d@      dd       Z eej                  d@      dd       Z eej                  d@      dd       Z eej                  d@      dIddd       Z	 	 d	 	 	 	 	 ddZ eej                  j                        ddd       Z eej                  j                        ddd       Z eej>                  j                        	 d	 	 	 dd       Z eej@                  j                        	 d	 	 	 dd       Z eej                  j                        	 	 	 d	 	 	 	 	 dd       Z eej                  j                        	 	 	 d	 	 	 	 	 dd       Zd Z eej                  j                        d        Z	 	 	 	 	 	 	 	 ddZ eej                   d@      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@      	 dd       Z		 	 	 	 	 	 	 	 	 	 ddZ
 eej                  d@      d        Zd Z eej6                  d@      	 	 	 	 dd       Z eej8                  d@      	 	 	 	 dd       Z eej:                  j                  d0֫      Z eej:                  d@      d         ZddZd Zd Zd Z eej$                  j                  d0֫      Z eej$                        d        Z eej(                  j                  d0֫      Z eej(                        d        Zd Z eej.                        d        Z eej0                        d	        Zd
 Z eej4                  j                        	 dd       Z eej,                  d@      	 	 	 	 	 dd       Z eej6                  d@      	 	 	 	 	 dd       Z eej8                  j                  d0֫       eej,                  j                  d0֫       eej6                  j                  d0֫      gZd Z eej.                  j                  d0֫      Z eej.                  d@      	 dd       Z eej@                  j                  d0֫      Z! eej@                  d@      	 dd       Z d Z"dddZ#dddZ$d Z% eejL                        dddd       Z&d Z'd Z(d Z)d Z* eejV                  ejV                  g      ddd0dd       Z, eejZ                        ddd0dd       Z-d Z.ed         Z/ eej`                  jb                  d0֫      Z2 eej`                  jf                  d0֫      Z4 eej`                  jj                  d0֫      Z6 eej`                  dI!      d"        Z0dd#Z7 eejp                        d$        Z8 eejr                  d@      dd%       Z9ed&        Z:ed'        Z; eejx                  dI!      dd(       Z= eej|                  gdI!      d)        Z>dd*Z? eejx                  gdI!      d+        Z@ eej                  ejx                  jp                  gdIe7j                  M      d,        Z< eej                  ej                  gdI!      d-        ZB eej                  ej                  g      dddd.       ZD eej                  j                        ZF eej                  j                        ZH eej                  j                        ZJ eej                  j                        ZL eej                  j                        ZN eej                        dd/       ZE eej                        dd0       ZG eej                        d1        ZI eej                  d@      dd2       ZK eej                  d@      dd3       ZM eej                        dddd4       ZO eej                        dd5       ZQ eej                  d@      dd6       ZS eej                  d@      dd7       ZU  eej                         e$d8               eej                         e$d9            ZX  eej                         e$d:            ZZ  eej                         e$d;e jL                  <            Z\  eej                         e$d=e jL                  <            Z^ eej                  dId>?      Z_ eej                  j                  d0֫      Z` eej                  j                  d@      ddd0d@dA       Za eej                  j                  d@      ddB       ZvddCZbddDZc ebej                        Zd ecej                        Ze ebej                        Zf ebej                        Zg eej                        Zh ecej                        Zi ecej                        Zj eej                        Zk eej                  dIE      Zl eej                  dI!      ddFdG       ZmddHZn eej                  jf                  en        ecej                          ecej                          eej                        Zr eej                        Zs eej                        Zt eej                  dIJ      Zu eej                        Zv eej                        Zw eej                        Zx ebej                          ebej                        Zz  eej                  e7j                  @      ez        ebej                          ebej                          ebej                          ecej                          eej                  ddIe jX                  K      Z eej                   ddIe jX                  K      Z eej                  ddIe jX                  K      Z eej                  ddIe jX                  K      Z eej                        Z eej                        Z  eej
                        e         eej                        e        eej                        Z eej                        Zr ebej                        Z eej                          eej                  dLJ      Z eej                          eej                  e jX                  <         eej                        e        eej                  e jX                  <        eej                  e jX                  <        eej                  e jX                  <        eej                   e jX                  <      Z eej"                  e jX                  <        eej$                  e jX                  <        ebej&                          ebej(                          ebej*                          ebej,                          ebej.                          ebej0                          ebej2                          ebej4                          ebej6                          ebej8                          ebej:                          ebej<                          ebej>                          ebej@                          ebejB                          ebejD                         ddMlmZmZ dN ZeD ]S  Z eee      D ]  \  ZZZ eʐeeeeO         eee      D ]  \  ZZZ eʐeeeeO        U  eejV                  jX                  e_dIE      Z eejV                  jf                  e_dIE      Z eejV                  jp                  e_dIE        eej^                  jX                  e>      Z eej^                  jp                  e>        eej^                  jf                  e>      Z eejd                  jX                  el        eejd                  jf                  el        eejf                  j                  e        eejh                  j                  er        eejj                  jf                  e0        eejj                  jX                  e0        eejj                  jl                  e0        eejn                  jX                  e<      Z eejn                  jp                  e<        eejn                  jf                  e<      Z eejt                  ej        eejv                  ed        eejx                  jX                  e        eejx                  jf                  e        eejz                  jX                  e        eejz                  jf                  e        eej|                  jX                  e        eej|                  jf                  e        eej~                  jX                  e        eej~                  jf                  e        eej                  e        eej                  e        eej                  e      ZÐdP Z eej                  jX                  ejV                  jX                  e        eej                  jf                  ejV                  jf                  e        eej                  jX                  ej^                  jX                  e        eej                  jf                  ej^                  jf                  e        eej                  jX                  ejn                  jX                  e        eej                  jf                  ejn                  jf                  e        eej                  j                  ej                  j                  eë       dQ Z eej                  e_        eej                  es        eej                  et        eej                  eu        eej                  ev        eej                  ew        eej                  ex        eej                  e>        eej                  jp                  e<        eej                  j                  e=        eej                  e        eej                  e        eej                  e        eej                  e        eej                  el        eej                  eh        eej                  ei         eej                        es         eej                        et         eej                        ev         eej                        ew         eej                        ex        eej                  ej                          eej                  ej                          eej                  ej                          eej                  ej                          eej                  ej                          eej                        ddR       Z eej                  j                        dS        Z eej                  j                        dT        Z eej                        dU        Z e>j                         D ]  \  ZZ  e e?e            e         ee j                        dV        Z eej6                        dW        Z ee j                  j                  j                        dX        Z ee j                  j                  j                        dY        Z ee j                  j                  j                  j                        dZ        Z ee j                  j                  d[      r; ee j                  j                  jr                  j                        d\        Z ee j                  j                  j                        ddd]       Zd d^lmZ  ee        ee,      d_        Z ee j                  j                  j                  d@      	 	 dÐd`       Z ee j                  j                  j                  d@      dda       Z  ee j                  j                  j                   d@       ej                  edIb              ee j                  j                  j                  d@      dĐdc       ZdŐddZd delmZ  eed@      df        Z ee j                  j                  d@      ddgdƐdh       Z ee*d@      	 	 	 	 dǐdi       Z	 ee j                  j
                  j                  j                        dj        Z
 ee j                  j
                  j                  j                        dk        Z ee j                  j                  j                  d@      dl        ZddmlmZmZ  e         e         eePj                   d@      dn        Zdo Z eePj$                  d@      dp        ZddqlNmZ  eMe       ddrlNmZ  ej,                           ej.                          ddslNmZ  ej2                          ddtlNmZ  ej6                          ej8                  dȐdu       Zy(      )annotationsN)defaultdict)Callable
CollectionIterableSequence)AnycastOptionalTYPE_CHECKING	TypeGuardTypeVarUnion)	ParamSpec)patch)counters)associative_scan_op)triton_kernel_wrapper_mutation)FakeScriptObject)get_layout_constraint_tag)canonicalize_dimcanonicalize_dimscheckdtype_to_typeelementwise_dtypesELEMENTWISE_TYPE_PROMOTION_KINDget_computation_dtypeis_boolean_dtypeis_float_dtypeis_integer_dtypeNumber)magic_methodsmethod_to_operator)free_unbacked_symbolshas_free_unbacked_symbolsresolve_unbacked_bindings)
OrderedSet)CeilDivFloorDivIdentityModModularIndexing   )import_submodule   )configinductor_primsirtest_operators)decompositionsget_decompositions)BaseView	DtypeView
ExpandViewIndexingConstantIRNode	is_triton
MutableBoxOnlineSoftmaxReductionops_wrapperPermuteView	Pointwise	ReductionSqueezeView	TensorBoxvalidate_irView)ceildivdecode_device
is_dynamicis_gpuis_pointwise_useis_view,needs_fallback_due_to_atomic_add_limitationspad_listlike#register_op_dtype_propagation_rules#register_op_requires_libdevice_fp64sympy_productuse_scatter_fallback)opsV)ReductionType_T_Pztorchvision::roi_alignzaten::index_add8dict[Union[Callable[..., Any], str], Callable[..., Any]]	loweringsz9dict[torch._ops.OpOverload, Optional[Callable[..., Any]]]_maybe_layout_constraintsz2dict[torch._ops.OpOverload, torch._ops.OpOverload]inplaceable_foreach_opsc                     t         j                  j                  j                  D ]7  } | j                  D ]&  }|j                  dk(  r|j
                  t        v r%  y 9 y)Ncall_functionTF)rS   graphcurrent_nodeusersoptargetforeach_ops)nodeusers     i/var/www/vps2.regionflexible.com/Desarrollo/venv/lib/python3.12/site-packages/torch/_inductor/lowering.pycur_node_has_non_foreach_usersrf      sS    $$** JJ 	D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   rH   r0   #combo_kernel_foreach_dynamic_shapesrC   data
get_deviceappend)	arg_pairsoutunpack_argsiargsuse_foreachdevicets           re   group_foreach_argsrx      s     d
CKY' 54$)K7DD!!OV%O%O 	  	A!Y'**,	 !S#SS!GTV[!"))1d)45  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      %
&.GB!%
   )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9)!,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      sN     "ejj112 	JH?I%gb(&;<	J )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   rC   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   rC   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_input3  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   .  sP    
A -11SOC 1D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   rX   ro   )aten_fnr   r   other_fns       re   get_overloadsr   B  s    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   R  sQ     "ejj112B1111	B

--	.BGGI%%rg   c           	     $   t        | j                  t        j                        rt	        | j                               r| S | j                         D cg c]+  }t        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 c c}w )zB
    Copy cpu scalar if doesn't not match with given `device`
    cpur   r/   F)rk   rm   r2   ReinterpretViewr%   r   rS   r]   sizevarssize_hint_or_throwrn   typer   rC   
StorageBox
DeviceCopycreate)r   rv   ssize
cur_devices        re   maybe_copy_cpu_scalarr   \  s     affb0015N	

6 <=JJLIqAGG//2IDIJOOu$& Y!^D	Q47a<r}}';';Az5'QRSSH Js   0D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!!Ts   r   r   c                    t        | t              rt        |       S t        | t        j                        r"t        j                  | j
                        S | S )Nvaluer   rv   )rk   rC   to_dtyper2   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     3T!W3   c              3  (   K   | ]	  }|     y wr   r   )r   kkwargss     re   r   z!transform_args.<locals>.<genexpr>  s     71VAY7r   N)r   r	   r   r	   )rj   rk   rC   itemsr{   r   r!   r   r   r   extendvaluesr   rn   r   broadcast_tensorsri   	itertoolschainr   zipr   ranger2   r   r8   r   )rt   r   	broadcastr   convert_input_to_boolrs   r   args_indicesr   vkwargs_indicesr   promoting_argsr   broadcastedr   rv   r   s   ``              @@re   transform_argsr  p  s    #,D/N$!QZ95MANLN$*LLNODAqjI6NaONOV|3 JJE
 a&%++!6771g;N N  !!!TV]]_!TT&$7E &2Da!vnQ>O7P
*, 	  	=A+DGV<DG	=   	AA-fQi@F1I	A	 %))q
),2LLN;DAq!WQZ-;'3l377
 KN++-.k2EC4E&FG 	DAqDG	C4E4G(HI 	DAqF1I	 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   )rD   )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   rX   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)      
JLR9_DR1C DD
s    "c              3  &   K   | ]	  }|d k(    yw)rq   Nr   r   r   s     re   r   z6_register_lowering.<locals>.wrapped.<locals>.<genexpr>  s     6!1:6   zout= ops aren't yet supported)	ri   r  r   rk   r   allanyr  rD   )	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OnerS   r]   r   is_size_one_or_falsero   check_equalsr   expandfree_symbols)r   boutputr   ys        re   broadcast_symbolic_shapesr3  '  s    	aAQF%%hqk8A;%''++V 
!177003MM!WW2215MM!GG))!Q/5<<?//03u||A7S7S3TTa a 
! &!""rg   c                V   |	|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              }g }| D ]  }t        |t        t        f      ro|j                  t        j                  t        j                  ||j                         |j                               t!        |j#                                            t        |t$        j&                        rf|j                  t        j                  t)        ||j                         |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>I  s"     HAz!ekk367H   /1c              3  f   K   | ])  }t        |t        t        t        j                  f       + y wr   )rk   r   r6  r   r   r  s     re   r   z$promote_constants.<locals>.<genexpr>K  s"     
D:a#uekk23
Dr7  r   c                    t        | t        j                        r!t        j                  | t        d             S t        j                  | t        d             S )Nindexr   rv   r   )rk   r   r   r2   r9   rG   r   r   r   s    re   
const_funcz%promote_constants.<locals>.const_funcR  sJ    !U[[)**5t1D  {{%d@STTrg   c              3  l   K   | ],  }t        |t        t        t        j                  f      s)| . y wr   )rk   rC   r8   r2   r   r  s     re   r   z$promote_constants.<locals>.<genexpr>[  s$     WA:a)Z1U#VaWs   *44r   r:  )r   DEFAULTr  r  r   nextrk   r   r6  ro   r8   r   r2   r   r   get_device_or_errorri   r   r   r   r9   )inputsoverride_return_dtyper   r=  r   exrq   r   s          @re   promote_constantsrE  A  s    (,?,G OG $)<)D=EEHHH

DV
DD% 
);*
 !4*
	U (..!
1..	WW	WB
C a#u&JJ!!KKr||~b>T>T>V '	 5;;'JJ!!$r||~b>T>T>V '	 JJqM), J3 /s   F&c                .     d dd fd}|S )Nalphac           	       	
 t        d D              r	rJ   S t              r%| '| dk7  r"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                  j"                  j$                  d uxr4 t        j                  j"                  j$                  j'                  dd	      		xr v 	
f	d
}sSd }D ]7  }t)        |j+                         j,                        s'|j+                         } n |sd   j+                         }xs }t/        j0                  ||      S c c}w )Nc              3  V   K   | ]!  }t        |t              xr t        |       # y wr   rk   r:   r;   )r   r   s     re   r   z0make_pointwise.<locals>.inner.<locals>.<genexpr>  s(      /
;>JsF#6	#6/
   ')r/   r   z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 rN  F)use_compute_types)r   r{   r   rj   r   rR   r   ro   )r;  loadinputs_loaded	inp_indexrq   	inp_dtypedowncastr   emulate_output_castemulate_precision_castsr   rB  loaders	low_pr_fpoverride_fn_when_input_boolrangess          re   inner_fnz/make_pointwise.<locals>.inner.<locals>.inner_fn  s
   u:V,LE7!F8.LL,

"'B'N2W4UTT%[4UVV "'0'9 .OItu+C &y 1 ; ; =I.9	3I#&<<YRW#X!ll8Y?!((-. -(&  #||C%PH<<%88
# 5Vs   D rv   r   r]  r\  )r  rE  ri   mulmake_loaderr   r   rk   r2   BaseConstantr   r{   bfloat16float16rS   r]   r   r^   metagetrI   rn   r   r@   r   )rH  rB  r   otherr]  rv   rs   r   rW  rX  rY  rZ  r\  allow_alphar   override_devicer[  rC  triton_fallbacks    `     @@@@@@re   innerzmake_pointwise.<locals>.inner~  sL   &3 /
BH/
 ,
 #"?"F++"6+@A UaZf U3r
= =,23q1==?3##%%>)<)<)>ABZ 	AEeR__5V I : At1VHAenn.>-?@A 	A ^^U]]3	GG4 X6dBX$$))5X $$))--.OQVW	 	  6L%9:L	 	, F !,,.--.\\^F --/ !*F	
 	
o 4s   !I)rB  rC   r   )r   rC  rh  r[  rg  ri  rj  s   `````` re   make_pointwiserk  v  s     )- L
 L
\ Lrg   c                     ddd fd}|S )Nr/   rG  c                *    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   z1at least one input must be a list to a foreach opc                "    r | diS  |  S )NrH  r   )rt   rg  rH  pw_fns    re   apply_fnz7make_foreach_pointwise.<locals>.inner.<locals>.apply_fn  s     d0%00d|#rg   )r   rS   r]   r^   r_   ra   inplace_foreach_opsrf   rk   ri   r   ro   rx   r  foreach_group_loop)
rH  rB  realize_outputsa_list_inputinputbroadcast_inputsgroupsrp  rg  ro  s
   `       re   rj  z%make_foreach_pointwise.<locals>.inner  s   $$**+q0 0ww##**.AA0-/ 	  	E%$/$	 ' 	
?	
'
  	/EedE]3 ''#l2C(CD !''.	/ $C)9$:;	$ "&#l*;XWWrg   )rB  zlist[list[TensorBox]]r   )ro  rg  rj  s   `` re   make_foreach_pointwiserx    s    45 !X !XF 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>       .q}.   )r   rS   r]   has_featureBackendFeatureFOREACHrealizero   get_operation_nameregister_operation_listr  )rw  num_outputsrp  rs  outputsrv   ru   groupoperation_list
output_indrt   r1  s               re   rr  rr    s     f{"G(. <$u$& % 
	CJd^F"(GJ ##FN,B,BC# %%f&?&?&AB
	C GG++N;<" .g....Nrg   c                    | j                         k(  r|rt        |       S | S fd} t        |      |       S )Nc                4    t        j                  |       S )N)	src_dtype)rR   r   )r   r   r  s    re   	_to_dtypezto_dtype.<locals>._to_dtype  s    ||Au	::rg   rC  )r   clonerk  )r   r   copyr  r  s    `  @re   r   r     sC    IEuQx&Q&; B>)5A!DDrg   r   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>J  r{  r|  )subgraph_loweringr  graph_modulerS   r]   set_graph_handlerrungraph_outputsrx   r   r   r}  r~  r  r  ro   r  r  r  )subgraphrt   r   r  rB  gmpw_subgraphsub_outputsrw  r  rv   ru   r  r  r  r1  s                   re   _foreach_mapr  !  sA    =F			B+BAGGLK	
		[	) ! ! ++K;,Ffs;''G(. <$u$& 	C 
"(GJww""6>+A+AB{ %%f&?&?&AB	C GG++N;< .g....N3! !s   E		Ec                @   |j                   s| j                         j                   rk| j                         r/t        | |      }t        j
                  j                  ||        |S  t        t        j                  j                  d      | |      S t        | |d      S )Nr   Fadd_to_fallback_setTr  )
is_complexr   r   
empty_liker2   InplaceCopyFallbackr   fallback_handlerprimsconvert_element_typedefaultr   )r   r   dsts      re   _convert_element_typer  N  s    1;;=33::< Qe,C""))#q1J#**22  Au4((rg   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                  |       j                  S r   )is_floating_pointr{   finfobitsiinfor   s    re   _get_primitive_bitwidthz1to_dtype_bitcast.<locals>._get_primitive_bitwidthc  s5    "";;u%***;;u%***rg   )	r   r  r  atenviewr   rC   r7   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   rC   r   r2   ComplexViewr{   rR   r  r  r   r  r<  s     re   _view_dtyper  r  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   )rG   rn   r  rC   r   r2   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   Tc	                   |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)rC  r[  rg  ri  )r  r   r  )r   r  )__name__r>   rN   rk  r%  r   r  r   )
r   r   r  r   r  rC  r[  rg  ri  r   s
             re   register_pointwiser    s     #7##D	T	B'!#8 #.&12M&N#	
3$?'
B
	/3	

 	

B ud	
E4  $"7	
 		
 Irg   ldexp)r   rC  )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       @)rR   r   constantpowr_  )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   r  r{   r   rk  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>.frexp1  r  rg   r  c                 0     d   | i | d   | i |fS Nr   r/   r   )rt   r   pw_fnss     re   r   zregister_frexp.<locals>.fn  s.    vay$)&)96!9d+Ef+EEErg   Nr  )
r>   rk  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                8    t        ||      }t        | |      }|S )Nrg  )rx  r  )r   pointwise_lowering_fnrg  r   s       re   register_foreach_pointwiser    s"    
 
  5;	OB	#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   )rR   wherert   s    re   r   zwhere.<locals>.fn	  s    yy$rg   r/   r-   r  r   r  )rk   r6  r   constant_liker   r   r?  rj   rC   r  r   r  r   r2   r   r8   r   ri   r   rk  r   )	r   r   r0  r   rt   r   rs   r   indicess	            re   r  r    so     !eS\"M!Q!eS\"M!Q!Q<DQa.M.U.UE 't_ITQ
1i0HqIGIG.'0JQa0JKL 1Q3t9 Td1gr{{+ ''Qd71:6F6O6O6Q1RSDGT ;>"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     #AQAJJL#A   r   c              3     K   | ]Z  \  }}t         j                  j                  j                  |      t         j                  j                  j                  |      k7   \ y wr   )rS   r]   r   r,  r   r   r0  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  reducer3  r   r  r  r.  ro   )rB  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  8  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   )rS   r]   r   	guard_intr   ds     re   r   zsqueeze.<locals>.<genexpr>J  s%     >Q177##--a0>s   35r/   )rk   rC   rB   r   rm   r   r   r   rS   r]   r   r  r   r   r   r   r'   rj   guard_or_falseEqro   r  )r   r   dims	new_shaper  r   s         re   squeezer  A  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  X  s    C!!rg   c                    t        | |      }t        | t              sJ t        |t              sJ |j                  | _        | S r   )r  rk   rC   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>   rk  r   r   s     re   r  r  f  ?    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  n  r  rg   c                f    t        |       rt        |       S t        d      } t        |      |       S )Nceilr   r  r>   rk  r  s     re   r  r  v  s/    qQx	V	B>"a  rg   c                f    t        |       rt        |       S t        d      } t        |      |       S )Nfloorr  r  s     re   r  r  ~  /    qQx	W	B>"a  rg   c                f    t        |       rt        |       S t        d      } t        |      |       S )Nroundr  r  s     re   r  r    s/    qQx!!~b!!$$rg   c                f    t        |       rt        |       S t        d      } t        |      |       S )Ntruncr  r  s     re   r  r    r  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  )rE  rk   r2   ra  r8   r   r   rC   ri   r   r$   rS   r]   r   r   rP   
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  33M%4HI!" Z&&qvvuU|<==rg   c                    t        |      }|D ]  }d||<   	 | }t        |      D ]  \  }}|dk7  st        ||      } t        ||      S NrM  )ri   rj   	unsqueezer.  )r   shapebroadcast_dimensionsr   broadcast_dimensionr  idxr   s           re   broadcast_in_dimr&    sf    UA3 $!#
$ 	
AA, "Q7!S!A" !Urg   c                6    t        | |j                               S r   )r.  r   )r   r2  s     re   	expand_asr(    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
A   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|$ 	IAqzQA;!#$ww||E!H.uQxHQKHE!H	I rg   r^  )ri   r   r   r   r*  r+  r  r  emptyr   rn   r  r  r  r.  r$   rS   r]   r   r   rP   r  r`  r@   r   )	r   r1  new_sizezero_tensorrs   r]  old_size_productr0  r2  s	    `     @@re   repeatr7    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++>>}X?VWa LL  33M(4KL#$
 }}H||~kkmH~	 rg   c                T    t        t        j                  | j                  |            S r   )rC   rE   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   rC   ri   r   r?   r   rm   )r   r  s     re   permuter:    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      }
 |	|||      }|
||
|}}d}|s8t	        t        j                   j#                  | j$                  |||||            S |sJ  |t        j                  j                  j&                  t        j                  j(                  j*                  d         }|J t-        |      d	k  sJ |       d\  }}|j/                         D ]7  \  }}| |d
      t1        j2                  |      fk(  r|})| |d      fk(  s6|}9 |
|J t        j4                  ||||| j                         |         }t        j                  j7                  |      |_        t        j                  j;                  |       |}| j=                         | j?                          |
6|J | jA                         jB                  |
| jE                         |   z  z   }nt        jF                  ||| jA                         jB                  | jE                         |   | j                         |   d      }t        j                  j7                  |      |_        t        j                  j;                  |       |}tI        | j                               }tI        | jE                               }|||<   ||xx   |z  cc<   tK        | |||      S # t        $ r Y w 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                  | d            r |t        j                  | |            r| S  |t        j                  | d            r" |t        j                  | |             r| |z   S  |t        j
                  | |            r|S  |t        j                  | |             ryy )Nc                T    t         j                  j                  j                  |       S r   )rS   r]   r   r  r   s    re   <lambda>z5slice_.<locals>.compute_slice_index.<locals>.<lambda>+  s    qww''66q9 rg   r   )r   r.  GeLeLtGt)r;  r   r  r   s       re   compute_slice_indexz#slice_.<locals>.compute_slice_index'  s    =N9U#||D!ehhua !b%)>&?L"#588ED5+A(B4<%&K&'rg   NNFclampunbacked_bindingsr-   r   storage_offsetTr   )&%torch.fx.experimental.symbolic_shapesr<  r&   rk   rC   _validate_dimr   r   r.  r   rS   r]   r   statically_known_leq	TypeErrorr2   	SliceViewr   rm   	shape_envr^   rd  r   r   pytreeSequenceKeyDynamicSliceSizeregister_bufferr   register_operationmaybe_get_layoutr  
get_layoutoffset
get_strideDynamicSelectStorageOffsetri   
as_strided)r   r   startendsteprG  r<  r&   r   rD  start_index	end_indexambiguous_slicerH  sym_sizesym_storagesymkeypathb_sizer4  new_storage_offset	b_storage	new_sizesnew_stridess                           re   slice_rj    sp   
 a###
3
"C::<D<<DdEJJ'4!89T93QJ  55dC@	H
" (KO)%q9'T48	"y'<$i3E#O 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	 )"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   75N 	NN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   rC   rm   r2   r6   rn   r   unwrap_viewr  is_storage_and_layoutr   as_storage_and_layoutFixedLayoutrv   r   r.  r   )
r   r   striderI  
new_device	new_dtypestorage
old_layoutr   
new_layouts
             re   rZ  rZ    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   rC   rZ  rm   )r   r   rr  rI  s       re   as_strided_ry    s/    a###48==AFHrg   c                4    t        | |||      }t        |      S r   )rZ  r  )r   r   rr  rI  results        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   rM  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   r?  z1pointwise_cat.<locals>.inner_fn.<locals>.<lambda>  s    -N1-h7 rg           rM  r-   )rR   
index_exprr{   int64r  r   r  geltand_ro   ri   r*   maskedr  )r%  idx_dimmasksmasked_loadsr[  r\  
start_condend_condmasknext_valrs   r  r   rB  r  inputs_rangess             @@re   r]  zpointwise_cat.<locals>.inner_fn  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1 	AyyaQH	 rg   r^  )ro   r   r`  ri   r@   r   rn   r   )rB  r   prev_endr   r]  r4  r  r  s   ``    @@re   pointwise_catr    s    9;MH )h3<<>#3F(FGH $R() 4::Ccoo':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 Nr   )_create_constantsr{   r  r   rR   r   r  
reciprocalr  maximumminimum)r%  channel_idxru  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{   rb  r   r  r`  r@   r   rn   )ru  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   )r2   ComputedBufferforce_realizerR   device_assert_asyncr`  )r;  r   r   s    re   r]  z_assert_async.<locals>.inner_fn'  sO    ,,. 	K**+=4+;+;+=e+DcJ	K 	K 	K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  5      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  rR   r   sub)r%  r  ru  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   )ru  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  rR   r  r  r  r   )r%  r  r  ru  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{   rb  r   r  r`  r@   r   rn   r  r$  r6  r   r   )ru  r  r  r  r  r   r]  r  s      ``` @re   0quantized_decomposed_quantize_per_tensor_defaultr  t  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  rR   r  r   )r%  r  r  ru  r  r  r  s        re   r]  zDquantized_decomposed_dequantize_per_tensor_default.<locals>.inner_fn  s]    S!-eZu}}Uzggcll5%--8*EMll3	*
rg   r  r^  )r   r{   r  r`  r@   r   rn   r  r$  r6  r   r   )	ru  r  r  r  r  r   r  r]  r  s	         ` @re   2quantized_decomposed_dequantize_per_tensor_defaultr    s     ??% 
(/@AR@ST% MM	$$&L !""E%LS_
 ~~ rg   c                0  	 | 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                         		fd}t        j                  | j                         || j                               S )Nr  r   r/   expect scale as scalar tensor"expect zero_point as scalar tensorc                    	|       } 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                  |      z        |z   }t        
t        j                        \  }}t        j                  t        j                  ||      |      }t        j                  |      S )Nr/   r   r   r   )r   r   r   r{   r  rR   r   r  r  r  r  r  )r%  ru  _scale_zero_pointr  r  r  r  r   r  r  r  r  scale_loaderr  zero_point_loaders           re   r]  zAquantized_decomposed_quantize_per_tensor_tensor.<locals>.inner_fn  s    S!c%..*:&;q&@dbI'ENN4D0E0JPRS;;%--'\\&%--8Fu}},,,{EMMBKiiv 667+E&y)5==Q
d++ckk#t4d;||GU++rg   r^  )r   r{   rb  r   r  r   r   r`  r@   r   rn   )
ru  r  r  r  r  r   r]  r  r  r  s
    ````` @@@re   /quantized_decomposed_quantize_per_tensor_tensorr    sf    ENN*.??- 
FuGXFYZ- u~~ A%ENN"u~~'7':a'?'&'  z""$%*J!"a'J,?,?,A!,D,I,+,  $$&L$$&L"..0, , !~~	 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  rR   r   r  )r%  ru  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  )ru  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        ||       } }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 |	}
t        fd| D              xr t!        fd| D               }|
s|r|st'        | |      S t)        t*        j,                  j/                  | |            S c c}w )Nr   r   c              3  |   K   | ]4  }|j                         t        j                  t        j                  fv  6 y wr   )r   r{   int8uint8r   ru  s     re   r   zcat.<locals>.<genexpr>'  s.      ;@ejj%++66s   :<c              3  T   K   | ]   }t        |j                               d k(   " yw)r   N)r   r   r  s     re   r   zcat.<locals>.<genexpr>.  s!     >es5>>#$)>s   &(r/   r   c                   t        | t              rJt        | j                  t        j                        r| j                  j                         S | j                  S t        | t        j                        r| j                  S | S r   )rk   rC   rm   r2   r6   rn  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   r2   r  rm   rA   )rw   s    re   is_reductionzcat.<locals>.is_reductionG  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   )rS   r]   
get_buffer)r   readcan_fuse_reductions     re   r   z2cat.<locals>.can_fuse_reduction.<locals>.<genexpr>P  s-       #177#5#5d#;<s   /2)rk   rC   r2   r   r@   r  get_read_names)rw   r  r  r  s    re   r  zcat.<locals>.can_fuse_reductionJ  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>W  s     Ba.q1B   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)	r2   ro  rp  ConcatKernelcan_realize_into_without_copyrk   rC   r   r@   )r   ru  _should_lower_cat_inputr  s      re   r  z#cat.<locals>.should_lower_cat_inputY  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   rC   r2   r   r@   inner_fn_opcountnum_opsr  rS   r]   r  )r   countr  op_countr  s      re   r  zcat.<locals>.op_countp  s    a)R]]34M!,-- !R\\*""$,,$$& 	8DXagg00677E	8 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     Cq22Cs   c              3  6   K   | ]  }t        |        y wr   )rJ   )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>  s     >&s+>r  c              3  .   K   | ]  } |        y wr   r   r  s     re   r   zcat.<locals>.<genexpr>  s      "
,/"3'"
r  c              3  .   K   | ]  } |        y wr   r   r  s     re   r   zcat.<locals>.<genexpr>  s     <(+<r  )r   zUnion[TensorBox, ir.StorageBox]r   	ir.IRNoder   r   r`   torch._ops.OpOverload)rn   r   r  r  require_channels_lastr  r  r  r  r   r  rK  r   r   r?  r   r  r0   force_pointwise_catr  rC   r2   r  r   max_pointwise_cat_inputsrS   r^   r_   )rB  r   
cpu_deviceru  r  r   r   fusable_reductionMAX_COMPLEX_POINTWISE_CATpointwise_usesfuse_pointwise_usehorizontal_fuse_catr  r  r  r  r  r  r  s               @@@@@@@re   r  r  $  sI   %%',,5Jc DJ 
  	EMMO	>v>>-dhh@@IFA1 0 01&#>>
6{aVAY
q	3
*C	%D%L%LE /55shsE"5F5
U

 B6BB !!VS)) //<==$ !"F 6{//	V77	7CFCC 
~~++
 
 >v>>Q> 	 " "
39"
 
 =<V<<< 	 "5>O --R__++FC899K 6s   5I1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)r%  rankc                     d  d S )Nz(diagonal dimensions cannot be identical z, r   )dim1dim2s   re   r?  zdiagonal.<locals>.<lambda>  s     HbQUPVW rg   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 )NrM  r   r/   r-   )r   r  )
r%  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   rS   r]   r   evaluate_exprr   rB  evaluate_maxevaluate_minrj   ro   rC   r2   GenericViewr   )ru  rW  r  r  offset_negative	diag_sizers   r   r  r  r  r  r  s     ``      @@@re   diagonalr"    s   ^^%N>"H84D84D	W gg&&44UXXfa5HIOGG$$11GG))t$v-~d/C 	
	 GG$$11GG))t$nT&:V&C 	
	 HGQ<v;$^4N41a$8MQNEN	LL   R^^**5%CDD' Os    F:/F:c                0    t        t        | |||            S r   )r  r"  )ru  rW  r  r  s       re   diagonal_copyr$    s    %t455rg   c                P    t        |       }t        ||||      }t        ||       |S r   )r  r"  	mutate_to)ru  srcrW  r  r  r1  ra   s          re   diagonal_scatterr(    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  rH  )r   r.  r   rS   r]   r   r  rB  r@  r%   r  rX  rV  rW  rZ  rj  r  r&   rO  r^   rd  r   r@  iterr   r2   rY  rS  r   rT  )r   r   r%  r   actual_indexr4  
new_striderf  slice_resultrH  unbacked_offset_symr  buffers                re   selectr0    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/   rM  FrF  )rK  rk   ri   r   r   rS   r]   r   r  r)   ro   rj  )
r   r  r   sizes_x_sizechunksr{  r[  r   r\  s
             re   splitr5  &  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   )r5  )r   r  r   s      re   split_with_sizesr7  A  s    E3rg   c                    t        | |d      }t        j                  j                  j	                  | j                         |         }t        |      D cg c]  }t        | ||       }}|S c c}w r  )rK  rS   r]   r   r  r   r  r0  )r   r   r3  rs   r{  s        re   unbindr9  F  s_    
3
"CWW''

S(9:F).v7AfQQ7F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 )NrM  r/   r   )r%  dim_idxr   r]  s     re   r  zunfold.<locals>.reindexerd  s?    b'CHtO+8Tc8G8c#'B&788rg   )r   r   r   rj  r!  rS   r]   r   	check_leqcheck_ltr)   r   r  r(   rC   r2   r  r   )r   	dimensionr   r]  r  ndimdim_sizer   new_dim_sizeout_sizer  r   s      `       @re   unfoldrD  N  s   JJLEu:D
4
+Cqyi1o4u==SzHwwHtX&aHtOT2Q6L""8,q0	''t0CX(NO	
 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  )rK  ri   r   insertr   r*  r+  r  )r   r   r  s      re   r!  r!  k  sE    
3
"CQZZ\"IS%''++&9rg   c                    t        | |      }t        | t              sJ t        |t              sJ |j                  | _        | S r   )r!  rk   rC   rm   r  s      re   
unsqueeze_rH  s  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  )	rS   r]   r   rO  r  r   sympifyr   r   )r   r   rW  r@  s       re   rK  rK  |  sr    
''


$
$
2
25==3E
FCqzz|D
Qwtf}#dVm#####Jrg   rM  c                   t        | |d      }t        j                  j                  j	                  | j                         |         dz  }t        | |d|d      }t        | |||dz  d      }t        |t        |            S )Nr   r-   FrF  )	rK  rS   r]   r   r  r   rj  r_  sigmoid)r   r   new_lenr   r0  s        re   glurN    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                d    t        | t        j                        rt        j                  |       S | S r   rk   r2   r:   rC   r   r   s    re   wrap_tensorsz7fallback_handler.<locals>.handler.<locals>.wrap_tensors  %    *4Q		*B9##A&IIrg   )rP  tree_mapr2   FallbackKernelr   )rt   r   rS  kernels      re   handlerz!fallback_handler.<locals>.handler  s<    	J "++226KDKFK
 	
rg   T)r  r   _is_fallback_handler)rW  r  rX  s   `  re   r  r    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]    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_mmrK   )rw   rc   s     re   unsupported_input_tensorrc    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{   rR   r  r  r  ra   r  rc  is_cpur0   disable_cpp_codegen)rw   rc   supported_complex_viewss      re   unsupported_output_tensorrh    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{   fxNoderd  rP  tree_leaves_subclasses
FakeTensorrh  rc  )inp_out_node	is_outputrd  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		  rg   )rr  T)
ra   r  view_as_complexr  r`   lift_fresh_copyrP  arg_tree_leavesrt   r   )rc   allow_cpu_inputsrs  r   s   `   re   %fallback_node_due_to_unsupported_typerx    s    {{d**222ww- {{d**222* %%tyy@DKK@ u5  55rg   c                   t         j                  xrE | t        j                  t        j                  j
                  t        j                  j
                  hv }| t        vs|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.                  j0                        r*| j3                         D ]  }t5        | |      } ||        y t-        | t        j.                  j6                  t        j.                  j8                  f      r	 ||        y t;        d|  d	t=        |              )
Nz*both a fallback and a decomp for same op: CIFznA make_fallback error occurred in suppress_errors config, and suppress_errors is being disabled to surface it.zmake_fallback(a.  ): a decomposition exists, we should switch to it. To fix this error, either add a decomposition to core_aten_decompositions (preferred) or inductor_decompositions, and delete the corresponding `make_fallback` line. Get help from the inductor team if unsure, don't pick arbitrarily to unblock yourself.c                n    t        |        t        |         t        | d       t        |             S Nr  )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 )r0   rX  r  addcmul_foreach_addcmulScalar_foreach_addcdivr4   r   osgetenvr5   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skip_decomp_for_precisionr  olr}  s    `      re   make_fallbackr  	  s    !' > > !2$$$$J D
 ^#:S 
4RD9S 	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   tensorr{   r  )r"  numelr   s      re   philox_rand_offsetr  <	  s2     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   )rR   r   r{   r  r   r  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_fnU	  s     ,,{2DLLr):EKKH''NN:e,ekk:<M
 
 ||FE**rg   r^  )
r2   rq  FlexibleLayoutcontiguous_stridesmake_indexerr`  r@   r   ri   r  )r   seedrW  rr  rv   r   r]  random_values_nodeoffset_noder  r  r  s        `   @@@re   philox_randr  H	  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)r0   r  rP  rU  rC   r   r2   rV  r  native_dropoutr  r   )r   ptrains      re   r  r  o	  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   )r0   r  rn   r{   rv   r  r   rk   r6  r  
bernoulli_Tensorr2   InplaceBernoulliFallback)r   rt   r}  s      re   r  r  z	  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  )r0   r  rn   r{   rv   r  r  )r   rt   s     re   bernoulli_pr  	  sG    !!Q\\^u||E7J%J ^J eAh&&&rg   c                    t         r   r   )r  s    re   _foobarr  	  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  rS   r]   creation_timer   rg   re   warn_triton_randomr  	  s    --.rg   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)re  fallback_rand_generatorr0   r  popfallback_rand_defaultr   rt   r   s     re   r  r  	  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  )re  fallback_randn_generatorr0   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   )r2   get_stride_orderExternKernelrequire_stride_order)input_tensorrr  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  rC   r   r2   RandomSeedsrG   )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   )rR   	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   rW  c                 	 t         j                  rJ dv sJ g | } t        j                  }|j	                         }t        j                  ||| t
        j                  j                  |       |      j                         |j                         		fd}t        j                  |||g |       }|j                          |S )N)r  r  r  c                     t        t               g       t        j                   |       t        j                              S r   )r   rR   r  r{   r  )r;  moder  r  s    re   r]  z!inductor_random.<locals>.inner_fn	  s7    !wsD!ONN:e,ekk:
 	
rg   r^  )r0   r  r{   r  rA  r2   rq  r  r  r  r`  r@   r   r  )
r   r  r  rW  r   rv   r]  r{  r  r  s
     `     @@re   inductor_randomr  	  s    %%%%$$$$T7DMME%%'FtR..AA$GPVln  ""$K
 w	F NNMrg   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   )rR   	randint64r  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^  )r0   r  r{   r  rA  r2   rq  r  r  r  r`  r@   r   )
r  r  r   r  rW  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r,  r   )r   r   sts      re   r   z%_boundaries_helper.<locals>.<genexpr>
  s     Aeaa!er\As   rM  )r   rX  sumr  r  )tbr   rr  
max_offsets       re   _boundaries_helperr  
  s^    
 ;;=D]]_FAs4/@AAF2JNJ
Rr
	 rg   c                H    | j                         | j                         d   fS r   )r  rX  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   )rS   r]   r}  r~  	BUCKETIZEr  s    re   r?  zsearchsorted.<locals>.<lambda>5
  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rR   	bucketizer  r  )r%  r  index_dtyper  sorted_sequencer  values_loaders     re   r]  zsearchsorted.<locals>.inner_fnZ
  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>q
  s     &UAq1u&U   rM  )rX  rR   r  r  r  operatorr   r  )r  stridesr%  r  s     re   get_flattened_indexz;searchsorted.<locals>.inner_fn.<locals>.get_flattened_indexm
  sU    --/~~$$ &UWSb\3sPR89T&U  	 rg   r  )r  rC   r  )r%  r  r  r  r  r  r  r  s   `  re   r]  zsearchsorted.<locals>.inner_fnh
  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"  )r  selfr  r  r  r  validate_bucketizer]  rv   r{  r  r  s   `  ` `    @@re   r  r  +
  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  r
  c                \     |       }t        j                  |t              d      }|S r  )rR   r  r  )r;  r  r  
boundariesr  r  r  s      re   r]  zbucketize.<locals>.inner_fn
  s6    5!--z*
 rg   r^  )r   r   rS   r]   r}  r~  r  r  r  r  r  r  rn   r`  r{   r  r  r@   r   )	ru  r  r  r  rv   r]  r{  r  r  s	    ` `   @@re   r  r  
  s     z""$%*** 	
E>#;#;<GG
N,D,DEQ 5 55Q:%
 	
 F$$&L!*%++K
 ~~	F NNMrg   c                    t        j                  t        j                  t        j                  j
                  ||f      \  }}||fS r   )rP  tree_map_onlyr2   r:   r  require_stride1r  rt   r   s      re   require_denser  
  s;    ''
		2??22T6NLD& <rg   c                    t        j                  t        j                  t        j                  j
                  ||f      \  }}||fS r   )rP  r  r2   r:   r  require_contiguousr  s      re   r  r  
  s;    ''
		2??55f~LD& <rg   c                    t        j                  t        j                  t        j                  j
                  ||f      \  }}||fS r   )rP  r  r2   r:   r  r   r  s      re   r   r   
  s=     ''
		2??==f~LD& <rg   c                    t        j                  t        j                  t        j                  j
                  ||f      \  }}||fS r   )rP  r  r2   r:   r  r  r  s      re   r  r  
  s;    ''
		2??884.LD& <rg   c           	        || S t        |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!      
19!S$Q,
r-  )rk   r   r2   r:   r  require_exact_stridesrr  r  r  r   ri   r   r  )r   fake_argkeys      re   r  r  
  s    
(,-
#ryy!44S(//:KLL#tRUV3-c#hFFVV	C%	'tCy 
=@h=O
 
 	
 J Ws   3C
c           
         t        d t        | |      D              } |j                         D ci c]  \  }}|t        |||          }}}| |fS c c}}w )Nc              3  :   K   | ]  \  }}t        ||        y wr   r  )r   r   r  s      re   r   z,constrain_to_fake_tensors.<locals>.<genexpr>
  s#      C 	!h/r-  )r   r  r   r  )rt   r   	fake_argsfake_kwargsr   r  s         re   r   r   
  sb      y1 D JPXAa)![^<<XF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        | t        j                        rst        j                  |j                  d   j                         t        j                  j                  j                        }t        j                  j                  | |      S t        | t              r| D ci c]  }| | |   ||          c}S | S c c}w Nr  )rk   r2   r:   r  rd  rr  rS   r]   r   rO  r  r  r  )r   fx_argr  r  apply_constraints       re   r'  z1constrain_to_fx_strides.<locals>.apply_constraint  s    c299%..E"))+QWW-=-=-G-GL ??77\JJc4 LOPSC)#c(F3K@@PP
 Qs   #B?c              3  6   K   | ]  \  }} ||        y wr   r   )r   r   r&  r'  s      re   r   z*constrain_to_fx_strides.<locals>.<genexpr>  s"      *5#vf%r   )r   r  rt   r   r   fx_nodert   r   r   r  r'  s        @re   r   r     ss      9<T7<<9P D EKLLNSDAqa!!W^^A%677SFS< Ts   !A*c                      fdt        fdt        t        | j                              D              }|j	                         D ci c]  \  }}| d| j
                  |          }}}||fS c c}}w )Nc           
     
   t        |t        j                        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}|j.                  s t        j0                  j3                  ||      S dj&                  t
        j4                  j(                  j6                  j,                  k(  xr | dk(  }t        |t8              sJ t#        |j%                               d	vr|S t        j:                  |      }	|	r3t        j<                  t        j0                  j?                  |      |      S t        |t              rE|jA                         5|	r3t        j<                  t        j0                  j?                  |      |      S |rt        |j%                               }
g }|jA                         }t!        t#        |j%                               d
z
        D ]q  }t        j                  j                  jC                  ||   d      s1|3t        j                  j                  jC                  ||   d      sa|jE                  |       s dgt#        |
      z  }d
|d<   d
}t!        t#        |
      dz
  dd      D ]  }||d
z      dk7  r||
|d
z      z  }||v r9t        j                  j                  jC                  ||d
z      z  d      rd||<   Vt        j                  j                  jC                  |z  d      stG        |      z  }|||<    t        j0                  jI                  ||      S |	r3t        j<                  t        j0                  j?                  |      |      S t        |t              rE|jA                         5|	r3t        j<                  t        j0                  j?                  |      |      S fd}t        |jJ                  t        jL                        rQ ||      sI ||jO                               r3t        j<                  t        j0                  j?                  |      |      S t        j0                  j3                  ||      S c c}w )Nr  rM  r   )r   r   r   )r   r/   r-   r   r   r   r   r   r/   r-   c                    t         j                  j                  j                  t	        j
                  t        | j                         d         d            S )NrM  r   )rS   r]   r   r  r   r   r+   r   )r   	ALIGNMENTs    re   
is_alignedz=sdpa_constraint.<locals>.apply_constraint.<locals>.is_aligned  s@    77##22QZZ\"-y91= rg   )(rk   r2   r:   rd  rr  r{   SymIntrc   exprrS   r]   r   rO  r  ri   r)  r  r   r   ra   r  0_scaled_dot_product_efficient_attention_backwardr  is_cudar  r  rR   '_scaled_dot_product_efficient_attentionrC   is_aligned_realized_tensortry_match_insignificant_stridesrealize_inputmaybe_get_stridestatically_known_equalsro   rF   r  rm   r6   rn  )r%  r   r&  meta_valr   meta_stride_exprrO  r  effn_attn_fwd_biasis_aligned_tensorrC  expanded_dimsmaybe_striders   out_stridesrr  r0  r/  r*  s                    @re   r'  z)sdpa_constraint.<locals>.apply_constraint  s   #ryy)J;;u%GOGX
BC:a6AFFKKA=
 
 GG$$..	**8??+<iHL,1s3<<>/B)C DEL NNDDLLMv|$)))
 (L??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!((+, $X.KKOF3x=1,b"5 (q1u%*#hq1uo5F %ww''??#AE*Y6 *+A ww''??@RTUV$VY7)CF!'A!($ ??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Fo
s    9Uc              3  >   K   | ]  \  }\  }} |||        y wr   r   )r   r%  r   r&  r'  s       re   r   z"sdpa_constraint.<locals>.<genexpr>  s*      C#v 	c6*s   rM  )r   rj   r  rt   r   r   r)  s   `    @re   sdpa_constraintrC    s}    |G|  "+Cgll,C"D D IOW1a!"a):;;WFW< Xs   "A5)r\  )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 Nr*  )rk   r2   r:   r  r   rn   r  r   r   r.  r  )r  r'  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   rF  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 )NrM  rl  )rk   rC   rm   r2   r   ro   rV  r  )r   reinterpret_view_layoutsrm  s      re   clone_preserve_reinterpret_viewrK    s    !!YJqvvr7I7I$JFFB../$++ALLN;A B../ aLaAFF.tt4 	:F""&9A	:aLHrg   ru  c               \    fd}t        j                  t        |      || g      S )Nc                D    t        j                  | d   z  z         S )Nr   r   rR   r  )r;  r   r[  r]  s    re   r   ziota.<locals>.fn  s!    ~~dU1Xo5UCCrg   r^  )r@   r   rG   )lengthr[  r]  r   rv   requires_gradr   s    ```   re   iotarQ    s1    D V$x	 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   )rR   r  eqr  r{   r  )r%  r   r;  
src_loaderr2  s    re   r]  z select_scatter.<locals>.inner_fn  sX    yyFFs3x5uekk2 sOSM
 	
rg   r^  )r   r   r`  rK  rS   r]   r   r  r   rB  r   r@  r  r  select_scatterr  r=  r>  r.  r!  r@   r   rn   ri   )r   r'  r   r;  r]  rU  r2  s     `` @@re   rV  rV    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_idxrU  s   re   r?  z1slice_scatter.<locals>.inner_fn.<locals>.<lambda>  s    Jw' rg   r  )rR   r  r{   r  ri   r)   ro   r  r   r.  r  rT  r,   r  r  r  r  r  r   r  )r%  r  r  src_valrZ  r   rA  r\  rU  r[  r]  r   r2  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`  rK  r   r2   rN  normalize_start_endri   r)   r.  r@   r   rn   )r   r'  r   r[  r\  r]  src_sizer]  rA  rU  r2  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`  0  s-    !dE]#A
qt}Hrg   r   rv   rm  
pin_memoryc                   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=rb  c                0    t        j                        S r   rN  r;  rm   r   s    re   r]  ztensor.<locals>.inner_fnC  s    >>$..rg   c                0    t        j                        S r   rR   r  rf  s    re   r]  ztensor.<locals>.inner_fnH  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   )rR   r  r  r  r  r{   r  )r[  r\  midbinary_searchrm   r   r;  s      re   rl  z/tensor.<locals>.inner_fn.<locals>.binary_searchP  s    s{"{;!#<<UU;;U{q(50yyFFuQx=S%++6 "%-!#s+ rg   r   )r   rR   r  )r;  rl  rm   r   s   `@re   r]  ztensor.<locals>.inner_fnO  s7     4yA~||Au-- CI..rg   r*  r^  )r   r{   stridedrk   r`  r   r  get_default_dtyper   r   r6  r   ro   IntegerrS   r]   add_tensor_constantr  r@   r   rG   )rm   r   rv   rm  rb  r\  r]  s   ``     re   r  r  6  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 rE  )rk   rC   r   r  r  )rm   r   rv   s      re   	as_tensorrr  o  sD    $	"D%(DT6*D$eF33rg   c                8    t        | t        j                        S r  r  )rm   s    re   long_tensorrt  z  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 )NrH  r/   r  )r&   rS   r]   r   rO  r^   rd  r   r@  r*  r   r2   DynamicScalarrS  r   rT  rk   r{   r1  SymFloatSymBoolrc   r2  r   rJ  )rm   rH  binding_symrd  r/  r  s         re   _local_scalar_denserz    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   rm  c                    y r   r   )r   r   rr  r   rv   rm  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   rh  r;  r   r   s    re   r]  z_full.<locals>.inner_fn  s    <<u--rg   c                0    t        j                        S r   rN  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   r6  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   rm  rb  rF  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 tensorsrd  rb  r/   r   )r   r{   rm  rG   rn  r   rk   ri   r   Sizer1  r   r.  r  )	r  r   rv   rm  rb  rF  r   r   r  s	           re   rj  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  	3A!!U\\222	3)-.AQ..Z55 /s   C8r   )r  rj  s   ` re   r  r    s#     6. Lrg   )r  r   rm  rv   rb  rF  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   rm  rv   rb  )	r   rG   r   rk   ri   r   r{   r  empty_strided)r  r   rm  rv   rb  rF  r   s          re   r3  r3    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   rm  rb  rF  c                  t        | d       t        |d t        j                  fv d|        || j                         }nt	        |      }|xs | j                         }t        | j                               } |||||      S )Nrb  rd  ra  )r   r{   rm  r   r   rn   ri   r   )r   r   rv   rm  rb  rF  r   creation_fns          re   _constant_likez*create_tensor_like.<locals>._constant_like  s     	z><06dEMM22gfX4FG=KKME 'E)1<<>AJJL!fV

 	
rg   r   )r  r  s   ` re   r  r    s     duTX
 rg   c                *    t        t        |             S r   r  )r  s    re   r  r  %  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 )Nrb  rd  )rk   ri   r   r   r{   rm  r   r   rn   r   ro  r  rG   )r   r   r   rm  rv   rb  r   r  s          re   _new_constantz#new_constant.<locals>._new_constant/  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  .  s    tDT
E rg   r  c          	     ~    || j                         }|| j                         }t        |d ||t        |      |      S Nr  r   rn   r  rG   )r   r   r   rm  rv   rb  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 )Nrd  r  r   )r  rv   r   r   )r\  F)rv   r   r   rr  	is_pinned)rk   ri   r   r   r   r{   rm  r   rn  r  rv   rG   r  r  rm   dataclassesreplacer   r2   r  r   r.  r  r  rq  rm  )	r   rr  r   rm  rv   rb  	pointwiser/  r   s	            re   r  r  N  sw    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%)*ELLO*D*  #))Qa)11$7 
 NN%FM  +)s   3F:F?c          	     ~    || j                         }|| j                         }t        ||||t        |      |      S r  r  )r   r   rr  r   rm  rv   rb  s          re   new_empty_stridedr  o  sH     }~V$ rg   c                
   |D cg c]+  }t         j                  j                  j                  |      - }}t	        t        t        |            |j                        }t        j                  j                  | |      S c c}w )N)r  )rS   r]   r   r   sortedr  r   __getitem__r2   r  r  )r   rr  r   r  s       re   copy_stridedr    sb    >DEagg11!4EFE%F,&2D2DEL??//<@@ Fs   0B c                V    |j                  d      J d        t        |      | fi |S )Nr   z(dtype should be handled by decomposition)re  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   rR   indirect_indexingr   )r%  
gather_idxr   index_loaderr   r2  s     re   r   zgather.<locals>.fn  sL    3i**<+<d3iH
s8q=,C!CH}rg   r^  )rk   rC   	get_numelr  r   r   rK  r.  r`  r@   r   rn   r   )	r   r   r;  sparse_gradrW  r   r  r   r2  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   rR   r  )r%  	var_index
weight_idxindices_loaderindices_ndimr4  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   rC   strr   r`  r   r   r@   r   rn   )weightr  padding_idxscale_grad_by_freqsparser   r  r  r4  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      = 	
%++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>  s1      78q}%**ekk22s   ?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   rC   r   r  r   rn   )r  rv   rs   r   
valid_idxsnew_indicess         re   check_and_broadcast_indicesr    s/      
 4G4eqWXWdQ[[]4e3fg   <C  ""=>> )' 2O1jI6N!OJOz?QD DD&3w<'KJ 1
3S1GAJ3S TU 1 <<>V#%&VWWA 
""# 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   rM  r/   r   wrap_neg)r   r  ro   rR   r  )r%  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  r2  s            re   r   z*index_output_size_and_inner_fn.<locals>.fn  sN   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   )r3  r  r  r  r  r  r2  r   r  previouscurrentrs   r  r  r   r  r  s    ````````      @@re   index_output_size_and_inner_fnr    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 rH  )index_impl_helperr@   r   rn   r   )r   r  r   r  r]  r  s         re   
index_implr  8  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   )r%  index_inner_fnr2  s    re   r]  z#index_impl_helper.<locals>.inner_fnb  s    s+,,rg   )rk   ri   r   r`  r  rn   r   r   r  
IndexErrorr  )r   r  r   r  r  rs   r  r  r3  r  r  r]  r  r2  s               @@re   r  r  C  sS   ge}---}}H9'1<<>RG^~"F$FF"KRSa!-q}}TASOS w~a01::<=KZZ\F',S\':U!gaj>TF1IUL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;  h  sT    
!WD11 
			M

 1 1uMw
 	

s    AAAc                    t        | |d      S )NFr  )r  r  s     re   _unsafe_indexr  t  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 )NFr  r  r  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   r&  r  )r  r  r   r  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_rS   r]   r^   ra   _overloadnamer0   graph_partitionrd  re  disable_cudagraphs_reasonr2   IndexPutFallback)	r  r  r   r  r  r}  r*  r   r  s	            re   index_put_fallbackr    s    C$//177+?+?+F+F+T+TUK
 gg""G""8A\!,,**=$??;?E*;-8C,/)T7FJGKrg   c                $    t        | |||dd      S )NTr  r  r  r  r   r  s       re   r  r    s    gvz4 rg   c                $    t        | |||dd      S )NFTr  r  r  s       re   _unsafe_index_put_r    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)Nr*  F)rk   rC   rm   r2   r6   rn  r   r  r   r*  ra   r{   rR   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   r  r	  s     re   r   z"index_put_impl_.<locals>.<genexpr>  s      H
39&v.H
r  r/   r   rM  r  
atomic_addrv   r   r]  r\  output_indexerscatter_moder   rm  rm   )%r2   try_get_namer  r  r  r  r   r   r{   r   r  r  r   r!  r  $are_deterministic_algorithms_enabledr  rL   r  r   r  rn   r   r`  rk   rC   ri   r  r.  Scatterr  MutationLayoutSHOULDREMOVErS   r]   rS  r   rT  )r  r  r   r  r   r  r  r  r;  r3  x_ndimr  rs   r  r  r  expected_vals_sizer]  rv   scatterr/  r	  s                        @re   r  r    s   	 ??4 F$9$9$;;C H
=DH
 E
 NN 	aLAAJ  "uzz5;;&??qzs4==?+S-AB 	'AT2&D	''tffjII 113!$DD  I!2uzz5;;6O!O%dGVZHHI ]]_F[FB4>>CSTQ;qc?D!$DQ;b>Dfdnn./FE"=T__&#
 LSSa!-q}}TASO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"~Ke  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_fnr%  self_loaders   re   r?  z8_unsafe_masked_index.<locals>.inner_fn.<locals>.<lambda>I  s    K8H8M,N rg   )r   r{   r   rR   r   r  )r%  mask_valr  fillr  mask_loaderr  s   ` re   r]  z&_unsafe_masked_index.<locals>.inner_fnD  sK    ::#||K$4ejjAH"3'Hzz($NPTUUrg   r^  )r  r`  r@   r   rn   r   )
r  r  r  r  r\  r  r]  r  r  r  s
    ` `   @@@re   _unsafe_masked_indexr  <  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)r  )r  r   r  r   rG  r  )r   r  r  r   masked_valuer"  rs   clamped_indicess           re   #_unsafe_masked_index_put_accumulater#  S  s    vq)LJJLE s7|$ 7>ajgaj58)U1X\2dJO  Q$OOs   (A-c                V    t        j                  |t        j                  ||             S r   )rR   r  r  r   minmaxs      re   rG  rG  `  s    ;;sCKKQ/00rg   c                P    t        |       }t        ||||      }t        ||       |S r   )r  rZ  copy_)r  r'  r   rr  rI  r1  output_views          re   as_strided_scatterr+  e  s*    4[FVT6>BK	+sMrg   c                2    t        t        |       |||fi |S r   )scatter_r  )r   r   r;  r'  r   s        re   r  r  m  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 implr.  )rk   rC   rQ   r   r
   r{   r   r   rn   r2   ScatterFallback)r}  r  r   r;  r'  r  r/  src_is_tensors           re   scatter_fallbackr3  r  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   multiplyr4  r   r  r6  prod)
r   r  r-  rS   r]   r^   ra   r  r3  scatter_reduce_)r  r   r;  r'  r  r}  fallback_results          re   r-  r-    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;  r'  s       re   scatter_addr=    s    a#uc22rg   c                     t        | |||d      S )Nr  )r8  r<  s       re   r;  r;    s    1c5#u55rg   c                4    t        t        |       ||||fi |S r   )r8  r  )r   r   r;  r'  reduction_typer   s         re   scatter_reducerA    s    58S%nOOOrg   )r/  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  r7  meanamaxaminr/   r  zKaten.scatter_reduce_.two is not the unique overload of aten.scatter_reduce_r.  r   r   c                    j                         }t        |      }t        |       }t        j                   |       |dk(  rdn|   d      |<   |S )Nr   r/   F)r  )r   r   ri   rR   r  )r%  r"  r@  indirect_idxr   r  r  s       re   r  z'scatter_reduce_.<locals>.output_indexer  sU    5zCy11DAIq5:
S rg   c                `    r |       S t        j                  j                               S r   rR   r  r   )r%  r  r'  rU  s    re   r   zscatter_reduce_.<locals>.fn  s*    c?" <<T^^%566rg   c                    | dk(  ry| J y )Nr  r  r   r4  s    re   backend_reduce_strz+scatter_reduce_.<locals>.backend_reduce_str  s    U? >!>rg   c                L    t        j                  dj                               S r  rI  )r;  r  s    re   r?  z!scatter_reduce_.<locals>.<lambda>  s    3<<4>>3C#D rg   r  r  )r   r  r8  r   rk   r!   r  r3  r  rC   r  r   r   r  r  rK  r  r`  rn   r2   r  r  r  rS   r]   rS  r   rT  )r  r   r;  r'  r  r/  r9  r@  r  r   rK  rv   zero_outr/  r  r  rU  s   `` `           @@re   r8  r8    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  )
rR   r  r{   r  r   r  r_  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  )r%  r   r0  rs   r   r   i_sizes
inv_scalesr  rR  r2  s         re   r   zupsample_nearestnd.<locals>.fnJ  s`    H!HWaW#aW:UVVJAq$8Aq$'VW
 	
Vs   Ar^  )realize_hintr`  r   rS   r]   r   r  r   r  rj   r@   r   rn   r   )r   r  scales_xr  rQ  batchrs   o_sizesor  r   rT  rU  rR  r2  s      ``      @@@@re   upsample_nearestndr[  *  s    NN}}HjjlA23GJJL1"E6=>qww))!,>Gx=AG$'$9:DAq!a%:Jh' (5%KJqM(	;
 
 ||~kkm!!!	 9 ?
 ;s   0D'D$c                "    t        | ||fd      S )Nr/   r  r[  r   r  r  s      re   upsample_nearest1dr`  Y  s    avi1==rg   c                $    t        | ||fdd      S )Nr/   Tr  rQ  r^  r_  s      re   _upsample_nearest_exact1drc  ^  s    avi1DIIrg   c                $    t        | |||fd      S )Nr-   r]  r^  r   r  scales_hscales_ws       re   upsample_nearest2drh  c  s     ax.BaHHrg   c                &    t        | |||fdd      S )Nr-   Trb  r^  re  s       re   _upsample_nearest_exact2drj  j  s     ax.BatTTrg   c                &    t        | ||||fd      S )Nr   r]  r^  r   r  scales_drf  rg  s        re   upsample_nearest3drn  q  s     ax8.LPQRRrg   c                (    t        | ||||fdd      S )Nr   Trb  r^  rl  s        re   _upsample_nearest_exact3drp  |  s#     	;8X6!4 rg   c                ,     t         fd|D              S )Nc              3  J   K   | ]  }t        j                  |        y wr   rh  )r   r   r   s     re   r   z$_create_constants.<locals>.<genexpr>  s     6Aa'6s    #)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   )r%  r   r  r  r2  s     re   r  zrev.<locals>.loader  sW    3i3x3u:%%% 	3Cc
Q#c(2CH	3 }rg   r^  )r`  r   r@   r   rn   r   )r   r  r  r  r2  s    ` @@re   revru    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)rS   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   r2   rC   rm   r   r  r0   can_inplace_pad_graph_inputInputBufferr   freeze_layoutrp  rr  r   rS   r]   buffer_to_padded_sizerZ  rW  rj  fill_r   )r   paddingr  rz  r  rm  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                   t        |      dz  dk(  sJ t        d |D              rt        |       S t        j                  rt        | |      }|r|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   r  s     re   r   z"constant_pad_nd.<locals>.<genexpr>  s     
#a16
#r  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;  r2  s   re   r?  z/constant_pad_nd.<locals>.mask.<locals>.<lambda>(  s     rg   )	r  ro   range_mask_lowrange_mask_highr  r  rR   r  r  )r;  r  r%  r  r  rO  boundsr  
mask_sizesr  r2  s   `     re   r  zconstant_pad_nd.<locals>.mask   s    (+E!"Ivz(J 	:$C#tfaxN323qyOC89		:
 $/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  r%  r  _highbounds_precompr  r  s        re   	offset_fnz"constant_pad_nd.<locals>.offset_fn*  si    rO	!$U12Y!? 	(C#uS3Y'	(9~U+++Irg   r^  )r   r  r  r0   r|  r  r   ri   r)  r  ro   rS   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  r2  s     `         @@@@@@re   r  r    s   L1"""

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


cTUA}}H$A39I
 
* Krg   dilationc                  |dgt        |      z  }t        | d||   z  z   ||   ||   dz
  z  z
  ||   dz
  z   ||         }|rt        | d||   z  z   ||   ||   dz
  z  z
  d||   dz
  z  z   ||         }t        j                  j                  j                  |dz
  ||   z  | z
  ||   z
        dk\  r>|dz  }t        j                  j                  j                  d|||   z  | z
  ||   z
         t        j                  j                  j                  ||z
        dk(  r0t        j                  j                  j                  ||       d}||fS |}||fS )Nr/   r-   r   F)r   r)   rS   r]   r   	size_hintr=  r-  )	r   rs   kernel_sizerr  r  	ceil_moder  x_outx_alts	            re   pooling_sizer  o  s   3W%	A
NXa[KNQ,>??6!9q=Qq	E
 '!*nqk[^a/01 6!9q=!" 1I
 77%%uqyF1I&=&AGAJ&NOSTTQJEGG&&q%&)*;a*?'!**LM77%%eem49GG))%7I ) E)rg   c               l    t        | |      } t        j                  t        j                  |       }|dkD  S )N   )rM   r  r  r  r_  )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  )rM   rk   rC   r   r   r  )r   r  rr  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r,  r   r  s     re   r   z)_max_pool_with_offsets.<locals>.<genexpr>  s     ,EqQU,Er  r  c                    | d   }|  d  }t              D cg c]  }||   	|   z  ||   |   z  z   |   z
  ! }} 
g ||      S c c}w r   r  )r%  reduction_idxr  bhrs   r  r  r  r  rr  r2  s         re   fn_innerz(_max_pool_with_offsets.<locals>.fn_inner  s    Wuf%\ 5\
 UVAY=#3hqk#ABWQZO
 
 &2''	
s   $Ar'  r@  
input_noderv   	dst_dtyper  r]  r\  reduction_rangesargmax)rV  r"  r  r  r  r   r{   r   r  r6  r  r&  ri   r  r  r`  rA   r   rn   r  rk   rm   r  )r   r  rr  r  r  r  r  rX  dhwr  dhw_outr   	min_valuer4  r  r{  offsetsr2  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  r0   r   r  r   r{   r  )
r   r  rr  r  r  r  r  r  r{  r  s
             re   !_low_memory_max_pool_with_offsetsr    s     E 1@	1-K(A 
"	5 
50
 x44
5 
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   )rR   r  r1   _flattened_index_to_ndr  _flatten_indexr{   r  )r%  rW  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   rJ  r  r  r  r_  r@   r   rn   r{   r  r   )	r  r  r  r  r  r  r  r  r  s	    ```  @@@re   _pool_offsets_to_indicesr    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  )r%  r  r  rs   r  r  r  rr  s       re   r  zD_low_memory_max_pool_offsets_to_indices.<locals>.increments_to_indexF  s_    %\ 5\
 UVAY=#3hqk#ABWQZO
 	
 
s   $<)r   r  )r  r  r  rr  r  r  r  r  s      ``` @re   '_low_memory_max_pool_offsets_to_indicesr  =  s-     E
 $j*= rg   c           	         t        | |||||      \  }}}}}t        | ||||||      \  }}	t        |	|| j                  | d  |||      }
||
fS )Nr  )r  r  r  r"  )r   r  rr  r  r  r  r  r  rq   r  r  s              re   _max_pool_with_indicesr  R  s}     1@	;1-K(A *	;9ELC 6	G <rg   c           	     &    t        | |||||d      S Nr-   r  r  r   r  rr  r  r  r  s         re   max_pool2d_with_indicesr  p        "	;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-   r-  )rv   r   r   r  c              3  &   K   | ]	  }|d k7    ywr,  r   r  s     re   r   z3max_pool2d_with_indices_backward.<locals>.<genexpr>  s     
$a16
$r  c              3     K   | ]=  }t        t        |d          t        d t        |d    z
  d                z
  d       ? ywr   r/   Nr'  r)   r   r  r  rr  s     re   r   z3max_pool2d_with_indices_backward.<locals>.<genexpr>  M       	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   Nr  r   wr  rr  s     re   r   z3max_pool2d_with_indices_backward.<locals>.<genexpr>  r  r  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  rM  r  )rR   r  r{   r  r)   r  r  r  r  r   r  r  rT  r  r  r  r  )r%  r  r  r  
index_testphstartpwstartphendpwendgradientph_pw_phpw
grad_indexindex_actual	grad_partr   r  grad_loaderh_window_sizer  indices_sizer  r  pooled_heightpooled_widthrr  w_window_sizewidths                      re   r   z,max_pool2d_with_indices_backward.<locals>.fn  sA   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' #	WC]+ "W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"W#	WH ###rg   r^  )rk   rC   r   r   rV  r9  rm   r@   rn   r2   r  r  r   decide_layoutrX  r  )fallback_max_pool2d_with_indices_backwardr`  ri   r'  r  r   r  r  )grad_outputr   r  rr  r  r  r  r  	gO_striderm   rv   x_bufferx_strideis_channels_last_batch_heightr  r4  r  r   rq   r  r	  r  r
  r  r  r  r  s     ```                @@@@@@@@re    max_pool2d_with_indices_backwardr    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_indexr2  s   re   r?  z3pad_adaptive_loader.<locals>.load.<locals>.<lambda>:  s&    HNvN}r'9N=2;MNO rg   )rR   r  r  r  r{   r  r  )r  
incrementsstart_indicesend_indicesh_end_indexw_end_indexr  r  r  r  r  pad_valr2  s   `      @@@@re   rR  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   r$  rR  r2  s    ` @re   pad_adaptive_loaderr&  %  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_outr  r"  r  r#  s
             re    compute_indices_adaptive_poolingr.  A  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  )r%  r  r  r  bwr  r"  r  r#  r{  r  r  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>.fnZ  s    R(,$R((,$R(''l1o(>lSTo@VW 
	1FBR.k*	C ~#C0
	1 rg   r.  )r^  r_  r5  in_sizes	out_sizesr6  r*  r+  r,  r-  r   r3  r4  r7  r8  s     `  `     @@@@re   _adaptive_pooling_fnr<  K  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   r1  r  rR   r  r{   r  r  gt)r%  r  r  r  r2  r  r"  r  r#  maxvalmaxindexr  r  r  r;  r3  r4  r5  r6  r7  r+  r8  s                  re   r   z)_adaptive_pooling_fn_with_idx.<locals>.fn  s   R(,$R((,$R(''l1o(>lSTo@VW 	1FBR.k*	C NN#t+m;b@%++E  99SVVC%8%J~#C0)	1, rg   r9  )r^  r_  r5  r:  r;  r6  r*  r,  r-  r   r3  r4  r7  r+  r8  s     `  `    @@@@@re   _adaptive_pooling_fn_with_idxrB  t  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   r*  r/   r  c                     t        | |z  |      S r   r)   r;  r(  r)  s      re   r^  z)_adaptive_avg_pool2d.<locals>.start_index      733rg   c                2    t        | dz   |z  |z   dz
  |      S r  rE  rF  s      re   r_  z'_adaptive_avg_pool2d.<locals>.end_index  "    g-7!;WEErg   r^  r_  r5  r:  r;  r6  c                ^    t        j                   | t                     |             S r   )rR   truedivr&  )r%  fn_sumones_loaderr   s    re   r   z _adaptive_avg_pool2d.<locals>.fn  s-    {{3+A./[1I
 	
rg   r^  )r   r{   r  r  rk   rC   r   rV  r   rS   r]   r   r  r  r3  rn   r)   
avg_pool2drF   ri   fallback_adaptive_avg_pool2dr<  rR   r   r&  	ones_liker@   r   )r   r  rX  r*  r+  r,  r-  o_sizer  h_kernel_maxw_kernel_maxr4  r   r  r^  r_  r   rvrM  rN  s   `                 @@re   _adaptive_avg_pool2drV    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   r*  r/   r  c                     t        | |z  |      S r   rE  rF  s      re   r^  z(adaptive_max_pool2d.<locals>.start_index  rG  rg   c                2    t        | dz   |z  |z   dz
  |      S r  rE  rF  s      re   r_  z&adaptive_max_pool2d.<locals>.end_index!  rI  rg   rJ  c           	     <     | t        t        d                  S Nr  r&  r6  )r%  inner_func_max_valr   s    re   inner_fn_max_valz-adaptive_max_pool2d.<locals>.inner_fn_max_val6      !#':1eFm'LMMrg   c           	     <     | t        t        d                  S r[  r\  )r%  inner_func_max_idxr   s    re   inner_fn_max_idxz-adaptive_max_pool2d.<locals>.inner_fn_max_idx9  r_  rg   r^  )r   r{   r  r  rk   rC   r   rV  r   rS   r]   r   r  r3  rn   
ValueErrorrF   ri   fallback_adaptive_max_pool2dr<  rR   r  rB  r@   r   )r   r  rX  r*  r+  r,  r-  rR  rS  rT  r4  r   r  r^  r_  r^  rb  rU  rira  r]  s   `                  @@re   adaptive_max_pool2drf    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   rR   r  r   r{   r  rL  r   float64r  rT  r  r  r  r   rJ  )r  rs   samples_shapesamplei_exprdiffout_sz_exprrH  seq_ir  r   in_sz	kernel_szndimsout_szsamplessamples_loaders             re   rR  z)_fractional_pooling_offsets.<locals>.loadQ  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   r%  )rt  rp  rs  rq  r   rr  rR  ru  s   `````` @re   _fractional_pooling_offsetsrv  K  s@    C[F#JE#I((*N!Y !YF Krg   c                "    t        | |||d      S r  _fractional_max_poolr   r  r  random_sampless       re   fractional_max_pool2dr|  w      ;^STUUrg   c                "    t        | |||d      S r  rx  rz  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  )rt  rp  rs  rq  rr  r   c                6    | d   } g | | |            S r   r   )r%  r  r  r  r  r2  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  )r%  r  r  bdhwr  dhw_index_fnr  s        re   r  z1_fractional_max_pool.<locals>.increments_to_index  s^    5&\Fvw<D u  QQ0=3CC  s   :r'  r  r  )rV  r"  r0   r   r  rv  r`  ri   r   rA   r   rn   r{   r  rk   rC   rm   r  r  )r   r  r  r{  r  rX  inp_dhwr  r  r4  r   r{  r  r  r  r  r2  s       `         @@@re   ry  ry    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{= =

= =s   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   rJ  rF  s      re   r^  z0upsample_nearest2d_backward.<locals>.start_index  s    uwg(>??rg   c                     | dz   ||      S r  r   )r;  r(  r)  r^  s      re   r_  z.upsample_nearest2d_backward.<locals>.end_index  s    EAI99rg   rJ  c                (     | t                    S r   )r&  )r%  rM  r   s    re   r   z'upsample_nearest2d_backward.<locals>.fn  s    c.q122rg   r^  )rV  r   rS   r]   r   r  rO  r)   rF   r<  rR   r   r@   r   rn   r   ri   )r   r  r  rf  rg  r  inp_hinp_wout_hout_wrS  rT  r_  r   rU  rM  r^  s   `              @@re   upsample_nearest2d_backwardr    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  rr  r  r  count_include_padr  s          re   rO  rO    )     		 	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  )
r%  r  r  r  r  rs   r   r  rr  r2  s
         re   r  z_avg_poolnd.<locals>.fn_inner]  sq    Usd#Z>CCjIbefQi"Q%''!*4II&2'' Js   A	r  c              3     K   | ]D  \  }}t         j                  j                  j                  t	        j
                  ||             F y wr   )rS   r]   r   statically_known_truer   Ne)r   r   r   s      re   r   z_avg_poolnd.<locals>.<genexpr>f  s;       Aq 	
..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   MinMaxrR   r  r{   r  ro   r  r  r_  )r%  r  divide_factorsrs   hstarthendfactorr  r   r  r  r  rr  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^  )#rM   rk   rC   r   r   rV  r  r  r  r  r  r`  ri   r   r   r   r?  r  r  r  r_  fallbacks_avg_poolndr0   r   
contextlibnullcontextrn   rA   r   r   rm   r  div_primr@   r   )r   r  rr  r  r  r  r  r   rX  rs   r,  
ceil_modeshad_paddingr4  r   output_dtyper  r  fallbackcontextrv   rU  divisorr{  r  divide_factorr  r2  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-   r-  r/   c              3     K   | ]=  }t        t        |d          t        d t        |d    z
  d                z
  d       ? ywr  r  r  s     re   r   z&avg_pool2d_backward.<locals>.<genexpr>  r  r  c              3     K   | ]=  }t        t        |d          t        dt        |d    z
  d                z
  d        ? ywr  r  r  s     re   r   z&avg_pool2d_backward.<locals>.<genexpr>  r  r  r  c           	     @   t        j                  d   t        j                        }t        j                  d   t        j                        }t        j                  d   t        j                        }t        j                  d   t        j                        }t        j                  d   t        j                        }t        j                  d   t        j                        }t        j                  t        j
                  | |      |      }t        j                  t        j
                  ||      |      }	t        j                  t        j                  ||      t        j                  t        j                  t        j                        |            }
t        j                  t        j                  |	|      t        j                  t        j                  t        j                        |            }t        j                  |t        j                  dt        j                              }t        j                  |	t        j                  dt        j                              }	t        j                  |
t        j                  t        j                              }
t        j                  |t        j                  t        j                              }t        j
                  t        j                  |
|      t        j                  ||	            }|S )z{
        This computes the scaling factor that we will divide an element
        by when `count_include_pad=False`
        r   r/   )
rR   r  r{   r  r  r_  r  r   r  r  )r  r  stride_hstride_wpad_hpad_wkernel_hkernel_wr  wstartr  wendr  heightr  r  rr  r  s                re   !compute_pool_size_without_paddingz>avg_pool2d_backward.<locals>.compute_pool_size_without_padding  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  )rR   r  r)   r{   r  r  r  r  r  r   rL  r  r  r  r  r  r  )r%  r  r  r  r  r  r  r  r   r  r  r  r  r  partr  r  r  r  r  r	  r  r  r  r  r  rr  r  s                   re   r   zavg_pool2d_backward.<locals>.fn  s2   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' *	RC]+ )R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)R*	RV ###rg   r^  )rk   rC   r   r   rV  r  r`  ri   r   r'  r  fallback_avg_pool2d_backwardr@   r   rn   )r  r   r  rr  r  r  r  r  r  _h_out
ceil_mode1_w_out
ceil_mode2r4  r   r  r   rU  r  r  r	  r  r  r  r  r  r  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r  )r'  )r   r  rs   r  rr  s     re   r   z0avg_pool3d_backward.<locals>.<genexpr>.<genexpr>  sF      
 VAYQ[^);q	(I!JJAN
s   58r-   N)r'  r  )r   rs   r  rr  s    @re   r   z&avg_pool3d_backward.<locals>.<genexpr>  s;      3
 	 	 
;q>A-.
 	
3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   rR   r  r{   r  )r   r   s     re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  s     'UQ(D'U   .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     Ms||Au{{;Mr  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$      (
-.CLLEKK((
r  c              3  x   K   | ]2  \  }}}t        j                  t        j                  ||      |       4 y wr   )rR   r  r_  )r   r  r   pads       re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  s3      "
1c GGCGGAqM3'"
s   8:c           
   3     K   | ]i  \  }}}}t        j                  t        j                  ||      t        j                  t        j                  |t        j
                        |             k y wr   )rR   r  r   r  r{   r  )r   r[  r   r   r  s        re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  sU      

 #q#s KKq!3773>>#u{{+KS#Q

s   A/A1c              3     K   | ]<  }t        j                  |t        j                  d t        j                               > ywr  rR   r  r  r{   r  )r   r[  s     re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  s2      "
 KKs||Au{{;<"
   AAc              3     K   | ]?  \  }}t        j                  |t        j                  |t        j                               A y wr   rR   r  r  r{   r  )r   r\  r   s      re   r   zQavg_pool3d_backward.<locals>.compute_pool_size_without_padding.<locals>.<genexpr>  s6      
S KKS^^C=>
   AA)r  rR   r_  r  )pdr  r  stride_dr  r  pad_dr  r  kernel_dr  r  dstartr  r  dendr  r  r  depthr  r  r  rr  r  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>  s     Avq#1s7Ar  c              3     K   | ]<  \  }}}t        j                  t        ||z
  |z   |      t        j                         > y wr   rR   r  r)   r{   r  )r   r  r   r   s       re   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>  s<      %
1a NN8AEAIq15;;?%
r  c              3     K   | ]8  \  }}t        j                  t        ||      d z   t        j                         : ywr,  r  )r   r  r   s      re   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>  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>  s2      %
 KKQ <=%
r  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>  s6      
 j KKcnnZEF
r  c              3     K   | ]?  \  }}t        j                  |t        j                  |t        j                               A y wr   )rR   r   r  r{   r  )r   r  p_s      re   r   z2avg_pool3d_backward.<locals>.fn.<locals>.<genexpr>  s6      "&FB R(EF"r  r   r/   r-   Fr  r  )r  r  rR   rL  r  r  r  r  r{   r  r  r  r  r  r   )#r%  r  r  r  r  pdstartr  r  pdendr  r  r   pd_r  r  r  r  r  r  r  r  r  r  d_window_sizer  r  r	  r  r  r  pooled_depthr  r  rr  r  s#                        re   r   zavg_pool3d_backward.<locals>.fn  s   AqAaAY)@A1a%
1ay+v>%
!'

Q1Iv.
ue
%
"GW5%
!'
$'u%m\'R%
ue ' 8	VC]+ 7V / 6V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6V7V8	Vr ###rg   r^  )rk   rC   r   r   rV  r  r`  r  ri   r   r  fallback_avg_pool3d_backwardr@   r   rn   )!r  r   r  rr  r  r  r  r  r  _d_outceil_mode_dr  ceil_mode_hr  ceil_mode_wr  r4  r   r  r   rU  r  r  r  r  r	  r  r  r  r  r  r  r  s!     ``` ``             @@@@@@@@@@@@re   avg_pool3d_backwardr  ^  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  rM  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_axisr  &  s    ::<D$vSY
4yA~T{//H>$1HH/	:D3t9 M7Q;GCIs4y14GDG'c$i'CINtAw!|LLM 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  rR   r  r{   r  )r;  reduction_indexrs   r  r%  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]  sL   ?#s;'7777u:T***'/0!U1X0E05zS]***Fc%j3+??@	!% #k?"C
 	!HC !IcN	! Y' (/>?!ell1o?F?  J1c&k* G'-*::VAYF
G 3>>*ekkBCC+ 1 @s   D4D9)rv   r  r  r]  r\  r  )r   r   r'   r   r  r  r   ro   r;   rk   rm   r?   r2   r   r   BufferrV  is_transposedis_contiguousri   r   r*  r+  r`  r  rn   r   )r   r  r  r   rC  r@  
kept_sizesrs   rm  r  r4  r  r  r  r  r  r   s     `        @@@@@@re   _make_reduction_innerr  8  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                   t        | |||      }t        j                  d| d|}t        |j                  j                  t              r|j                          |S )N)r  r  r   rC  r@  )r@  r  r   )r  rA   r   rk   rm   r  )r   r  r  r   r   r{  rC  r@  s         re   rj  zmake_reduction.<locals>.inner  se    &"7)
 !!XAXQWXKK
 NNrg   r   r   )r@  rC  rj  s   `` re   make_reductionr    s    T " Lrg   c                   |t        | |      } t        | |      }t        | j                         | j	                         f| j                         f| j                         |      S )N)rv   dtypes	inner_fnsr   r  )r   rK  r  rn   r   r`  r   )r   r  r   s      re   _make_scan_innerr    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>       0a$q'0r   r:  )r   r   r  r   r{   rc  rb  r6  sum_rP   r2   r9   rn   r8   r   ri   div)r   r  keepdimr   r  
sum_resultdenomr   s          @re   rC  rC    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)r  c              3  (   K   | ]	  }|     y wr   r   r  s     re   r   z var_mean_sum_.<locals>.<genexpr>  r  r   r   r:  )r   r  rC  r  squarer  r  rP   r   r  r2   r9   r   rn   r8   r   ri   r  r  )r   r  
correctionr  return_meanx_meandiffsr  r  x_varr   s             @re   var_mean_sum_r%    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                    t        | |      }t        | ||d d       }|d   }t        |d         }t        |t        j
                        xr, t        |      t        j                  k  xr t        |      dk7  S )Nr  r  r   rC  r\  r  r/   )	r  r  rP   rk   r   ro  r   r0   r  )r   r  r  r   r\  reduction_numels         re   use_two_step_variancer)    s|    #At,D"	wd$F HF#F+=$>?O?EMM2 	' 6#E#EE	'&!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)r  r@  r   c              3  (   K   | ]	  }|     y wr   r   r  s     re   r   z$var_mean_welford_.<locals>.<genexpr>  s     1q471r   c                    t        | t        j                        rD| j                  s8t	        j
                  t	        j                  | t        j                        |      S t	        j                  | |      S r   )
rk   r   r   	is_numberrR   r   r  r{   r  r  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  )rR   r  r  )rm   cNzeror   r   r/  rnumels       re   rR  z#var_mean_welford_.<locals>.scale_fn  sE    &z59&vu5||Au%ckk$A...rg   r   )r  r  r2   WelfordReductionr   r   r  r   r  rP   rk  )r   r  r   r  r!  r   r  rC  m2r  rR  r  r   r/  r4  r   s     `         @@@@re   var_mean_welford_r7    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	        | ||      rt        di |n
t        di |}t        fd|D              }|s|d   S |S )NFr  )r   r  r   r  r!  )r  r  c              3  :   K   | ]  }t        |d         yw)Fr  N)r   )r   r   r  s     re   r   z#var_mean_helper_.<locals>.<genexpr>#  s     F!8Ayu55Fs   r   r   )r   r   r   r  r)  r%  r7  r   )	r   r  r   r  r!  r   r   r1  r  s	           @re   var_mean_helper_r:    s    I))4MM.A
F !w? 	(( 
 FvFFF'6!93V3rg   )r   r  c               "    t        | |||d      S )NFr  r   r  r!  r:  r   r  r   r  s       re   var_r?  '  s    	W% rg   c               "    t        | |||d      S )NTr<  r=  r>  s       re   var_meanrA  .  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_recursiverR   r  r  r_  )r   r2  r   r{  s       re   rC  rC  5  s    1uS^^A.E::Av||Au%%Av1a1fe,FWWVV$F	A!|#Mrg   c                .    t        j                  | |      S r   )rR   r  r   r0  s     re   
pow_nativerF  D  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 )
NrP  r/   c              3  r   K   | ]/  }t        |t        j                        s|j                          1 y wr   )rk   r2   rC   r   r  s     re   r   zpow.<locals>.<genexpr>\  s#     N1*Q2MNs   77i    r   c                F    t         |       j                               S r   )rC  r   )r%  r   r0  r  s    re   r   zpow.<locals>.fnf  s     a??rg   r^  r-   )rk   r6  r   r  r   sqrtr  r@  r    r`  r@   r   rn   r   r   r!   r  r   exp2fallback_pow_scalarfallback_pow_tensor_scalarfallback_pow_tensor_tensorrF  )r   r0  r   is_integer_powembed_exponentr   r  s   ``    @re   r  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   rC   rm   r2   r   r@   r   rn   r   r`  r   r6   r<   is_input_bufferis_module_buffer	NopKernelr  r  realize_into)changedr  rT  changed_datarc   s        re   r&  r&    s*   '9%||#y!hhc2==)%%'##%__&##%	
 $: 6777ii#r}}---,.$$&((*l''6 	HH!!..\ /  Nrg   c                .    t        | t        | |            S r   )r&  r  )r   r  s     re   r  r    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   r&  )r  r'  r  s      re   r)  r)    sR    
cz

C)
*C
3
(C
clln
%CS#rg   c                .    t        j                  | |      S r   )rR   floordivrE  s     re   r_  r_        <<1rg   c                .    t        j                  | |      S r   )rR   truncdivrE  s     re   rb  rb    r`  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 )Nr  z5floordiv operands can not be boolean at the same timer  z5truncdiv operands can not be boolean at the same time)r   r   r_  r  r  rb  r  )r   r0  rounding_modeboth_integerboth_booleans        re   div_moderg    s    "1%</!*<L"1%</!*<L X!XX!-x1~C5Q3CC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  r_  r  rk  )r   r0  	both_boolr   s       re   r_  r_    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   r2   r<   get_constant_valuerm   r6   rn  r   Loopsr{   	_inductorops_handlerExtractConstantsHandlerrn   rS   set_ops_handlerr   objectr  r]  inner_fn_argsvirtualizedOpsValuer   )r   rX  rq   s      re   rm  rm    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"     O1oa(>OA,>>Os   "$r   r   infr  c                 &    t        j                  |  S r   )rR   rL  r  s    re   r   zdiv_prim.<locals>.fn  s    {{D!!rg   )r  rb  rm  rn   r   r   mathcopysignr6  r_  rk  )r   r0  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|  )rE  r   INT_TO_FLOATr  rE  s     re   r  r    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   )rR   modrE  s     re   r   zfmod.<locals>.fn&  s    771a= rg   c                .    t        j                  | |      S r   )rR   fmodrE  s     re   r   zfmod.<locals>.fn+  s    88Aq>!rg   )r   r   rk  )r   r0  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  r  r   r  r  r   r   s        re   r  r  1  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   rM  Tr  c                @    | \  }|\  }t        j                  ||      fS r   )rR   r   a_tupleb_tupler   r0  s       re   
combine_fnzcumsum.<locals>.combine_fnO  #    1rg   r  r   r  r   r   r   )r    r   r   r{   r  r   r   r   r  r2   Scanr   fallback_cumsumr   r  r   r  r   r{  s         re   cumsumr  C  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   )rR   r_  r  s       re   r  zcumprod.<locals>.combine_fng  r  rg   r  r  r  r   )r    r   r   r{   r  r   r   r   r  r2   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   )rR   r  r  r  r  log1pexp)r  r  r   r0  min_vmax_vr  s          re   log_add_exp_helperz(logcumsumexp.<locals>.log_add_exp_helperu  s{    Aq!Aq!CIIe$4#45		$		#''%%-*@ AE I1MOOrg   r   r  r  r  r  r   )	r   r   r   r  r  r2   r  r   fallback_logcumsumexp)r   r   r  r   r   r{  s         re   logcumsumexpr  s  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  r  c                R    t        j                  |    t        j                        S r   rR   r  r{   r  r%  r  s    re   r?  zcummax.<locals>.<lambda>      CNN3t9ekk: rg   r  r  r  r   )r   r   r  r  r{   r  r   r2   get_reduction_combine_fnr  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   r  Fr  r  r  c                R    t        j                  |    t        j                        S r   r  r  s    re   r?  zcummin.<locals>.<lambda>  r  rg   r  r  r  r   )r   r   r  r  r{   r  r   r2   r  r  r`  r  r   fallback_cumminr  s    `     re   cumminr    r  rg   c                   t        | j                               st        | j                               r|t        j                  }t        d|      } || |||      S )Nr7  r  r   r  r  s        re   r7  r7    sK     	'+;AKKM+J
-	e	<Baxu--rg   c                ^    t        | t        j                        }  t        d      | ||      S )Nr  r  r  )r   r{   r   r  r   r   r  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    :    
g6!#8
 	

 qtg66rg   c                Z    |t        | ||      t        | ||      fS t        | d |      S r  )reduce_aminreduce_argminr  s      re   
reduce_minr    r  rg   xor_sumr'  r&  r  r  r  
logical_or)rg  r[  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                  j                  |t        j                  t        j                        j                        st        | |||      S t!        |ddt        j                  |d      }dgt        |      z  }t        |      r|||<   t#        ||      }t%        ||      }t&        j(                  j+                  || j,                  |j,                  f| j/                         |j/                         f||||      \  }	}|	t        | |||      S |J |	t1        |t        j                        fS )NFr   r/   r  )r[  r]  r   rv   rP  )rv   r  r  r   r  r  r  )r   rn   r   r   r  r  r{   r  rS   r]   r   statically_known_ltr  int16r'  sort_fallbackrQ  r  r.  r2   Sortr   r   r`  r   )
r   r  r   r  r"  rv   rA  r  
view_shaper   s
             re   sort_stabler    s   ~JJLE\\^F
3u:s
+C
5zQQxq&%++u=== ZuSzQH77//%++ekk:R:V:VWQv3:NNVSX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        | |t        j                  |      S )Nr   r   ri  )r  r   r  )r`   r   ri  s      re   register_pointwise_numericr    s!    
;HH'	 rg   c                b    t        | j                         t        | t        j                        S r|  )rO   r  r  r   r  r  s    re    register_pointwise_numeric_ldf64r  %  s'    '4
;HH rg   r  )r   c                 	 t         j                  st        S t        | ||t        j
                        | j                         |j                         |j                         j                  xr t        j                  j                   		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)

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

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

    When emulate_precision_casts is False, we return NotImplemented to use the
    decomposition instead.
    r  c                    |       } |       } 	|       }dk(  r
rt        j                  |||      S 
rt        j                  ||      }nt        j                  ||      }t	        t
        j                        rt        j                        }nt        j                        }
rt        j                  |||      S t        j                  |t        j                  ||            S r  )
rR   fmamul_rnr_  rk   r   r   r  r  r   )r%  self_valt1_valt2_valt1_times_t2
value_exprr   r  	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^  )r0   rX  NotImplementedr   r   r?  r`  r  r{   versionhipr@   r   rn   r   )
r  tensor1tensor2r   r]  r   r  r  r  r  s
      ` @@@@@re   r  r  8  s    & ));CC	E ""$K##%I##%I %%?emm.?.?*?GG G<  }}	 rg   c                d   t         j                  st        S t        t        j
                  j                  j                        dk(  xs6 t        j
                  j                  j                  t        v xs
 t               }t        t        | ||            }fd}t        |t        |       ||      S )z
    Foreach version of addcmul with scalar value parameter.
    Uses foreach_group_loop for consistent grouping behavior.

    When emulate_precision_casts is False, we return NotImplemented to use the
    decomposition instead.
    r   c                    t        | diS )Nr   )r  )rt   r   s    re   rp  z)_foreach_addcmul_scalar.<locals>.apply_fn  s    *E**rg   )r0   rX  r  r   rS   r]   r^   r_   ra   rq  rf   rx   r  rr  )r  r  r  r   rs  rw  rp  s      `   re   _foreach_addcmul_scalarr    s     )) 	AGG  &&'1, 	,77&&*==	,)+   D'7 ;<F+ fc$i?KKrg   logical_not)r[  )r   r  rC  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   )tritonr  )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lln 	IFV$Bd..0DR0HHH	I $**,@,DDDs   BB!r  c                d    | t         |<   t        j                  |        fd}t        | |       y )Nc                      | i |}g }t        | d   |      D ]"  \  }}|j                  t        ||d             $ |S )Nr   TrS  )r  ro   r&  )rt   r   resultsmut_resultsr   r{  outplace_ops         re   r   z$register_foreach_inplace.<locals>.fnF  sV    t.v.tAw0 	JKCyf4HI	J rg   )rZ   rq  r   r  )aten_opoutplace_aten_opr  r   s     ` re   register_foreach_inplacer  B  s.    07,-G$ w+rg   c                2    t        | d       fd       }|S )Nr  c                 l     | i |}t        || d   j                               }t        | d   |      S r  )r   r   r&  )rt   r   r{  r  s      re   r   zregister_inplace.<locals>.fni  s<    d-f-&$q'"3"3"56a&))rg   )r%  )r  r  r   s    ` re   register_inplacer  h  s#    wD9* :*
 Irg   c                     y r   r   r%  s      re   sym_constrain_ranger        rg   c                    t         j                  j                  j                  d   }t	        |t
        j                        r|j                  j                  S t        |      S r%  
rS   r]   r^   rd  rk   r{   r1  rc   r2  r   r   r   r  s      re   ra  ra    A    
''


#
#E
*C#u||$xx}}3xrg   c                    t         j                  j                  j                  d   }t	        |t
        j                        r|j                  j                  S t        |      S r%  r  r  s      re   
sym_strider    r  rg   c                "    | j                         S r   )r  r   s    re   	sym_numelr	    s    ;;=rg   c                &    t        j                  |  S r   )r   Addr  s    re   sym_sumr    s    99drg   c                    t        d      )NzHelpful for debuggingr   )r  rt   r   s      re   foobarr    s    
5
66rg   c                8    | j                          t        |       S r   )r  r  r   s    re   _realizer    s    IIK8Org   c                R    | j                          t        j                  | |       | S r   )r  r2   ResizeStorageBytes)variabler4  s     re   resize_storage_bytes_r    s#    (H-Org   c                    | j                          |j                          t        j                  t        j                  | |            S r   )r  rC   r   r2   SetSourceTensorKernel)r  source_tensors     re   set__source_tensorr    s5    LLNB44T=IJJrg   r)  c                    | |u r| S t        || j                               }t        || j                               }t	        || j                               }t        | |      S r   r]  )r  r'  s     re   
fsdp_copy_r    sR    #:JS^^-.sCMMO,S#,,.)c""rg   c                 	
 t        | t              sJ t        |t        t        f      sJ |t        j
                  }|t        j                  k(  rt        d|       |t        j                  k(  rt        |      dk(  sJ |t        j                  k(  rt        |      dk(  sJ | j                         
| j                         }| j                         }t        | j                  t        j                         r| j                  j#                         | _        t	        j$                         rit        j&                  j(                  j*                  rEt-        |      rt/        d      n0t1        |      r t	        j2                  |      j4                  ndndt6        j8                  j:                  j=                  
d      rt?        |||      S tA        | 
gd	g      }|jC                         	t        jD                  jG                  ||      }t        jH                  ||||      jK                         	
fd
}tM        jN                  |||t        |            }|S )Nzunsupported memory format: r   r   nanTr  r   r*  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   r?  z*resize.<locals>.inner_fn.<locals>.<lambda>  s    ZL(A rg   )rR   r  r{   r  r  r  )	r%  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   rC   ri   r   r{   contiguous_formatpreserve_formatr  channels_lastr   channels_last_3dr  r   rA  rm   r2   r6   rn  r  r  deterministicfill_uninitialized_memoryr   r6  r    r  r'  rS   r]   r   r:  r  rZ  r`  r   stride_ordered_for_memory_formatrq  r  r@   r   )r   r   rF  r   rv   x_flat
out_strider]  rq   r   r#  r$  r%  s            @@@@re   resizer/    s   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		
 	
F $$&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_wrapr2  get_constant_argsr2   UserDefinedTritonKernelr   rk   rC   )	r3  constant_args_idxr4  r5  r   r2  constant_argsr  r  s	            re   triton_kernel_wrap_r<  $  sj     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   rK  r  s     re   r   zcond.<locals>.<genexpr>?  s#     
Ma:a 1Yq\1
MrL  z"control flow operator: torch.cond.r  r  )r  rS   r]   r^   rd  re  r  r2   Conditionalr   ri   maprC   )predtrue_fnfalse_fnoperandsr   r  r{  s          re   r   r   9  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   rK  r  s     re   r   zwhile_loop.<locals>.<genexpr>M  s+      * 	1f.)A,.*rL  z(control flow operator: torch.while_loop.r  r  )r0   r  r  rS   r]   r^   rd  re  r  r2   	WhileLoopr   rk   r   ri   r@  _maybe_wrap_as_tensor_box)cond_fnbody_fncarried_inputsadditional_inputsstack_outputr   r  r{  s           re   
while_looprN  I  s     !!c *"33* ' 9''..3377tLL;LE*;-8C,/)\\  .*;\F fh'''BLL::FCDDrg   )rM  c                    t        j                  j                  | g| }t        t	        t
        j                  |            S r   )r2   InvokeSubgraphr   ri   r@  rC   )subgraph_fn
identifierrD  r{  s       re   invoke_subgraphrS  b  s5    %%k=H=FI$$f-..rg   c                X   d}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        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

    Nrj  r1  zNo output node found in graph)rj   r]   nodesr`   rS   envfetch_args_kwargs_from_envr{   rl  Interpreterr1  r^   run_noder  )r  rt   r1  rs   rc   output_argsr   saved_current_nodes           re   process_subgraph_nodesr\  h  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$!:$ ~:;;M (:$s   AFF))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-   rj  r  N)rk   r:   r  ro   r  rS   r]   r^   rt   r   
operationsr  
find_nodesr\  ri   operation_nameadditional_buffer_depsr   )additional_depsrQ  rt   	dep_namesdeporiginal_args
arg_offsetoperation_lenr1  r`   dep_nameop_names               re   control_deps_op_loweringrk    sV    I )#v&() GG((--MJt9z!S%7777**+M{''--88M8JKsSWyXXX $K$<$<d4jIF/11 gg  0 B! 	BH''G&&&GG**7377A	BB 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_optionsrj  r1  )rS   r]   r^   rd  re  rj   r  rU  r`   rV  rW  r   r  r   r  codegen_low_precisionlow_precision_codegen_opsr   r  invoke_quant_opsr{   rl  rX  r1  rY  )
rQ  rl  rD  r1  rn  rs   rc   rt   r   r  s
             re   invoke_quant_tracerrr    sl   FGG((--11/4HM$$$[55;;AAB 7477m# (AGGKKWW 77==dCLD&__T6==?; E		 66GG5599!:N:N:PQ((,,Q-A-A-CDE XX))00$fMF ! 0 0 6AGGKK%7( 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 argumentsr*  c                d     g t        j                  |       t        j                  |       S r   )rP  rn  )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     7qQ[[]7r  r  c              3  <   K   | ]  }|j                           y wr   r%  r  s     re   r   z#associative_scan.<locals>.<genexpr>  s     <A<r  r  F)r  can_fallback_to_atenz/Unable to generate code for associative_scan opr   )r  rt  ru  r   r  r   r  r   rn   r  r   r2   r  r   )r  xsrL  rt  ru  r   subgraph_inputsrz  r   r{  ry  s             @re   associative_scanr    s     M
!a
 	
 R( 	akkmALLNCO  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   r?  zwith_effects.<locals>.<lambda>*  s    !))+ rg   c                d    t        | t        j                        rt        j                  |       S | S r   rR  r   s    re   rS  z"with_effects.<locals>.wrap_tensors-  rT  rg   zCNo operation nodes were generated when lowering effectful operator .c                      yr   r   r   rg   re   r?  zwith_effects.<locals>.<lambda>=  s    rg   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   r2   TorchBindObject	get_valuerC   rm   r   r  AttributeErrorr   )r   ru  s     re   convert_ir_to_valuez)with_effects.<locals>.convert_ir_to_valueH  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{   rR   higher_orderrS  torch._guardsr  r  try_gethop_dispatch_set_cache	get_cacherk   get_effectsr   r@  r*  rS   r]   r_  rX   rP  r  rC   rU  r2   rV  r   effectful_opsre  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effectsrh  r{  rS  prev_effect_buffernew_oprj  r  schema_argsschema_kwargsschemae	error_msgs                           re   with_effectsr  	  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	J "++222GGG
 qww!!-.12Q6 
MbTQRS6 WW2266{Cgg((8 	YF&2F#! //+,,W599:L:U:U:WX	Y  	
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                t   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                  t!        j"                  d
             t%        | |d      }	t'        t(        j*                     t-        | |	            }
t/        |
|d      }|	|fS )zn
    Lowering inductor_prims.prepare_softmax_online to compute max/sum in one pass if no split is needed.
    TNr'  r  online_softmax_reduce)r@  r(  r/   r-   )r  
num_outputreduction_hintz
            Online softmax is disabled on the fly since Inductor decides to
            split the reduction. Cut an issue to PyTorch if this is an
            important use case and you want to speed it up with online
            softmax.
            )r  r   )r  rS   r]   r   simplifyrP   r2   rA   
num_splitsstatically_known_geqr0   r  r=   r   r[  r\  textwrapdedentr  rX   r  r  r  r  )r   r   r   r  r4  hint	num_split
max_tensor
sum_tensorrD  r  xsums               re   prepare_softmax_onliner  w  s5   
 #	d$dF 01WW&&}5E'FGFll-- 
.OD) A~!''**??22 "8!>!> "
Qt"
?E"

J :%% 	OO		
 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{   cudais_availableget_device_capabilityr   rg   re   _is_sm100_or_laterr    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)r  r   r   r{   r  rc  rb  rc  r   r  r$  rR   inline_asm_elementwiseuint16rk  r  )r   r   r   r{  s       re   cvt_e8m0_rceil_loweringr    s     !T
 	
 MMOEU]]EMM5>>BBOPUwW
 	

 sEMM*			""7ll
B  ^B$FFEKK((rg   )rW  )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}   rX   re  r%  r  r  )r`   old_handlers     re   force_fallbackr    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Optional[Callable[..., Any]])r   ztorch._C.Tagr   z(Optional[Callable[..., tuple[Any, Any]]])r   r   r   r  r   None)r   zUnion[Collection[Union[torch._ops.OpOverload, torch._ops.OpOverloadPacket]], torch._ops.OpOverload, torch._ops.OpOverloadPacket]r   zOptional[list[Any]])r   z9Union[torch._ops.OpOverloadPacket, torch._ops.OpOverload]r   zCallable[..., tuple[Any, Any]]r   r  )r   zUnion[int, torch.dtype]r   torch.dtype)r   r	   r   z,TypeGuard[Union[TensorBox, sympy.Expr, int]])r   r	   r   z!TypeGuard[Union[TensorBox, bool]])rt   r	   r   r   r   r   r   r  )r`   z>Union[Any, torch._ops.OpOverloadPacket, torch._ops.OpOverload]r   r  r   r   )r   rC   rv   torch.devicer   rC   )rt   	list[Any]r   zdict[str, Any]r  r   r   )Optional[ELEMENTWISE_TYPE_PROMOTION_KIND]r  r   r   z tuple[list[Any], dict[str, Any]])r   r  r  r  r   r  )
r  r  r  r   r   r  r  r   r!  rW   )r   r  r   z.Callable[[Callable[_P, _T]], Callable[_P, _T]]rE  )NNNFN)F)r   rC   r   r  r  r   )r   rC   r   r  )r   rC   rv   r  )r   rC   r  rC   r   )r   rC   r  zSequence[sympy.Expr]r   rC   )r   r   l            r/   Tr  )ru  rC   r  rC   r  rC   r  r   r  r   r  r   r   r  r   rC   )ru  rC   r  rC   r  rC   r  r   r  r   r  r   r   r  r  Optional[torch.dtype]r   rC   )ru  rC   r  r6  r  r   r  r   r  r   r   r  r   rC   )ru  rC   r  r6  r  r   r  r   r  r   r   r  r  r  r   rC   )ru  rC   r  rC   r  rC   r  r   r  r   r   r  r   rC   )ru  rC   r  rC   r  rC   r  r   r  r   r   r  r  r  r   rC   )r   r   r/   )rW  r   r  r   r  r   r  )T)rw   ztorch.Tensor)rc   ztorch.fx.Node)NTF)rv   r  )r   	list[int]r  rC   r  r  rW  r   )
r  r   r  r   r   r  r  rC   rW  r   )r  rC   r   z.tuple[str, sympy.Expr, sympy.Expr, sympy.Expr])r  rC   r   ztuple[str, sympy.Expr])r  rC   r  rC   r  r   r  r   r  Optional[str]r  Optional[TensorBox]r   rC   )ru  rC   r  rC   r  r   r  r   )r   r   r;  r   )r   NNr/   )NNN)rM  FF)r   r   )r}  r  r   r   r  r  r/  r   )r   r   r  r  )r   r   r/  r   )r-   F)rW  ztuple[Optional[float], ...]r  r   rQ  r   )r  Optional[float])rf  r  rg  r  )rm  r  rf  r  rg  r  )r   rC   r  zSequence[int]r  r6  r   r  )rs   
sympy.Exprr  zUnion[sympy.Expr, int])rs   r  r  r  )rs   r  r  r  r  r  )Nr  N)
r  rC   r  "Sequence[Union[int, torch.SymInt]]r  r  r  zxCallable[[Sequence[Union[int, torch.SymInt]], Sequence[Union[int, torch.SymInt]]], torch._inductor.virtualized.OpsValue]r   rC   )Nr   r/   F)r  )NNNN)r   r   FTN)r@  rT   r   )r   r  r   zOptional[ir.Constant])rM  F)r`   ztorch._ops.OpOverloadPacket)r/   )r   z3list[Union[ir.TensorBox, ir.ShapeAsConstantBuffer]])rQ  ir.SubgraphrR  r  )r  ztorch.fx.GraphModulert   r  )rQ  r  )r  r  rL  ztuple[torch.Tensor]r  (  
__future__r   r  r  r  r   loggingr{  r  r  r  r[  collectionsr   collections.abcr   r   r   r   typingr	   r
   r   r   r   r   r   typing_extensionsr   unittest.mockr   r   r{   $torch.ao.quantization.fx._decomposedtorch.fxtorch.utils._pytreer  _pytreerP  torch._dynamo.utilsr   (torch._higher_order_ops.associative_scanr   r7  r   "torch._library.fake_class_registryr   torch._library.utilsr   torch._prims_commonr   r   r   r   r   r   r   r   r   r    r!   torch.fx.experimental.sym_noder"   r#   rJ  r$   r%   r&   torch.utils._ordered_setr'   torch.utils._sympy.functionsr(   r)   r*   r+   r,   _dynamo.utilsr.    r0   r1   r2   r3   decompositionr4   r5   r6   r7   r8   r9   r:   r;   r<   r=   r>   r?   r@   rA   rB   rC   rD   rE   rF   rG   rH   rI   rJ   rK   rL   rM   rN   rO   rP   rQ   ru  rR   rS   rp  rT   rU   rV   FALLBACK_ALLOW_LIST	getLoggerr  r  rX   __annotations__rY   r|   r}   r  r  tr_c10dr  r   _higher_order_opsr  rb   rq  rZ   quantized_decomposedrf   rx   r   r~   r   r   r   rZ  r|  rO  r  bmmconvolutionconvolution_backwardr  r  r  rx  rh  rj  _int_mmr  r  r  r  r  rc  r  ri  	complex32	complex64r   rb  r   r   r   r   r   r   r   r   r  r  r"  r?  r%  r3  rE  rk  rx  rr  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  r  r  r  r  r.  r&  r(  r7  _unsafe_viewreshaper:  slicerj  ry  r  quantize_per_channelr  r  r   r  _functional_assert_asyncr  dequantize_per_channelr  quantize_per_tensorr  dequantize_per_tensorr  r  r  r  r  r"  r$  r(  r0  r5  r7  r9  rD  r!  rH  rK  rN  r  cacher]  rc  rh  rx  r  r  rngprimsr  r  r  	bernoullir  r  r  	lru_cacher  r  r  r  r  r  r  r  r  randintstreamsrecord_event
wait_eventforce_stride_orderr  r  r  r  r  lookup_seedr  randomr  r  r  r  r  r  r  	NO_OPMATHr  r  r   r  r  r   r   rC  _adaptive_avg_pool3dadaptive_max_pool3d*_scaled_dot_product_attention_math_for_mpsuniformexponential_pdist_forwardsoft_margin_loss_backward_fused_rms_normxpur  embedding_dense_backwardmtia_is_compilednative_layer_norm_cdist_forward_cdist_backward
_trilinearsegment_reduce_segment_reduce_backwardhistc	histogrambin_ct_histogramdd_bin_edges_histogramdd_from_bin_ctsaddbmm_addmm_activation_grouped_mm
_cudnn_rnn_cudnn_rnn_backward
miopen_rnnmiopen_rnn_backward_embedding_bag_embedding_bag_forward_only_embedding_bag_backward*_embedding_bag_per_sample_weights_backward_fused_moving_avg_obs_fq_helper*_fused_moving_avg_obs_fq_helper_functional max_pool3d_with_indices_backward_adaptive_avg_pool2d_backward_adaptive_avg_pool3d_backwardadaptive_max_pool2d_backwardadaptive_max_pool3d_backwardfractional_max_pool2d_backwardfractional_max_pool3d_backwardreplication_pad1d_backwardreplication_pad2d_backwardupsample_linear1d_backwardupsample_bicubic2d_backwardupsample_trilinear3d_backwardgrid_sampler_2d_backward_pdist_backwardr  r  kthvaluetopkr  median	nanmedianr  resize_
resize_as__linalg_detlinalg_householder_productlinalg_inv_exlinalg_ldl_factor_exlinalg_ldl_solve	linalg_lulinalg_lu_factor_exlinalg_lu_solvelinalg_matrix_exp	linalg_qr_linalg_slogdet_linalg_solve_exlinalg_solve_triangular_linalg_svd	lu_unpackormqr_linalg_check_errorslinalg_pinvatol_rtol_tensor_linalg_eightriangular_solvelinalg_cholesky_excholesky_inversecholesky_solvegeqrf_fft_r2cnonzerogcd_thnn_fused_lstm_cell_prims	rng_primsrun_and_save_rng_staterun_with_rng_stategraphsafe_run_with_rng_statemasked_scattermasked_scatter_backwardrt  angle_efficientzerotensor(_sparse_coo_tensor_with_dims_and_tensors	to_sparse
_to_sparser   r  r5  r3  #_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  rK  ru  rQ  rV  r^  r`  scalar_tensorrr  
LongTensorrt  rz  r|  r~  r  r  r  r3  r  r  r  rQ  
zeros_liker  r  r  r  r  r  r  r  r  r  r  r  r;  r  r  r  r  r  r  r  r  r  fallback__unsafe_masked_indexr#  ,fallback__unsafe_masked_index_put_accumulaterG  r+  r  r3  r-  r=  r;  rA  r8  r[  r`  rc  rn  rp  r  ru  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r&  r.  r<  rB  rV  rP  rf  rd  rv  r|  r  ry  r  r  
avg_pool1dr  r  r  r  r  r  r  r  r  rC  r%  r)  r7  r:  r  r?  rA  rC  rF  r  Tensor_TensorrO  r  rM  Tensor_ScalarrN  r&  r  r)  r_  rb  r  rg  r_  rm  r  true_divider  r  r  r  r  r  r  r  r  r  r  r  r  r7  r  r  r'  r  r&  r  r  rD  r  rE  r  r  r  r  r  r   r  r  r  r  rsqrtr  rL  expm1relurL  rK  r  r  r  r  r  cossinabsbitwise_andbitwise_left_shiftbitwise_not
bitwise_orbitwise_right_shiftbitwise_xorlgammaerfspecial_erfr  tantanhri  r  r  logical_xorr  r  	clamp_min	clamp_maxnegr  	remaindersignsignbit	_neg_viewler  r  r?  rT  necoshsinhacosacoshasinasinhatan2atanatanhr|  erfcerfinvhypotlog10log2	nextaftercodegen.commonr~  r  r  r   r`   r   ri  _foreach_addListforeach_add_listforeach_add_scalar_foreach_mulforeach_mul_listforeach_mul_scalar_foreach_sub_foreach_neg_foreach_abs_foreach_powScalarAndTensor_foreach_divforeach_div_listforeach_div_scalar_foreach_sqrt_foreach_rsqrt_foreach_maximum_foreach_minimum_foreach_clamp_min_foreach_clamp_max_foreach_reciprocal_foreach_sign_foreach_copyforeach_copyr  _foreach_add__foreach_mul__foreach_div__foreach_copy_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  ra  r   r  r	  r   methodfuncr  r  _inductor_testr  r  r{  r  set_source_Tensorr  fsdpr  r/  *torch._higher_order_ops.auto_functionalizer0  r<  r  r   rN  while_loop_stack_outputr$  rS  r\  .torch._inductor.fx_passes.control_dependenciesr]  rk  invoke_quantrr  r  r  r  r  comm_loweringr  r  r  r  cvt_e8m0_rceilr  rW  r  register_quantized_opsregister_woq_mm_opsr  register_onednn_fusion_opsr  register_jagged_opscontextmanagerr  r   rg   re   <module>r	
     s9   "        	   # D D P P P '    +  $ $ ( H U ? :    M 
 0  - 8 8 =    $      * T]t_ !   g!FH	C H     .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A3.3 
3   !!$$$$--&&, {{zz{{{{{{}}}}}}

."# "'CC8C C 	C( FSV	(O
OO O C	O
  O &Od"/A42!2 2 C	2
  2 L2n  	(// 4,#42n  $Vr$N@E 5**77TR) S)X 5--4H) I) @E 5( 499??=& >& ;@e K 5##>F ?F 
7?? $*Z $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;;1  1h 4$$$?499$74<<T:1 ; 8 @1 4<<T:> ;> 4::48uE 9u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*** * 	*
 * * **Z ..554 (,,,, , 	,
 , , %, ,,^ 488y: y:x 4==d;3E <3El 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%6P7Nt	, 599%%11tL#+ M#+L 4&&DAG BG 4??= > 4>>##>' ?' 4<<  ! QB B/
 )):):; *499+>+>? )$***<*<= +DJJ,@,@A  dll  eii,,44 5 eii**22 3 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CRS  D6 >))tDLM	(09FI E8 . 4$$++F
 "&^^
^ 	^
 ^ ^  ^ ^ G^B NN/N/X/X 777 	7
 77t $FT 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 dii  dii  dmm  dii  dii  dkk  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
 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&&DA  B F 4%%4@A AAH ELL$"4"456tDU 5 7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tn !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;DH : <:" 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3RVI-I@OI 4I 41199:RVU-U@OU ;U 4**223 !% $ $S S 	S
 S 4S 41199: !% $ $	 	 		
 	 ;	7 599$$% &*\\(\6;\\~ 4''TB4 C4n :>@ MQ < JN@:HV 5::PTU 5 V5D3 3	 > 	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 (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$L !!1!1uM 499##>!re $2 ?$2N 499$$$?H @H 	#4::.&txx0!$)),"4::.$))$
*4<<
8'		2	DKK	(t4 4<<40-. F 1FRL2 400779P Q ! *   *" !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+		

 5Mt5 
0 	 3+		


, .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)$*<*<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'') 8LFD1(01$78 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/8/ J/ 599))44$OE PE& 	II22)JT24
 599))99tT/ U/
 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1 T1hW
 >00dK") L")L     ! ! +  * * , '  ' ' )   ,  + + -   %  $ $ &  rg   