
    ^j,                      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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Zd dl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	m#Z# d dl$m%Z%m&Z&m'Z'm(Z(m)Z)m*Z*m+Z+m,Z,m-Z-m.Z.m/Z/ d dl0m1Z1m2Z2m3Z3 d d	lm4Z4 d dl5Z5d dl6Z6d dl7m8c m9Z: d d
l;m<Z< d dl=m>Z> d dl?m@Z@ d dlAmBZB d dlCmDZD d dl7mEZEmFZF d dlGmHZH ddgZId dlJmKZK d dlLmMZMmNZNmOZOmPZPmQZQ e,rhd dlmRZRmSZSmTZT d dlUmVZV d dl6mWZWmXZXmYZY d dlZm[Z[ d dl\m]Z] d dl^m_Z_ d dl`maZa ddlbmcZc ddldmeZe ddlfmgZg dd lhmiZi dd!ljmkZkmlZlmmZmmnZnmoZompZp dd"lqmrZr dd#lsmtZtmuZu g d$Zv e/d%      Zwej                  d_d&       Zyd d'lzm{Z{ d d(l|m}Z} d d)l~mZ d d*lmZ d d+lmZ d d,lmZ d d-lmZmZmZmZmZmZ d d.lmZmZ d d/lmZmZ dd0lmZ dd1lmZ ej                  d2k(  Z ej2                  e      Zd3Z e/d4      Zee5j>                  e5j>                  f   Ze6jB                  ez  e6j                  z  dz  Zerd5n ejH                  d6d7      Zd8d9d:e d;Zd<Zd<Zd<Zd=Z eDe6jV                  e6jX                  e6jZ                  e6j\                  e6j^                  e6j`                  e6jb                  e6jd                  e6jf                  e6jh                  e6jj                  e6jl                  e6jn                  e6jp                  e6jr                  g      Zd>ed?<   d@Zeedz
  z  d k(  redAk\  sJ dB       d`dCZdadDZ G dE dFe5j~                        Z ej                  dGH       G dI dJ             ZdbdcdKZ	 	 	 dd	 	 	 	 	 	 	 	 	 dedMZ	 	 	 	 	 	 	 	 	 	 dfdNZ	 	 	 dd	 	 	 	 	 	 	 	 	 dedOZej                  dgdP       ZǐdhdQZȐdidRZɐdjdSZ	 	 	 	 	 	 dkdTZ	 	 	 	 	 	 dldUZ̐dmdVZ͐dndWZdodXZ	 	 	 	 dpdYZϐdqdZZАdrd[Z	 	 	 	 dsd\ZҐdtd]Zd^ f	 	 	 	 	 dud_Z	 	 	 	 	 	 	 	 dvdaZՐdwdxdbZ	 	 dy	 	 	 	 	 	 	 	 	 dzdcZ	 	 	 	 	 d{	 	 	 	 	 	 	 	 	 	 	 	 	 d|ddZؐd}deZِd~dfZڐddgZېddhZܐddiZ e2dj      Z e/dkdGl      Zee'e%ef   ef   Z G dm dne+e(eef         ZddoZ	 	 	 	 ddpZ	 	 	 	 ddqZ	 	 	 	 ddrZ	 	 	 	 	 	 ddsZ	 	 	 	 	 	 ddtZ	 d	 	 	 	 	 dduZ	 	 	 	 	 	 ddvZddwZddxZddyZddzZdd{Zdd|Zdd}Zdd~ZddZddZ eg de6j                  j                  dgng z         Z	 	 	 	 ddZddZddZd dlZddZg Zd`ed<   ddZddZej                   dd       Zej                   	 	 	 d	 	 	 	 	 	 	 dd       ZeZeZeZdLdddZdLd	 	 	 	 	 	 	 ddZ ejD                  dA      dd       Z G d de*      Z	ej                   G d d             Z
 G d d      Z G d de      Z G d de      ZddZ G d de      Zej                   dd       Z G d d      Z G d de      Zej                  ddd       ZejD                  dd       ZejD                  dgd       ZdgdZddZ	 d	 	 	 	 	 	 	 ddZdddZ	 	 	 	 	 	 ddZddZddZddZddZdLdLdGd	 	 	 	 	 	 	 	 	 ddZddLd	 	 	 	 	 	 	 ddZ 	 d	 	 	 	 	 ddZ!dLd	 	 	 	 	 	 	 ddZ"dLd	 	 	 	 	 	 	 ddZ#	 	 	 	 	 	 	 	 ddZ$ ejD                  d      dgd       Z% ejD                  d      dgd       Z&d Z' ejD                  d      dgd       Z(	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddZ)ddZ*	 	 d	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddZ+ddZ,ee5j>                  z  Z-ded<   ej                  	 d	 	 	 	 	 	 	 	 	 dd       Z.ej                  ddÄ       Z/ej                  ddĄ       Z0ej                  ddń       Z1ej                  ddƄ       Z2ddǄZ3ddȄZ4ddɄZ5ddʄZ6dd˄Z7	 	 	 	 	 	 	 	 dd̄Z8	 	 	 	 d	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dd̈́Z9dgd΄Z: G dτ dЫ      Z;	 	 	 	 	 	 	 	 ddфZ<	 	 	 	 	 	 	 	 dd҄Z=ddӄZ>ddԄZ?ddՄZ@	 	 	 	 	 	 	 	 ddքZA	 	 	 	 	 	 	 	 ddׄZBej                   	 	 	 	 	 	 dd؄       ZC	 d	 	 	 	 	 ddلZDddڄZEddۄZFdd܄ZGdd݄ZHddބZIdd߄ZJej                   dd       ZKddZLej                  dd       ZMej                  dd       ZNej                  dd       ZOddZPddZQddZRddZSdgdZTdgdZUddZVddZW G d dej                        ZY	 	 	 	 	 	 	 	 	 	 ddZZddZ[	 	 	 	 ddZ\	 d	 	 	 	 	 ddZ]ddZ^	 d	 	 	 	 	 ddZ_ddZ`	 	 	 	 	 	 ddZa	 	 	 	 	 	 	 	 ddZbd f	 	 	 	 	 	 	 	 	 	 	 ddZcd f	 	 	 	 	 	 	 	 	 	 	 ddZdddZeddZfddZgej                   G d d              Zhej                   dڐd       ZidېdZjdܐdZkdgdZl ejD                  d      	 	 	 	 	 	 	 	 dݐd       Zm	 ddddd	 	 	 	 	 	 	 	 	 	 	 dސdZndߐdZodd	Zp	 	 	 	 	 	 	 	 	 	 	 	 	 	 dd
ZqddZrddZsddZtddZu	 	 	 	 	 	 	 	 ddZvddZw	 	 	 	 	 	 	 	 ddZxddZy	 d	 	 	 	 	 	 	 ddZz	 	 	 	 	 	 ddZ{ddZ|	 	 	 	 	 	 ddZ}dgdZ~dېdZddddddddZej                         D  ci c]  \  } }|| 
 c}} Z ej                  d       Zdd!Zdd"Zdd#Zdd$Zej                  dd%       Zej                   G d& d'             Zi Zd(ed)<   	 	 	 	 	 	 	 	 dd*Z eD       Zd+ed,<   dd-Zddd.Zdd/Zdd0Z e/d1      Z e/d2      Z G d3 d4eeef         Z e1dG5      ddGdHdd6       Zdd7Z G d8 d9ej                        Zej                  dd:       Zdgd;Zdd<Zdd=Zdd>Zdd?Zd_d@ZddAZdgdBZddCZdDZddEZddFZd dGZ	 	 d	 	 	 	 	 	 	 	 	 ddHZddIZddJZdgdKZddLZddMZ ej                  dGH       G dN dO             ZedPe%f   Zeeegef   Z G dQ dR      Z e       ZddSZddTZddUZd	dVZd
dWZ eDg dX      ZddYZe"ddZ       Zdd[Z	 	 	 	 	 	 	 	 	 	 	 	 	 	 dd\Z	 	 	 	 	 	 dd]Z	 d	 	 	 	 	 	 	 dd^Zyc c}} w (      )annotationsN)Callable
Collection	GeneratorIteratorMappingMutableMapping
MutableSet)datetime)	lru_cache)StringIO)AnycastConcatenateGenericLiteral
NamedTupleProtocolTYPE_CHECKING	TypeAlias	TypeGuardTypeVar)dataclass_transform	ParamSpecSelf)mock)datasheet_tops)DeviceProperties)_needs_inductor_compile)dtype_abbrs)
OrderedSet)tree_flattentree_map_only)has_triton_package!activation_quantization_aten_passinductor_autotune_lookup_table_sympy_subs)free_symbolsfree_unbacked_symbolsGuardOnDataDependentSymNodeIterateExprsShapeEnv)IterableSequence
ValuesView)Path)SymBoolSymFloatSymInt)ELEMENTWISE_TYPE_PROMOTION_KIND)GraphModule)Node)ScalingType   )WorkspaceArgPythonWrapperCodegen)DepGraphLowering)BufferExternKernelIRNodeLayout	OperationReinterpretViewCompiledFxGraph)BaseSchedulerNodeSchedulerBuffer)cudampsxpumtiaTc                     t         D  cg c]#  } t        t        |       j                         s"| % }} t	        |      dk  sJ t	        |      dk(  rd}|S |j                         }|S c c} w )Nr9   r   rJ   )	GPU_TYPESgetattrtorchis_availablelenpop)x
avail_gpusgpu_types      `/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_inductor/utils.pyget_gpu_typerZ   k   sg    &K'%*;*H*H*J!KJKz?aZA-vHO 4>>>3CHO Ls
   #A'A')get_interface_for_device)detect_fake_mode)
DeviceType)	EventList)GraphTransformObserver)	ShapeProp)CeilDivCleanDivFloorDivIdentityMaxModularIndexing)make_symbolSymT)bound_sympyValueRangesconfig)ceildivwin32!inductor_do_bench_using_profiling_TspvTORCHINDUCTOR_XPU_KERNEL_FORMATzebinz.cubinz.hsaco.)rJ   hiprL         zOrderedSet[torch.dtype]_TMA_SUPPORTED_DTYPES@      zmust be power of 2c                *    | t         z   dz
  t          z  S )z/Round up to the nearest multiple of ALIGN_BYTESr9   )ALIGN_BYTES)nbytess    rY   _alignr~      s    [ 1$44    c                   t        | t        j                  t        j                  t        f      r#t	        t        t        | j                              S t        | t              xs! t        j                  | t              t        k(  S )z:v can be statically proven to be a multiple of ALIGN_BYTES)
isinstancesympyAddre   allmap_is_alignedargsaligngcdr|   )vs    rY   r   r      sS    !eiiC013{AFF+,,aK599Q#<#KKr   c                  *    e Zd ZdZdZdZedd       Zy)r   z<Symbolically round up to the nearest multiple of ALIGN_BYTESr9   Tc                    t        |t        t        j                  f      rt	        t        |            S t        |      r|S y N)r   intr   Integerr~   r   )clsvalues     rY   evalz
align.eval   s6    ec5==12#e*%%uL r   N)r   
sympy.Exprreturnzsympy.Expr | None)__name__
__module____qualname____doc__nargs
is_integerclassmethodr    r   rY   r   r      s!    FEJ r   r   T)frozenc                  :    e Zd ZU dZded<   ded<   ded<   ded<   y	)
GraphPartitionMapzP
    Mapping from the partition info (e.g., input/output) to the graph info
    r   idzlist[int | None]input_index_mappingoutput_index_mapping	list[str]constant_namesNr   r   r   r   __annotations__r   r   rY   r   r      s$    
 	G *)** r   r   c           
         |         t         j                  j                          t        j                  t	        d      t         j
                  d      }t         j                  j                  d      }t         j                  j                  d      }|j                          t        d      D ]  }|j                           |          |j                          t         j                  j                          |j                  |      dz  }t        dt	        ||z              }t        dt	        ||z              }	t        |      D ]	  } |          t        |	      D cg c]"  }t         j                  j                  d      $ }}t        |	      D cg c]"  }t         j                  j                  d      $ }}t         j                  j                  t         j                  j                  j                  g      5 }
t         j                  j                          t        |	      D ]q  }|j                          ||   j                          t         j                  j                   j                  d	      5   |         d
d
d
       ||   j                          s t         j                  j                          t        j"                  t%        ||      D cg c]  \  }}|j                  |       c}}      }d
d
d
       t        j&                        j)                         }t*        j-                  d       t*        j-                  
j/                         j1                  dd             t3        |
j5                         D cg c]A  }|j6                  t8        j                  k(  r"t;        j<                  d|j>                        |C c}      }|r"|tA        j&                  d |D              dz  z  }t*        j-                  d|       |S c c}w c c}w # 1 sw Y   xY wc c}}w # 1 sw Y   3xY wc c}w )aR  
    Returns benchmark results by examining torch profiler events.
    This could be more accurate as it doesn't count CPU side overhead.
    However, this also requires manually excluding irrelevant event, e.g.
    vectorized_elementwise_kernel which is used to fill L2 cache,
    various CUDA events, etc, so could also be fragile.
        ArJ   dtypedeviceTenable_timing   r9   
activitiesRunCudaModuleN
raw eventsself_device_time_totalsort_by	row_limitzfused_abs_max_\dc              3  4   K   | ]  }|j                     y wr   )device_time_total).0events     rY   	<genexpr>zfp8_bench.<locals>.<genexpr>-  s     QE33Q        @@profiling results: %s ms)!rR   rJ   synchronizeemptyr   float16Eventrecordrangezero_elapsed_timemaxprofilerprofileProfilerActivityCUDAnvtxtensorzipmeanitemlogdebugkey_averagestabler^   eventsdevice_typer]   rematchname
statistics)fnwarmuprepcachestart_event	end_event_estimate_msn_warmupn_repeatpisetimesresr   filtered_eventss                     rY   	fp8_benchr      sT    D	JJKKJu}}VLE **"""6K

  t 4I1X 
 	JJ**959K 1c&;./0H1c#+,-H 8_ 
 BGxQA5::##$#7QKQ?DXO!!!!5OIO			NN++00
 
  
 
 


 x 	"AKKMN!!#&&7 aL!	" 	

 +.{I+FG41aQ^^AG

" **U

 
 
"CIIlIIann$$-EQS$TU 	
!!Z__4HH0%**=I	 	
	O OOQQQ	

 II(#.JO RO 
 H
 
*	
sE   "'P'PA9P2=PAP2P,9P2AP?P)$P22P<Fc                L    ddl m}m}  |t              }  ||      | |||      S )Nr   )gpu_benchmark_lockmay_distort_benchmarking_result)$torch._inductor.runtime.benchmarkingr   r   _do_bench_using_profiling)r   r   r   is_vetted_benchmarkingr   r   locked_benchs          rY   do_bench_using_profilingr   5  s3    "
 &&?@L8*<8
FC/ r   c           
       
 t               
d
fd|D cg c]4  }|j                  t        k(  r|j                  t        j
                  k(  r|6 }}t        |      |k7  r"t        d| dt         dt        |       d      |D ]
  } |        d}| D ]  }|j                         }|j                         }|j                         }	|j                         |k(  sG|	dk7  sM|
v s|dk(  sW|
v s\|j                         d	k7  sp||j                         |j                         z
  d
z  z  } |dk  rt        dt         d      |d
z  |z  S c c}w )Nc                    | j                   t        j                  k7  ry j                  | j                         | j
                  D ]
  } |        y r   )r   r]   CPUaddr   cpu_children)r   childbenchmark_event_idscollect_cpu_event_idss     rY   r   z;_get_do_bench_profile_result.<locals>.collect_cpu_event_idsY  sG    
.)'' 	)E!%(	)r   z	Expected  z profiling events. Found z events.        gpu_user_annotationr   zContext Syncr   z$Failed to capture device events for rt   )r   r   r   None)r!   r   _DO_BENCH_PROFILE_EVENT_NAMEr   r]   r   rT   RuntimeErrorlinked_correlation_idcorrelation_idactivity_typeend_nsstart_ns)kineto_eventsprofiler_eventsr   expected_device_typer   benchmark_eventsdevice_time_usr  r  r  r   r   s             @@rY   _get_do_bench_profile_resultr  Q  s    ,6<) %::55/ 	  (z#?"@ A)*+85
 	

 " %e$% N K % ; ; =--/++-#77!66%)<<)Q.>EX3X 

.u||~0@@FJJNK" 23O2PPQR
 	
 F"X--Ms   9Ec           	     f   |sddl m}  |        t               }|j                         }t	        |      } |         |j                          t        j                  t        d      t        j                  |      }|j                  d      }	|j                  d      }
|	j                          t        d      D ]  }|j                           |          |
j                          |j                          |	j                  |
      dz  }t        dt        ||z              }t        dt        ||z              }t        |      D ]	  } |          |j                          t        t        j                   j"                  |      }t        j                   j%                  t        j                   j"                  j&                  |g	      5 }t        |      D ]E  }|j                          t        j                   j)                  t*              5   |         d
d
d
       G |j                          d
d
d
       t,        j/                  d       t,        j1                  t2        j4                        r4t,        j/                  j7                         j9                  dd             t;        j                   j<                  j?                         |j?                         |t        t@        |            }t,        j/                  d       t,        j/                  d|       |S # 1 sw Y   OxY w# 1 sw Y   xY w)aB  
    Returns benchmark results by examining torch profiler events.
    This could be more accurate as it doesn't count CPU side overhead.
    The benchmarked function is wrapped in a profiler record_function so
    cache-clearing kernels can be excluded without relying on raw CUDA event
    grouping or kernel names.
    r   )may_ban_benchmarkingr   r   Tr   r   r9   r   Nr   r   r   r   zprofiling time breakdownr   )!r   r  rZ   upperr[   r   rR   r   r   r   r   r   r   r   r   rQ   r   r   r   r   record_functionr  r   r   isEnabledForloggingDEBUGr   r   r  kineto_resultsr   r]   )r   r   r   r   r  r   device_type_upperdevice_interfacer   r   r   r   r   r   r   profile_activityr   results                     rY   r   r     s    "M.K#))+/<D  "KKJuyyME #((t(<K &&T&:I1X 
   "**959K 1c&;./0H1c#+,-H 8_ 
   "u~~>>@QR			NN++//
 
  
 '
 
x 	AKKM//0LM  		 	$$&'  IIl
&		NN""+Cr"R	
 *	

!!((*	

-.	F II()II(&1M) ' 's%   	AL'LL'L$L''L0c                    	 ddl m}  t        j                  j	                  dd       | d uxr% t        t        t        j                  dd       d      S # t        $ r Y yt        $ r}dt        |      v sJ Y d }~yd }~ww xY w)	Nr   )	roi_alignztorchvision::nmsMetatorchvisionr  Fztorchvision::nms does not exist)torchvision.opsr  rR   _C%_dispatch_has_kernel_for_dispatch_keyhasattrrQ   opsImportErrorr  str)r  r   s     rY   has_torchvision_roi_alignr'    s|    -667I6R$ 
EII}d3[*
 	
   0CF:::s   AA 	A?A?&A::A?c                b   | t        j                  d      j                  S t        | t              rt        j                  |       } | j
                  dvrZ| j                  Nt        | j
                        }t        j                  | j
                  |j                  j                               S | S )Nr   )cpumeta)index)
rR   r   r   r   r&  typer+  r[   Workercurrent_devicer   r  s     rY   decode_devicer0    s    ~||C '''&#f%{{/)fll.B3FKK@||FKK/?/F/F/U/U/WXXMr   c                |    t        j                  t        j                  | t        j
                  j                        S r   )	functoolsreduceoperatormulr   SOne)its    rY   sympy_productr9    s#    HLL"eggkk::r   c           	         t        |       t        |      k(  sJ t        j                  t        d t	        | |      D                    S )Nc              3  ,   K   | ]  \  }}||z    y wr   r   )r   abs      rY   r   zsympy_dot.<locals>.<genexpr>  s     >daAE>s   )rT   r   expandsumr   )seq1seq2s     rY   	sympy_dotrB    s8    t9D	!!!<<>c$o>>??r   c                    t        |       t        |      k(  sJ t        j                  j                  }t	        | |      D ]  \  }}||z  |z   } |S )z7Row-major flatten: per-dimension indices -> flat index.)rT   r   r6  Zeror   )indicessizesflatr+  sizes        rY   flatten_indexrI    sT    
 w<3u:%%%77<<D7E* #td{U"#Kr   c                ~    t        |      D cg c]!  \  }}t        | t        ||dz   d       |      # c}}S c c}}w )z=Row-major decomposition: flat index -> per-dimension indices.r9   N)	enumeraterf   r9  )r+  rF  r   rH  s       rY   decompose_indexrL    sE     !'At 	}U1q57^<dC  s   &9c                \    | D ci c]  }t        |      | c}j                         S c c}w r   )r   values)r8  rV   s     rY   uniquerO    s'     !BqE1H!((**!s   )c           
     n   t        | t        j                        st        |t        j                        r2t        t        j                  |       t        j                  |            S t        | t
              rt        |t
              s$J |  dt        |        d| dt        |              t        | |      S )Nz: , )r   r   Exprra   sympifyr   r,  runtime_ceildiv)numberdenoms     rY   rm   rm     s    &%**%E5::)Fu}}V,emmE.BCC fc"z%'= ("T&\N"UG2d5k];= 65))r   c                j   | yt        |       j                  d      d   }i dddddd	d
ddddddd	dddddddddddddddddd d!d"d#d$dd%d&d'd(}|j                  t        |j	                               D ci c]  }|| c}       t        | t               r| S d)||    S c c}w )*Nz*i8rt   r   booli1
float8e4nvfp8e4nvfloat8e5fp8e5float8e4b15fp8e4b15float8e4b15x4
fp8e4b15x4float8_e4m3fnfloat8_e5m2float8_e4m3fnuzfp8e4b8float8_e5m2fnuzfp8e5b16float8_e8m0fnuu8float4_e2m1fn_x2r   fp16bfloat16bf16float32fp32float64fp64int8i8int16i16i32i64u16u32u64)int32int64uint8uint16uint32uint64*)r&  splitupdatelistrN  r   )key	dtype_strtysr   s       rY   _type_ofr  &  sX   
 {Cs#B'Ii 	G 	z	
 	 	 	w 	9 	: 	$ 	D 	6 	F  	6!" 	6#$ 	%& 	'( 3C8 JJd3::<01112S#&3@aI/?,@@ 2s   
B0c                R    | D cg c]  }t        j                  |       c}S c c}w )z
    Gets the shape and stride of a tensor. For non-symbolic tensors, this is
    trivial. But for symbolic tensors, we need to map from SymIntNode into
    sympy.Expr.
    )r   rS  lstr   s     rY   convert_shape_to_inductorr  N  s!     '**EMM!***s   $c                f    t        | t        j                        r| j                  j                  S | S )z
    Convert SymInt to sympy.Expr, leave int as is.

    Unlike sympy.sympify() which converts int to sympy.Integer,
    this function preserves int as int and only converts SymInt to Expr.
    )r   rR   r4   nodeexprvals    rY   convert_symint_to_exprr  Y  s%     #u||$xx}}Jr   c                    ddl m} t        | t              r| S t        | t        j
                        rt        |       S |j                  j                  j                  j                  | d      S )zL
    Like convert_shape_to_symint, but operates on a single expression.
    r9   VN)hint)
virtualizedr  r   r   r   r   graphsizevars	shape_envcreate_symintnode)r   r  s     rY   convert_to_symintr  e  se      a 	

 !U]]+ F	 !!++==ad=Kr   c                >    | D cg c]  }t        |       c}S c c}w )zz
    Takes a list of shapes from Inductor and converts them into symints (or just
    ints if all shapes are static).
    )r  r  s     rY   convert_shape_to_symintr  v  s     +..Qa ...s   c                N    t        d | j                  j                  D              S )z-
    Does this op overload have aliasing
    c              3  8   K   | ]  }|j                   d u  y wr   )
alias_infor   r<  s     rY   r   zis_view.<locals>.<genexpr>  s     FAq||4'Fs   )any_schema	argumentsops    rY   is_viewr    s     F1E1EFFFr   c                     yNFr   )r   s    rY   <lambda>r        r   c                   | j                   dk7  ryt        | j                  t        j                  j
                        s| j                  t        j                  u syt        t        j                  j
                  | j                        }|t        j                  u st        |      rt        fd| j                  D              S t        j                  j                  |j                  v xs  |      S )z
    Do all uses of this op have torch.Tag.pointwise or return True for optional `is_pointwise_fn`

    Uses in views ops will follow the views uses
    call_functionFc              3  6   K   | ]  }t        |        y wr   )is_pointwise_use)r   uis_pointwise_fns     rY   r   z#is_pointwise_use.<locals>.<genexpr>  s     KA#A7Ks   )r  r   targetrR   _ops
OpOverloadr4  getitemr   r  r   usersTag	pointwisetags)user  r  s    ` rY   r  r    s     vv 3::uzz445xGWGW9W%**''4F!!!WV_KKKK99&++-H1HHr   	list[Any]c           	        t         j                  j                         g dfd} j                  | gt	        t         j
                  |||f       }t        | j                  j                        dk(  r2t        | j                  j                  d   j                        dk(  r|f}j                  |       t         j                  j                  i       }|fS )Nc                `    j                  |        j                  dt                     S )Narg)appendplaceholderrT   )r  g
graph_argss    rY   add_tensor_argz)gen_gm_and_inputs.<locals>.add_tensor_arg  s,    #}}s3z?"3455r   r9   r   Tensor)r  torch.Tensorr   r7   )rR   fxGraphr  r#   r  rT   r  returnsr&  r,  outputr6   )r  r   kwargsr  r  gmr  r  s         @@rY   gen_gm_and_inputsr    s     	A%'J6 1??u||^dF^LD 	FNN""#q(&&q)../8;wHHTN			b!	$Bz>r   c                h    | dk(  ry t        |       }|j                         r|j                          y y Nr)  )r[   rS   r   r/  s     rY   r   r     s4    /7$$&$$& 'r   c                    t        |       t        j                  d       t        j                         }t        |      D ]  } | | }t        |        t        j                         }J ||z
  S )Ni9  )r   rR   manual_seedtimeperf_counterr   )modelexample_inputsr   r   t0r   r  t1s           rY   timedr    sr     	d				B5\ 'F 
			B7Nr   c                    t        j                  t        |      D cg c]  }t        | |||       c}      }t        j                  |      |z  }t        ||z  d       |j                         S c c}w )Nz.6f)rR   r   r   r  medianprintr   )	r  r  r   repeatbaseliner   r   timingstooks	            rY   print_performancer    sg     ll>CFmLuneV	4LG << 5(D	TH_S!#99;	 	Ms   A1c                H     t        | |             t        | |fd       y)zKReplace obj.method() with a new method that returns a precomputed constant.c                      S r   r   )r  s   rY   r  z#precompute_method.<locals>.<lambda>  s     r   N)rQ   setattr)objmethodr  s     @rY   precompute_methodr    s     !WS&!#FC(r   c                *    |D ]  }t        | |        y)zFReplace methods with new methods that returns a precomputed constants.N)r  )r  methodsr  s      rY   precompute_methodsr    s     '#v&'r   c                <    t        | |kD        t        | |k        z
  S r   )r   r<  r=  s     rY   cmpr    s    q1u:AE
""r   c                ~    t        | t              r| g|z  S t        |       dk(  r t        |       | d   g      |z  S | S )Nr9   r   )r   r   rT   r,  )rV   rH  s     rY   pad_listliker    sC    !SsTz
1v{tAw!v%%Hr   c                D    t        |       dk(  rg S dd}t        | |      S )Nr   c                n    t        | t              r| S ddlm} t        | |      sJ | j	                         S )Nr9   )rH   )r   r&  	schedulerrH   get_name)elemrH   s     rY   	sort_funcztuple_sorted.<locals>.sort_func  s1    dC K0$ 1222}}r   r  )r  rp   r   r&  )rT   sorted)rV   r  s     rY   tuple_sortedr    s&    
1v{	 !##r   PRV)	covariantc                  &    e Zd Zedd       ZddZy)CachedMethodc                     y r   r   )r   s    rY   clear_cachezCachedMethod.clear_cache  s    ),r   c                     y r   r   selfr   r  s      rY   __call__zCachedMethod.__call__  r  r   N)r   r   r   r  )r   P.argsr  P.kwargsr   r  )r   r   r   staticmethodr  r  r   r   rY   r  r    s    , ,Dr   r  c           	         | j                   }d| dd| i}t        d| d d dj                         |        t        j                  |       || d         }d
fd	}||_        |S )N___cacher   z        def zC_cache_on_self(self):
            try:
                return self.zy
            except AttributeError:
                pass
            rv = fn(self)
            object.__setattr__(self, "z%", rv)
            return rv
        _cache_on_selfc                8    t        |       rt        |        y y r   r#  delattrr  r  s    rY   r  z"cache_on_self.<locals>.clear_cache,  s    4D# r   r  r   r   r  r   execlstripr2  wrapsr  )r   r   ctxwrapperr  r  s        @rY   cache_on_selfr    s    ;;DtfF
C *CF  E "' (+e ,			 FH "ioob!#n&=">?G &GNr   c                    t        |       S )z]
    Variant of cache_on_self for properties. The only difference is the type signature.
    )r  )r   s    rY   cache_property_on_selfr  4  s     r   c                     	 	 	 	 d fd}|S )Nc           	         d d| j                    dd| i}t        d d d dj                         |        t        j                  |       |d	         }dfd
}||_        |S )Nr  r   r  r   z            def inner(self: Any, *args: P.args, **kwargs: P.kwargs) -> RV:
                args_kwargs = (args, tuple(sorted(kwargs.items())))

                if not hasattr(self, "z2"):
                    object.__setattr__(self, "z%", {})

                cache = self.z

                try:
                    return cache[args_kwargs]
                except KeyError:
                    pass

                rv = fn(self, *args, **kwargs)

                cache[args_kwargs] = rv
                return rv
            innerc                8    t        |       rt        |        y y r   r  r  s    rY   r  z<cache_on_self_and_args.<locals>.wrapper.<locals>.clear_cacheb  s    tS!c" "r   r  r  )r   r  r  r  r  
class_names       @rY   r  z'cache_on_self_and_args.<locals>.wrapperC  s     :,a}F3 Rj' (+e ,//2e 4!U #$ )	
, $	#CL1	# (r   )r   FN_TYPE[P, RV]r   r  r   )r  r  s   ` rY   cache_on_self_and_argsr  >  s     
$$	$L Nr   c           
     ^   ddl m} t        | t              rgt	        j
                  t        j                  | D cg c]0  }t        |d      r"|j                  r|j                  j                  2 c}t                     S t        | |j                        r| j                  S t               S c c}w )Nr9   irr  ) r!  r   r  r2  r3  r4  or_r#  r  originsr!   rA   )node_scheduler!  r  s      rY   aggregate_originsr&  l  s     -&LL * 4(TYY 		!! L	
 		
 
M2??	3$$$|s   5B*
c                   t        |       }|dk(  rYd }|D cg c]6  }|j                  dk(  r%d|j                  v r|j                  d    ||      8 }}t        t	        |            }n|dk(  rg }|D ]  }|j                  dk(  sd }d}d|j                  v r|j                  d   d   }n"d|j                  v r|j                  d   d   }d	}|s]t        |d
   t              r|j                  |d
   |z          |j                  |d
   j                  |z           t        t	        |            }n5|dk(  r*|D cg c]  }|j                  dk(  s|j                    }}nt        dj                  dg|z         S c c}w c c}w )Noriginal_atenc                   | j                   d   }d}t        |t        j                  j                        r|j
                  j                  }|S t        |t        j                  j                        rt        |j                               }|S )Nr(  r"  )
r*  r   rR   r  r  _overloadpacketr   HigherOrderOperatorr&  r   )originr(  r  s      rY   get_origin_meta_strz2get_fused_kernel_name.<locals>.get_origin_meta_str  so    "KK8MC-)>)>?#33<< J M5::+I+IJ-,,./Jr   r  rR   r"  source_fn_stackr   fwd_source_fn_stackbackwardr9   inductor_noder   fused)r&  r  r*  r  r!   r   r&  r  r   r   NotImplementedErrorjoin)r%  descriptive_namesall_originsr-  r,  sources	source_fnsuffixs           rY   get_fused_kernel_namer:    s    $M2KO+	 &
yyO+6;;.O,8	  '
 
 G,-	g	%! 	CFyyO+ 	$3 &,= >r BI*fkk9 &,A B2 FI'F ilC0NN9Q<&#89NN9Q<#8#86#AB	C" G,-	o	-&1
"VYY/5QFKK
 
 "!88WI'((G
<
s   ;E=<FFc                   !" t        |       }|D cg c]  }|j                  dk(  s| }}t        j                  t              }t        j                  t              }d!|rt        d |D              }t        |      dk(  r_|d   j                  !t        !d      s/t        !j                        D 	ci c]  \  }}	|	|
 }
}}	|
!_        |j                  !fd       |D ]Z  }d	|j                  v r|j                  d	   |j                  d	   }d}t        |t        j                   j"                        rt%        |j&                        }n=t        |t        j                   j(                        rt%        |j+                               }|r||   j-                  |j*                         d
|j                  v r<|j                  d
   d   j*                  }||   j-                  |j*                         |j                  j/                  d      dk(  s3||j*                     j-                  |j*                         ] !dnd}j0                   d| ddj3                  |j5                                ddj3                  |j5                                d}j0                   dg}t7        |j9                               D ]@  \  }}|j-                  j0                   d| ddj3                  t7        |                    B !ddlm  |j-                  j0                   d       t               }g }t        |  j>                        sddl m!} 	 	 	 	 	 	 d* fd}d+d"d,"fd}| D ]  }	t        |	d      r|	jD                  t        |	jD                  d      r|	jD                  jF                  |	jD                  jF                  D ]  }|j*                  |v r|jI                  |j*                         |j                  jK                  |j*                        }|U |||j*                        \  }}|j-                  j0                   d| d  ||       d!| d        t        |	jD                  d"      s|	jD                  jL                  )|	jD                  jL                  D ]T  }|j                  jK                  |j*                        }|+ |||j*                        \  }}|j-                  d#|z          V  |D ]y  }|jO                  d$%      }|Ct        jP                  jR                  r)|jU                  fd&|jW                         D               Z|j-                  j0                   d|        { |j-                  j0                   d'd(j3                  |              |d)j3                  |      fS c c}w c c}	}w )-aH  
    Retrieves metadata information for a kernel.
    Args:
        node_schedule (Union[Sequence[BaseSchedulerNode], ExternKernel]):
            Either a sequence of BaseSchedulerNode objects or an ExternKernel instance.
        wrapper (PythonWrapperCodegen):
            An instance of PythonWrapperCodegen, used to define the code comment format.
    Returns:
        tuple[str, str]:
            A tuple containing two strings:
                - The first string represents the kernel's metadata.
                - The second string represent the kernel's detailed metadata.
    r  Nc              3  4   K   | ]  }|j                     y wr   )r  )r   ns     rY   r   z&get_kernel_metadata.<locals>.<genexpr>  s     "Cq177"Cr   r9   r   )_inductor_kernel_metadata_node_to_idx_mapc                "    j                   |    S r   )r>  )r=  single_graphs    rY   r  z%get_kernel_metadata.<locals>.<lambda>  s    lTTUVW r   r  r(  	from_nodepartitioner_tagis_backwardzTopologically SortedUnsortedr   z Source Nodes: [rQ  z], Original ATen: []z" Source node to ATen node mapping:   z => r   z Graph fragment:r  c                >   t        | j                        rAt        | j                  j                        r!| j                  j                  j                  }n| j                  }||}n|j
                  }	 | j                         }||fS # t        $ r d }Y ||fS w xY wr   )r   	TensorBoxdata
StorageBoxorigin_noder   
get_layoutr3  )bufferrw_namerK  r   layoutr!  s        rY   get_buffer_infoz,get_kernel_metadata.<locals>.get_buffer_info  s     fbll3
KK9 #)++"2"2">">K"("4"4K&"D&++D"#..0F V|# + "!FV|#"s   7B BBc           	     d    ddj                  | D cg c]  }t        |       c}       dS c c}w )N[rQ  rE  )r4  r&  )shaperV   s     rY   stringify_shapez,get_kernel_metadata.<locals>.stringify_shape  s-    499e%<c!f%<=>a@@%<s   -
c                    | y | j                          } | j                         }| j                   }dt        | j                      | | | dS )Nr"  ")rH  strider   r    r   )rO  shape_annotationstride_annotationdevice_annotationrT  s       rY   stringfy_layoutz,get_kernel_metadata.<locals>.stringfy_layout  sl    >&5fkk&B%C '6v}}'E&F!'-}}o! FLL123C2D()*;)<A?r   read_writesreadsz   %z
 : Tensor z = PlaceHolder[target=writes%T)include_tensor_metadatac              3  @   K   | ]  }j                    d |   yw)rF  N)comment)r   liner  s     rY   r   z&get_kernel_metadata.<locals>.<genexpr>O  s)      ) 's4&1)s   z
   return ,
)rM  z-ir.TensorBox | ir.Buffer | ir.TorchBindObjectrN  r&  r   ztuple[str, ir.Layout | None])rS  zIterable[int]r   r&  )rO  zir.Layout | Noner   r&  ),r&  r  collectionsdefaultdictr  r!   rT   r  r#  rK  nodesr>  sortr*  r   rR   r  r  r&  r*  r+  r   r  getrb  r4  keysr  itemsr"  r!  rA   r  r  r\  r]  r   try_get_bufferr^  format_nodeversionru   extend
splitlines)#r%  r  r6  r,  inductor_nodesfrom_node_dictoriginal_aten_dictunique_graphsidxr=  node_to_idx_mapr  r(  r  sort_strmetadatadetailed_metadataoriginal_noderh  	all_reads
all_writesr  rP  r[  rrM  
input_namerO  woutput_namer   formatted_noder!  r@  rT  s#    `                              @@@rY   get_kernel_metadatar    sZ   $ $M2K+6W&)):VfWNW ,,T2N$006
 L""CN"CC}")!,22L<)TU8A,BTBT8U"Vfc11c6"V"VIXFW     8dii'DIIo,F,R IIo6MC-)>)>?-778M5::+I+IJ-,,./"3'..tyy9$))#))K(+00C3&&tyy1YY]],->499%,,TYY78  *6)A%zH??
1XJ&6tyyATATAV7W6X Y99%7%<%<%>?@	C  $OO,,NOP &~';';'= > 
u  s=/diiu6N5OP	

   GOO#44D!EF%/\	 "
-9&$E$PS$-$(A
 # =q-0AMM4I1=='2q}}7J7J7V]]00 66Y.$!aff-!"!7!7!?!>$-<VQVV-L*
F)00&/tJ<z.v677Mj\YZ\ AMM84,,8]]11 =!"!7!7!?!>$)8)HQ"))#*;<=-=< # 	RD!--d-KN)emm.?.? "(( ) . 9 9 ;) 
 "((GOO+<C?O)PQ	R 	  GOO#4Jsxx
?S>T!UVTYY0111Y X #Ws   W	W	8Wc                    t        |       } t        |       }| rV| j                         }|j                  D ]4  }|r	 ||      r||vs|j	                  |       | j                  |       6 | rV|S )zJReturns the set of nodes whose values depend on those within initial_queue)r  r!   rU   r  r   r  )initial_queueskip_filterdominated_setr  users        rY   dominated_nodesr  [  sz    
 'M}-M
  "JJ 	+D{40=(!!$'$$T*	+  r   c                4  	 ddl m d	fd	t        |      \  }}|D cg c]  } 	|      s|j                   }}t        |       \  }}|D cg c]  } 	|      s|j                   }}t	        t        j                  g ||       S c c}w c c}w )Nr9   r   c                F   t        | j                        r | j                        S t        | j                        r | j                        S t        | j                        xr9 t        | j
                  j                  j                  j                  f       S r   )	r   rH  rI  rJ  rB   ComputedBufferInputsKernelInputBufferTemplateBuffer)r=  r!  is_unrealized_nodes    rY   r  z*gather_origins.<locals>.is_unrealized_nodet  s    a&%aff--a'%aff--!RYY' 

!!!!	1
 -
 	
r   )r=  rB   r   rX  )r"  r!  r"   r$  r!   	itertoolschain)
r   r  kwargs_flattenr   r  kwargs_originsargs_flattenargs_originsr!  r  s
           @@rY   gather_originsr  o  s     
" %V,NA-;Wc?QRU?VckkWNW"4(OL!+7SC;Mc;RCKKSLSiooE|EnEFF XSs   BBB Bc                J    dddfddfddfd |       S )z
    Normal sympy str is very slow, this is a lot faster.  The result are
    somewhat worse, as it doesn't do as much simplification.  So don't
    use this for final codegen.
    c                    t        | t        j                        xr, t        | j                        dk(  xr | j                  d   dk(  S )N   r   r   )r   r   MulrT   r   )r  s    rY   is_neg_leadzsympy_str.<locals>.is_neg_lead  s:    tUYY'VC		Na,?VDIIaLTVDV	
r   c                `   t        | t        j                        rt        | j                        dk(  rO | j                  d         r: | j                  d          d | j                  d   j                  d          S dj                  t        | j                              S  |       S )Nr  r9   r   z - z + )r   r   r   rT   r   r4  r   )r  r  sympy_str_muls    rY   sympy_str_addz sympy_str.<locals>.sympy_str_add  s    dEII& 499~"{499Q<'@'		!56c-		RSHYHYZ[H\:]9^__zz#mTYY"?@@ &&r   c                    t        | t        j                        rE |       rd | j                  d          S dj	                  t        | j                              S  |       S )N-r9   z * )r   r   r  r   r4  r   )r  r  sympy_str_atoms    rY   r  z sympy_str.<locals>.sympy_str_mul  s[    dEII&4  >$))A,7899zz#ndii"@AA!$''r   c                   t        | t        j                        r| j                  S t        | t        j                  t        j
                  f      rd |        dS t        | t        t        t        t        f      rC| j                  j                   ddj                  t        t        | j                               dS t!        |       S )N()rQ  )r   r   Symbolr   r   r  rf   rb   rc   rd   funcr   r4  r   	sympy_strr   r&  )r  r  s    rY   r  z!sympy_str.<locals>.sympy_str_atom  s    dELL)99uyy%))45}T*+1--(HMNii(()499SDII5N+O*PPQRRt9r   )r  r   r   rX  r  r   r   r&  r   )r  r  r  r  r  s    @@@@rY   r  r    s$    

	'	( r   c                    ddl m} t        j                  r3t	        |j
                  dd       x}r|j                  dk7  rt        |       S t        j                         S )Nr9   r  current_node
index_expr)
r  r  rl   compute_all_boundsrQ   interpreterr  ri   rj   unknown)r+  r  fx_nodes      rY   get_bounds_index_exprr    sN     	!!~tDDWDNNl*5!!""$$r   c                    | d   dk(  S )Nr   r~  r   prefixs    rY   prefix_is_reductionr    s    !9r   c                    | d   dv S )Nr   )rV   yzr   r  s    rY   prefix_is_pointwiser    s    !9''r   c                J    | t         j                  k7  sJ t        | |dd      S )9
    Used to generate an integer-nonnegative symbol.
    Tintegernonnegative)rh   SIZErg   )r  rv  s     rY   sympy_index_symbol_with_prefixr    s)     TYY vsDdCCr   c                N    | xs t         j                  xr t         j                  S r   )rl   debug_index_assertsassert_indirect_indexing)checks    rY   generate_assertr    s    /V//TV5T5TTr   c                F    | d   dk7  sJ t        j                  | dd      S )r  r   r   Tr  )r   r  r   s    rY   sympy_index_symbolr    s)     7c>> <<d==r   c                    t        | |      S )z
    When the passed replacement symbol v is a string, it is converted to a symbol with name v that
    have the same replaced expression integer and nonnegative properties.
    r'   )r  replacementss     rY   
sympy_subsr    s    
 t\**r   c                    t        | t        j                        xs( t        | t        j                        xr | j                  S r   )r   rR   r4   r  _has_symbolic_sizes_strides)r<  s    rY   is_symbolicr    s3    a& 1ell#E(E(Er   c                 &    t        d | D              S )Nc              3  2   K   | ]  }t        |        y wr   )r  r  s     rY   r   z"any_is_symbolic.<locals>.<genexpr>  s     ,!{1~,   r  )r   s    rY   any_is_symbolicr    s    ,t,,,r   )z,aten._fused_moving_avg_obs_fq_helper.defaultz7aten._fused_moving_avg_obs_fq_helper_functional.defaultzfbgemm.dense_to_jagged.defaultz%fbgemm.jagged_to_padded_dense.defaultrun_and_save_rng_staterun_with_rng_statezaten._local_scalar_densezaten._assert_scalarzaten.topk.defaultc                    ddl m} | j                  j                  D ];  }t	        |      r|c S |j
                  j                  d      x}0 ||      s9|c S  y )Nr   )r*   r  )%torch.fx.experimental.symbolic_shapesr*   r  rh  is_cudagraph_unsafe_fx_noder*  rj  )r  r*   r  r  s       rY   %get_first_incompatible_cudagraph_noder    sV     L &t,K99==''C49Ns9SK r   c                    t        t        t        | j                  j                                    }|j
                  dk(  sJ |S )z$Get the output node from an FX graphr  )nextiterreversedr  rh  r  )r  	last_nodes     rY   output_noder  +  s6    T(288>>234I<<8###r   c                    | j                   j                  d      }t        d |D              }t        |       j                  d   }t        |t              r|n|f}t        d |D              }||z  S )Nr  r  c              3     K   | ]P  }t        |j                  j                  d       t        j                        r|j                  d    j
                   R ywr  N)r   r*  rj  rR   r  r   )r   r  s     rY   r   z"get_all_devices.<locals>.<genexpr>4  sB      9diimmE*ELL9 			%9s   AAr   c              3     K   | ]t  }t        |t        j                  j                        rNt        |j                  j                  d       t        j                        r|j                  d    j                   v ywr  )r   rR   r  r7   r*  rj  r  r   )r   r  s     rY   r   z"get_all_devices.<locals>.<genexpr><  sS      7c588==)sxx||E*ELL9 	7s   A:A<)r  
find_nodesr!   r  r   r   tuple)r  placeholder_nodesinput_devicesout_argout_argsout_devicess         rY   get_all_devicesr  2  s}    ++}+=.8 9%9 /M "o""1%G$We4w7*H,6 77 -K ;&&r   c                 |   t        t        j                  j                               D ]  } | j	                  d      st        j                  |    }|j
                  D ]  }|j	                  d      st        ||      }t        |t        j                  j                  j                  j                        sZ|j                  D ]i  }t        |t        j                  j                  j                  j                        s<|j                  j                   j"                  j%                          k  t        j                  | =  dt        j                  v rRt        j                  d   }t'        |j(                  j*                  j,                        `|j(                  j*                  `t1        j2                          y )Nz&torch._inductor.runtime.compile_tasks.triton_ztriton.runtime.driver)r  sysmodulesrk  
startswith__dict__rQ   r   rR   	_inductorruntimetriton_heuristicsCachingAutotunercompile_resultsTritonCompileResultkernelrunmod__del__r,  driveractiveutilsinstancegccollect)module_namem	attr_namer  r  r  s         rY   unload_xpu_triton_pydsr  H  sK   CKK,,./ %%%&NOKK$ 	<I##I. I.EOO33EEVV #)"8"8 <%"!OO33EEYY
 #MM--1199;<	< KK$#%( #++-kk12""(()2JJ#JJLr   _registered_cachesc                    t        | d      rt        | j                        st        |  d      t        j                  |        | S )zh
    Use this decorator to register any caches that should be cache_clear'd
    with fresh_cache().
    cache_clearz# does not have a cache_clear method)r#  callabler  AttributeErrorr  r  r  s    rY   clear_on_fresh_cacher  i  s?    
 3&hs.Gu$GHIIc"Jr   c                 :    t         D ]  } | j                           y)z&
    Clear all registered caches.
    N)r  r  r  s    rY   clear_cachesr  u  s     " r   c              #  \  K   t         j                  j                  |       }	 |t         j                  | <   d |!t         j                  j                  | d       y|t         j                  | <   y# |!t         j                  j                  | d       w |t         j                  | <   w xY ww)a  Thread-safe env var set/restore using atomic C-level lookups.

    We avoid mock.patch.dict(os.environ, ...) because it internally calls
    os.environ.copy(), which iterates all env var keys then fetches values in
    separate steps. That approach is not atomic and can race with background threads
    (e.g. Triton async compilation) modifying the environment, causing KeyError,
    so we use os.environ.get() for individual keys which is an atomic C-level lookup.
    N)osenvironrj  rU   )r  r   olds      rY   _set_envr  }  sz      **..
C"

3;JJNN3%!BJJsO ;JJNN3%!BJJsOs    B,A1 7B,18B))B,c              #    K   t                ddlm}  |t        j                  |            	 t        d      5  t        j                  d        |t        j                  j                  d            }t        d|      5  d t        | t              rt        |       dk(  sJ d	       t        j                  j                  |      rtt        j                  |      }| j!                  |D ci c]D  }d
|vr>|t        j                  j#                  t        j                  j                  ||            F c}       ddd       ddd       |rUt%               r(t&        j(                  j+                         r
t-                t/        j0                  t%               fd       t                yc c}w # 1 sw Y   xxY w# 1 sw Y   |xY w# t2        $ r t        j5                  d        w xY w# t                w xY ww)z
    Contextmanager that provides a clean tmp cachedir for pt2 caches.

    Optionally, pass a dict as 'cache_entries' to get a list of filenames and sizes
    generated with this cache instance.
    r   )normalize_path_separator)dirTORCHINDUCTOR_CACHE_DIRzUsing inductor cache dir %stritonTRITON_CACHE_DIRNz!expected empty cache_entries dictz.lockc                4    t         j                  d|      S )Nz*Failed to remove temporary cache dir at %s)exc_info)r   warning)r  pathr  inductor_cache_dirs      rY   r  zfresh_cache.<locals>.<lambda>  s    S[[@&% 6A 6 r   )ignore_errorsonerrorz(on error, temporary cache dir kept at %s)r  torch._inductor.cpp_builderr  tempfilemkdtempr  r   r   r  r  r4  r   dictrT   existslistdirr  getsize
is_windowsrR   rL   rS   r  shutilrmtree	Exceptionr  )cache_entriesr  deleter  triton_cache_dirfilesfr  s          @rY   fresh_cacher1    s     ND1(2B2Bs2KL'/1CD 	II35GH7/:  ,.>? mT2}-2W4WW2ww~~&67 "

+; <%,, */$%#*!#3 !"277??277<<@PRS3T#U U	$ |		 6 6 8&(MM" )l  	5 	 	D  >@RS 	sm   -HG A	GA-F:3A	F5<F:GAG *H5F::G	?GGG !G33G6 6HH)reversec                   | j                   }t        t        |             }t        t	        ||d            }|st        t        |            S |S )NTr  r2  )__getitem__r   rT   r  r  r  )seqr2  gettera_rsort_idxs        rY   argsortr:    sE    __F
C/C F3FD9:HHX&''Or   c          	     4    d fd}t        |      D cg c]9  \  }}|t        |t        j                        r|j                  j
                  n|f; }}}t        |t        j                  |      |      }|D cg c]  \  }}|	 }}}|S c c}}w c c}}w )z
    Return a symbolic sort order for optimization-only layout heuristics.

    Data-dependent comparisons can fall back to non-guarding optimization hints,
    so callers must not use this order to make correctness decisions.
    c                    | \  }|\  }	 	 	 	 	 	 dfd}d	fd |k  fd      ry |kD  fd      ry||k  ry||kD  ryy)
Nc                x    t        | t              r| S 	 j                  |       S # t        $ r
  |       cY S w xY wr   )r   rX  evaluate_exprr+   )r  fallbackr  s     rY   evaluatez*argsort_sym.<locals>.cmp.<locals>.evaluate  sA     $%" ..t44. "z!"s   & 99c                J    j                  |       j                  |      k  S r   )optimization_hint)lhsrhsr  s     rY   hint_ltz)argsort_sym.<locals>.cmp.<locals>.hint_lt  s'     ..s3i6Q6QRU6VVVr   c                             S r   r   a_valb_valrE  s   rY   r  z*argsort_sym.<locals>.cmp.<locals>.<lambda>      75%+@ r   r   c                             S r   r   rG  s   rY   r  z*argsort_sym.<locals>.cmp.<locals>.<lambda>  rJ  r   r9   r   )r  z bool | torch.SymInt | sympy.Exprr?  zCallable[[], bool]r   rX  )rC  r   rD  r   r   rX  r   )	r<  r=  a_idxb_idxr@  rH  rI  rE  r  s	        @@@rY   r  zargsort_sym.<locals>.cmp  sy    uu		"2		"(		" 		"	W EEM#@AEEM#@A
 5=5=r   r4  )r<  tuple[int, sympy.Expr]r=  rN  r   r   )	rK  r   rR   r4   r  r  r  r2  
cmp_to_key)	r  r6  r2  r  rv  r   exprsr   r  s	   `        rY   argsort_symrQ    s    !L  nC 
Z5<<8affkka@E  5i2237IE %&fc1c&F&M
 's   >B=Bc                t    | t         j                  k(  ryt        j                  d|       j                         S )Nrz   r   r   )rR   r  r   element_sizerS  s    rY   get_dtype_sizerU    s-     ;;r'4466r   c                      e Zd ZU ded<   y)LineContextr   contextNr   r   r   r   r   r   rY   rW  rW  &  s    Lr   rW  c                  "    e Zd ZU ded<   ded<   y)ValueWithLineMapr&  r   zlist[tuple[int, LineContext]]line_mapNrY  r   r   rY   r[  r[  *  s    J++r   r[  c                      e Zd ZdZdddZej                  dd       ZddZddZ	ddZ
d Zd dZd!d	Zdd
Zd dZd"dZd"dZd"dZd#d$dZd#d$dZd%dZd&d'dZd&d(dZd&d(dZd#d$dZd)dZddZd*dZd+dZy),IndentedBuffer   c                     g | _         || _        y r   )_lines_indent)r  initial_indents     rY   __init__zIndentedBuffer.__init__3  s    BD%r   c              #  b   K   | j                   }	 || _         d  || _         y # || _         w xY wwr   )tabwidth)r  rf  prevs      rY   set_tabwidthzIndentedBuffer.set_tabwidth7  s,     }}	!$DM DMDDMs   /# /	,/c                   t               }d}g }| j                  D ]  }t        |t              r
 |       }|1t        |t              r|j                  ||j                  f       K|}t        |t              sJ |j                  |       |j                  d       |d|j                  d      z   z  } t        |j                         |      S )Nr9   re  )r   ra  r   DeferredLineBaserW  r  rX  r&  writecountr[  getvalue)r  bufr   linemaplirc  s         rY   getvaluewithlinemapz"IndentedBuffer.getvaluewithlinemap@  s    j13++ 	&B"./t<B,2::/dC(((IIdOIIdOTZZ%%%A	&  88r   c                6    | j                         j                  S r   )rq  r   r  s    rY   rm  zIndentedBuffer.getvalueT  s    '')///r   c                f   t               }| j                  D ]  }t        |t              r
 |       }|t        |t              r.|}t        |t
              sJ |j                  d      r|j                  |d d        h|j                  |       |j                  d        |j                         S )N\r   re  )	r   ra  r   rj  rW  r&  endswithrk  rm  )r  rn  rp  rc  s       rY   getrawvaluezIndentedBuffer.getrawvalueW  s    j++ 	 B"./t<B,dC(((}}T"		$s)$		$		$	   ||~r   c                    | j                   S r   ra  rs  s    rY   get_lines_refzIndentedBuffer.get_lines_refk  s    {{r   c                8    | j                   j                          y r   )ra  clearrs  s    rY   r|  zIndentedBuffer.clearn  s    r   c                ,    t        | j                        S r   )rX  ra  rs  s    rY   __bool__zIndentedBuffer.__bool__q  s    DKK  r   c                :    d| j                   | j                  z  z  S )Nr   )rb  rf  rs  s    rY   r  zIndentedBuffer.prefixt  s    dllT]]233r   c                &    | j                  d       y )Nre  	writeliners  s    rY   newlinezIndentedBuffer.newlinew      tr   c                   t        |t              r| j                  j                  |       y t        |t              r9| j                  j                  |j                  | j                                      y |j                         r.| j                  j                  | j                          |        y | j                  j                  d       y Nr"  )r   rW  ra  r  rj  with_prefixr  stripr  rc  s     rY   r  zIndentedBuffer.writelinez  s    dK(KKt$./KKt//>?ZZ\KK$++-78KKr"r   c                &    | j                  |       y)zGWrite to JIT buffer only. On a plain IndentedBuffer, same as writeline.Nr  r  s     rY   writeline_jitzIndentedBuffer.writeline_jit  s    tr   c                     y)z;Write to AOTI buffer only. No-op on a plain IndentedBuffer.Nr   r  s     rY   writeline_aotzIndentedBuffer.writeline_aot  r  r   c                *    | j                  ||       y)zESplice to JIT buffer only. On a plain IndentedBuffer, same as splice.r  Nsplicer  
other_coder  s      rY   
splice_jitzIndentedBuffer.splice_jit  s    Je,r   c                     y)z<Splice to AOTI buffer only. No-op on a plain IndentedBuffer.Nr   r  s      rY   
splice_aotzIndentedBuffer.splice_aot  r  r   c                4    |D ]  }| j                  |        y r   r  r  linesrc  s      rY   
writelineszIndentedBuffer.writelines       	!DNN4 	!r   c                H     t         j                  d fd       } |       S )Nc               3     K   xj                    z  c_         	 d  xj                    z  c_         y # xj                    z  c_         w xY wwr   rb  offsetr  s   rY   r  z"IndentedBuffer.indent.<locals>.ctx  s9     LLF"L'&&s   A4 AAAr   Iterator[None]
contextlibcontextmanagerr  r  r  s   `` rY   indentzIndentedBuffer.indent  s$    		"	"	' 
#	' ur   c                .    | xj                   |z  c_         y r   r  r  r  s     rY   	do_indentzIndentedBuffer.do_indent      r   c                .    | xj                   |z  c_         y r   r  r  s     rY   do_unindentzIndentedBuffer.do_unindent  r  r   c           	        t        |t              rt        d      }|j                  D ]E  }t        |t              r|st        |t        |      t        |j                               z
        }G t        j                  |      rd}|j                  D ]P  }t        |t              r| j                  j                  |       /t        j                  | |t        |      d         R y t        j                  |      }|r|j                         }|sy |j                         }|j!                  d      D ]  }t        j                  | |        y )Ninfr   re  )r   r^  floatra  rW  minrT   r  mathisinfr  r  r   textwrapdedentrstripr  )r  r  r  r  rc  r   s         rY   r  zIndentedBuffer.splice  s   j.15\F")) I!$4 TS5G)GHFI zz&!")) HdK0KK&&t,",,T4F3FG	H "4J'..0
#**,J%%d+ 2((q12r   c                    t        | j                        }| j                  D cg c]
  } ||       c}|_        |S c c}w N)rc  )r^  rb  ra  )r  r  r   rc  s       rY   r   zIndentedBuffer.map  s4    DLL9-1[[9Td4j9

 :s   >c                @    t        |        d| j                          dS )Nr  r  )r,  rm  rs  s    rY   __repr__zIndentedBuffer.__repr__  s     t*Qt}}/q11r   c                    | j                   |j                   k(  sJ t        | j                         }|j                  | j                         |j                  |j                         |S r  )rb  r^  r  ra  )r  otherr   s      rY   __add__zIndentedBuffer.__add__  sK    ||u}},,,DLL9t{{#u||$
r   c                    || j                   v S r   ry  )r  new_lines     rY   containszIndentedBuffer.contains  s    4;;&&r   Nr   rc  r   r   r  )rf  r   r   r  )r   r[  r   r&  r   r  r   rX  )rc  z$LineContext | DeferredLineBase | strr   r  F)r  zIndentedBuffer | strr  rX  r   r  )r  z.Sequence[LineContext | DeferredLineBase | str]r   r  r   )r  r   r   'contextlib.AbstractContextManager[None]r  r   r   r  )r  zCallable[[Any], Any]r   r^  )r  r   r   r^  )r  z$DeferredLineBase | LineContext | strr   rX  )r   r   r   rf  rd  r  r  rh  rq  rm  rw  rz  r|  r~  r  r  r  r  r  r  r  r  r  r  r  r  r   r  r  r  r   r   rY   r^  r^  0  s    H& ! !9(0(!4#J-K!	20
2'r   r^  c                       e Zd ZdZdd fdZedd       Zd ZddZddZ	d Z
dddZddd	Zddd
ZdddZdddZdddZd fdZ xZS )DualIndentedBufferzIndentedBuffer that simultaneously accumulates JIT and AOTI output.

    The base class (self._lines) holds JIT content. self.aot holds AOTI content.
    By default, writeline/splice write to both. Use _jit/_aot variants for
    mode-specific writes.
    c                D    t         |   |       t        |      | _        y r   )superrd  r^  aot)r  rc  	__class__s     rY   rd  zDualIndentedBuffer.__init__  s    (!.1r   c                    | S )u   Read-only accessor for JIT buffer (for splicing into result).

        WARNING: Do NOT use jit for writing — writes would go to both buffers
        because self.writeline is overridden. Use writeline_jit/splice_jit instead.
        r   rs  s    rY   jitzDualIndentedBuffer.jit  s	     r   c                f    t         j                  | |       | j                  j                  |       y r   )r^  r  r  r  s     rY   r  zDualIndentedBuffer.writeline  s$      t,4 r   c                0    t         j                  | |       y r   )r^  r  r  s     rY   r  z DualIndentedBuffer.writeline_jit  s      t,r   c                :    | j                   j                  |       y r   )r  r  r  s     rY   r  z DualIndentedBuffer.writeline_aot  s    4 r   c                4    |D ]  }| j                  |        y r   r  r  s      rY   r  zDualIndentedBuffer.writelines  r  r   c                    t         j                  | ||       t        |t              r|j                  n|}| j                  j                  ||       y Nr  )r^  r  r   r  r  r  r  r  	aot_others       rY   r  zDualIndentedBuffer.splice  sG     	dJe<(5GHJNNj 	 		/r   c                4    t         j                  | ||       y r  )r^  r  r  s      rY   r  zDualIndentedBuffer.splice_jit  s    dJe<r   c                z    t        |t              r|j                  n|}| j                  j                  ||       y r  )r   r  r  r  r  s       rY   r  zDualIndentedBuffer.splice_aot  s0    (5GHJNNj 	 		/r   c                F     t         j                   fd       } |       S )Nc               3  Z  K   xj                    z  c_         j                  xj                    z  c_         	 d  xj                    z  c_         j                  xj                    z  c_         y # xj                    z  c_         j                  xj                    z  c_         w xY wwr   rb  r  r  s   rY   r  z&DualIndentedBuffer.indent.<locals>.ctx	  s{     LLF"LHH&+&  F*  &  F* s   5B+A2 5B+26B((B+r  r  s   `` rY   r  zDualIndentedBuffer.indent  s$    		"	"	+ 
#	+ ur   c                l    | xj                   |z  c_         | j                  xj                   |z  c_         y r   r  r  s     rY   r  zDualIndentedBuffer.do_indent  %    F"r   c                l    | xj                   |z  c_         | j                  xj                   |z  c_         y r   r  r  s     rY   r  zDualIndentedBuffer.do_unindent  r  r   c                V    t         |           | j                  j                          y r   )r  r|  r  r  r  s    rY   r|  zDualIndentedBuffer.clear  s    r   r  r  r   r^  r  r  r  rX  r   r  r   )r  r   r  )r   r   r   r   rd  propertyr  r  r  r  r  r  r  r  r  r  r  r|  __classcell__r  s   @rY   r  r    s]    2  !-!!0=0## r   r  c                  4    e Zd ZdZddZddZdd	dZdd	dZy)
AotOnlyBuffera*  IndentedBuffer for pure-AOTI codegen.

    Mirror of the base class's pure-JIT defaults: writeline_aot/splice_aot
    write to the buffer; writeline_jit/splice_jit are no-ops. Lets call
    sites use writeline_jit/writeline_aot uniformly across pure-JIT,
    pure-AOTI, and dual-wrapper modes.
    c                     y r   r   r  s     rY   r  zAotOnlyBuffer.writeline_jit+      r   c                &    | j                  |       y r   r  r  s     rY   r  zAotOnlyBuffer.writeline_aot.  r  r   c                     y r   r   r  s      rY   r  zAotOnlyBuffer.splice_jit1  r  r   c                *    | j                  ||       y r  r  r  s      rY   r  zAotOnlyBuffer.splice_aot4  s    Je,r   Nr  r  r  )r   r   r   r   r  r  r  r  r   r   rY   r  r  "  s    -r   r  c                 b    ddl m}  | j                  j                  r
t	               S t               S )zConstruct the IndentedBuffer subclass matching the current codegen mode.

    Pure AOTI -> AotOnlyBuffer (writeline_aot writes; writeline_jit drops).
    Pure JIT  -> IndentedBuffer  (writeline_jit writes; writeline_aot drops).
    r9   r  )r  r  r  aot_moder  r^  r  s    rY   make_codegen_bufferr  8  s%     wwr   c                  (     e Zd Zd fdZddZ xZS )FakeIndentedBufferc                "    t         |           y r   )r  rd  r  s    rY   rd  zFakeIndentedBuffer.__init__F  s    r   c                V    |dk(  rt         j                  | |      S t        d| d      )Nr  zTried to call self.z on FakeIndentedBuffer. This bufferis currently used on TritonTemplateKernel to prevent actualwrites to the body without explicitly specifying the body with`TritonTemplateKernel.set_subgraph_body(name)`)object__getattribute__r  )r  r   s     rY   r  z#FakeIndentedBuffer.__getattribute__I  s;    ;**466!$ (= =
 	
r   r  )r   r&  r   r   )r   r   r   rd  r  r  r  s   @rY   r  r  E  s    
r   r  c               #     K   t         j                  t         j                  }} 	 d  | |ct         _        t         _        y # | |ct         _        t         _        w xY wwr   )r  stdoutstderr)initial_stdoutinitial_stderrs     rY   restore_stdout_stderrr  T  s@     %(ZZNN@!/
CJ
CJs   !AA  A AAc                  P    e Zd ZdZddZddZddZddZddZddZ	ddZ
dd	Zy
)rj  z.A line that can be 'unwritten' at a later timec                6    |j                         sd}|| _        y r  )r  rc  r  s     rY   rd  zDeferredLineBase.__init__`  s    zz|D	r   c                    t         )zJReturns either self.line or None to indicate the line has been 'unwritten'r3  rs  s    rY   r  zDeferredLineBase.__call__e      !!r   c                    t         )z3Returns a new deferred line with the same conditionr  r  s     rY   	_new_linezDeferredLineBase._new_linei  r  r   c                @    | j                  | | j                         S r   r  rc  )r  r  s     rY   r  zDeferredLineBase.with_prefixm  s    ~~455r   c                T    | j                  | j                  j                               S r   )r  rc  r  rs  s    rY   r  zDeferredLineBase.lstripp  s    ~~dii..011r   c                >    | j                  | j                  |         S r   r  )r  r+  s     rY   r5  zDeferredLineBase.__getitem__s  s    ~~dii.//r   c                ,    t        | j                        S r   )rX  rc  rs  s    rY   r~  zDeferredLineBase.__bool__v  s    DIIr   c                ,    t        | j                        S r   )rT   rc  rs  s    rY   __len__zDeferredLineBase.__len__y  s    499~r   N)rc  r&  r   
str | None)rc  r&  r   r   )r  r&  r   r   )r   r   )r+  zint | slicer   r   r  r   r   )r   r   r   r   rd  r  r  r  r  r5  r~  r  r   r   rY   rj  rj  ]  s-    8
""620r   rj  c                  4     e Zd ZdZd fdZddZddZ xZS )DelayReplaceLinez6At end of codegen call `line.replace(key, value_fn())`c                @    t         |   |       || _        || _        y r   )r  rd  r  value_fn)r  r  r  rc  r  s       rY   rd  zDelayReplaceLine.__init__  s     r   c                j    | j                   j                  | j                  | j                               S r   )rc  replacer  r  rs  s    rY   r  zDelayReplaceLine.__call__  s#    yy  4==?;;r   c                D    t        | j                  | j                  |      S r   )r  r  r  r  s     rY   r  zDelayReplaceLine._new_line  s    $-->>r   )r  r&  r  zCallable[[], str]rc  r&  r  )rc  r&  r   r  )r   r   r   r   rd  r  r  r  r  s   @rY   r  r  }  s    @!
<?r   r  c                   t        | t        j                        r| }nt        j                  t               |       }t	        j
                  |      }t        j                  j                  rC|j                  J |j                  dk  s|j                  dk(  rt        j                  d       yy|j                  dk(  rdnd}|j                  }||k  rt        j                  d	||d
       yy)N	   
   z6GPU arch does not support max_autotune_gemm mode usageFTrL   rv   D   z,Not enough SMs to use max_autotune_gemm mode)min_sms	avail_sms)extra)r   rR   r   rZ   r   createro  ru   majorr   r  r,  multi_processor_count)index_or_devicer   propr  r  s        rY   
is_big_gpur"    s    /5<<0 lno>""6*D }}zz%%%::>TZZ2-KKPQKK5(bbG**I7:%I> 	 	
 r   c                     t         j                  j                         r(t         j                  j                         j                  S t         j
                  j                  d      j                  S )NrJ   )rR   rL   rS   get_device_propertiesgpu_subslice_countrJ   r  r   r   rY   get_max_num_smsr&    sF    yyyy..0CCC::++F3IIIr   c                     t         j                  j                         syt         j                  j                  t         j                  j	                               } | j
                  dk(  S )zEReturns true if the device is a NVIDIA B200, otherwise returns false.Fr  )rR   rJ   rS   r$  r.  r  )device_propertiess    rY   
using_b200r)    sJ     ::""$

889R9R9TU""b((r   c                     t         j                  j                         xr? t         j                  j                  du xr! t         j                  j                         dk\  S )zCReturn true for NVIDIA CUDA devices with compute capability SM100+.N)r  r   )rR   rJ   rS   ro  ru   get_device_capabilityr   r   rY   is_nvidia_sm100_or_laterr,    sJ     	

! 	:MM%	:JJ,,.'9r   c                     t         j                  j                         r
t               S t         j                  j                         } t               | | z
  S dz
  S )zFHandle experimental carveout if set otherwise return hardware SM countr   )rR   rL   rS   r&  r!  _get_sm_carveout_experimental)carveouts    rY   get_num_smsr0    sJ     yy  xx557HH,@HHaHHr   c                    ddl m}m} |
t               }|j	                  d      }|| z  t
        z  } |||| |j                               S )zKBuilds and returns a WorkspaceArg for the device side TMA workspace buffer.r9   )r:   WorkspaceZeroModeF)rl  	zero_moder   
outer_name)codegen.commonr:   r2  r0  	from_boolTMA_DESCRIPTOR_SIZEunique_name)num_tma_descriptorsr   num_programsr:   r2  r3  rH  s          rY   get_tma_workspace_argr;    sZ     @"}!++E2I--0CCD+<++-	 r   c                    t         j                  j                  sydt         j                  j	                  d      j
                  v r| dk  ryy)Nr   gfx942rv   r9   r  )rR   ro  ru   rJ   r$  gcnArchName)block_ks    rY   get_default_kpackr@    s<    ==5::33A6BBBwRT}r   c                    | j                   |vr!t        j                  d| j                   |       t        | j                  j
                        xr% | j                   |v xr t        | j                        S )NzDNot using template since dtype %s is not in allowed layout dtypes %s)r   r   r   is_gpur   r,  r"  )rO  allowed_layout_dtypess     rY   _use_template_for_gpurD    sf     ||00		RLL!	
 	v}}!!" 	&LL11	&v}}%r   c                    | j                         t        j                  j                         j                  d      D cg c]  }|j	                          c}v S c c}w Nrd  )r  rl   max_autotune_gemm_backendsr  r  backendrV   s     rY   _use_autotune_backendrJ    M    ==?!<<BBDJJ3O	      Ac                    | j                         t        j                  j                         j                  d      D cg c]  }|j	                          c}v S c c}w rF  )r  rl   max_autotune_conv_backendsr  r  rH  s     rY   _use_conv_autotune_backendrO    rK  rL  c                    | j                         t        j                  j                         j                  d      D cg c]  }|j	                          c}v S c c}w rF  )r  rl   %max_autotune_conv_bwd_weight_backendsr  r  rH  s     rY   %_use_conv_bwd_weight_autotune_backendrR    sP    ==?==CCEKKCP 	
	   rL  c                    | j                         t        j                  j                         j                  d      D cg c]  }|j	                          c}v S c c}w rF  )r  rl   $max_autotune_conv_bwd_input_backendsr  r  rH  s     rY   $_use_conv_bwd_input_autotune_backendrU    sP    ==?<<BBDJJ3O 	
	   rL  )enable_int32enable_float8check_max_autotunec                  ddl m}m} t        j                  t        j
                  t        j                  g}|r>t        j                  t        j
                  t        j                  t        j                  g}|r/|j                  t        j                  t        j                  g       t        | j                  j                        xr t        | |      xs) | j                  j                  dk(  xr | j                  |v xrS t         j"                  xs t         j$                  xs | xr* t'        d      xr  || j                  |j(                        S )Nr9   )BackendFeaturehas_backend_featurer)  TRITON)r5  rZ  r[  rR   r   rl  rn  r{  rp  rb  rc  rB  r   r,  rD  r   rl   max_autotunemax_autotune_gemmrJ  TRITON_TEMPLATES)rO  rV  rW  rX  rZ  r[  layout_dtypess          rY   use_triton_templatera    s    D]]ENNEMMBMu{{Se1153D3DEF v}}))* A)&-@O ""e+M0M
	P   VF$<$<VDV@V
	P "(+
	P  ~/N/NOr   output_layout
add_guardsc                   	 ddl m} ddlm dfddfd}d	fd	 	 	 	 	 	 	 	 dfd	 	 	 	 	 	 	 	 dfd		 |       xr t	        fd
|D              xr  ||       S )u^  
    Return True iff *all* supplied tensors satisfy the CUDA TMA constraints
    that Triton relies on today.
    * https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TENSOR__MEMORY.html

    A tensor is accepted when:
      * 1 ≤ rank ≤ 5 (cuTensorMapEncodeTiled)
      * dtype in _TMA_SUPPORTED_DTYPES (CUtensorMapDataType enum)
      * Base pointer 16-byte aligned
      * Exactly one contiguous ("inner") dim with stride 1
      * All "outer" dims have 16-byte aligned strides
      * Inner dim size × itemsize is a multiple of 16
      * For 1-byte dtypes (e.g. FP8), inner dim ≥ 32
    r   )has_triton_tma_devicer9   r  c                X    j                   j                  j                  | t              S r   )r  r  statically_known_multiple_ofTMA_ALIGNMENT)
expr_bytesr  s    rY   _alignedzcan_use_tma.<locals>._alignedD  s     ww<<ZWWr   c                    | y| j                   }| j                  }| j                  } | j                        sy |||      S )NTF)rH  rW  r   r  )rO  rF  stridesr   rk  _is_tma_compatibles       rY   _is_tma_compatible_layoutz.can_use_tma.<locals>._is_tma_compatible_layoutG  sF    >-- &!%%88r   c                   | j                         }| j                         }| j                         }| j                         j                  j
                  v ry| j                         x}|j                  dk(  r
 |||      S  |||      S )NFrL   )get_size
get_stride	get_dtyper  r  unaligned_buffers
get_devicer,  )r  rF  rm  r   m_devicer  rn  _is_tma_compatible_xpus        rY   _is_tma_compatible_matrixz.can_use_tma.<locals>._is_tma_compatible_matrixT  s{    

,,. ::<177444&H38N)%%@@!%%88r   c                R   t        |       }|j                  }|dk  s|dkD  ry|t        vryrKj                  j                  j                  |       }j                  j                  j                  |      }nd| D cg c]'  }j                  j                  j                  |      ) }}|D cg c]'  }j                  j                  j                  |      ) }}t        |      D 	cg c]-  \  }	}j                  j                  j                  |d      r|	/ }
}	}t        |
      dk7  ry|
d   }t        |      D ]  \  }	}|	|k(  r ||z        r y ||   } ||z        sy|dk(  r'j                  j                  j                  |d      syyc c}w c c}w c c}}	w )Nr9   r   Fr       T)
rT   itemsizerx   r  r  guard_int_seq!replace_backed_symbols_with_hintsrK  statically_known_equalsstatically_known_geq)rF  rm  r   rankr{  sizes_i	strides_ir   str   r  	inner_idx	inner_dimr  rk  rd  s                rY   rn  z'can_use_tma.<locals>._is_tma_compatibleb  s   
 5z>>!8tax--gg&&44U;G((66w?I PUJK  BB1EG  RYKM  BB2FI  #9-
2ww77A> 
 

 u:?!H	 y) 	EArI~BM*		 I&		H,- q=!1!1!F!FyRT!UC

s   >,F0,F,2F#c                T   |d   }j                   j                  j                  |      }j                   j                  j                  |d      syd}| D ]O  }j                   j                  j                  |      }j                   j                  j	                  ||      sO y y)Nr   r9   Fl    T)r  r  r}  r~  statically_known_gt)	rF  rm  r   last_stridelast_stride_hint
MAX_UINT32rH  	size_hintr  s	           rY   rw  z+can_use_tma.<locals>._is_tma_compatible_xpu  s     bk77++MM
 ww778H!L 
 	D((JJ4PIww33IzJ	
 r   c              3  .   K   | ]  } |        y wr   r   )r   r  rx  s     rY   r   zcan_use_tma.<locals>.<genexpr>  s     ?)!,?   )rj  int | sympy.Exprr   rX  )rO  Layout | Noner   rX  )r  rB   r   rX  )rF  Sequence[sympy.Expr]rm  zSequence[_IntLike]r   torch.dtyper   rX  )torch.utils._tritonrf  r  r  r   )
rc  rd  matricesrf  ro  r  rk  rn  rx  rw  s
    `   @@@@@rY   can_use_tmar  /  s    " :X993#3#3 3 
	3j##  
	. 	 	5?h??	5%m4r   c                   t        j                  t         j                        j                  }g }| D ]O  }t	        |t
        t        j                  f      r||kD  s) y|j                  t        j                  ||             Q |syddl
m} t        |      dk(  r|d   nt        j                  | }|r%|j                  j                  j!                  |      S |j                  j                  j#                  |      S )NFTr9   r  r   )rR   iinfor{  r   r   r   r   r   r  Ler  r  rT   Andr  r  guard_or_falsestatically_known_true)rF  rd  	int32_max
conditionsrH  r  	conditions          rY   _descriptor_shape_fits_in_int32r    s     EKK(,,IJ 9dS%--01iehhtY789 !$ZA!5
1599j;QI  	
''	2 WW33I>r   rd  c                x   t         j                  j                  syt        d |D              syt        fd|D              syt         j                  j                  rt        | j                        syt        j                  j                  yt         j                  j                  r| nd }t        ||dS )NFc              3  T   K   | ]   }t        |j                               d k(   " yw)r  N)rT   rq  )r   r  s     rY   r   z*use_triton_tma_template.<locals>.<genexpr>  s      8!s1::< A%8s   &(c              3  T   K   | ]  }t        |j                                 ! yw)r  N)r  rq  )r   r  rd  s     rY   r   z*use_triton_tma_template.<locals>.<genexpr>  s)       	(

LLs   %(r  Trb  )rl   r  enable_persistent_tma_matmulr   enable_template_tma_storer  rH  rR   ro  ru   r  )rc  rd  r  rO  s    `  rY   use_triton_tma_templater    s     ==558x88   }}..7Vz8  }}$$mmEE]4F:NNr   c                T    t        || |dsyddlm} ddlm}  |       xr  |       S )Nrb  Fr   )%has_triton_tensor_descriptor_host_tmar9   is_datacenter_blackwell_arch)r  r  r  codegen.cuda.cuda_envr  )rc  rd  r  r  r  s        rY   !use_triton_blackwell_tma_templater    s2     #	: IC 12U7S7UUr   c                    | |v xr ||v S r   r   )scale_option_ascale_option_bscaling_typess      rY   use_triton_scaling_templater    s    
 ]*N~/NNr   )maxsizec                 d    	 t         j                  j                  d      duS # t        $ r Y yw xY w)zCheck if CuTeDSL is importable; cache the result for reuse.

    Call ensure_cute_available.cache_clear() after installing CuTeDSL
    in the same interpreter to retry the import.
    cutlassNF	importlibutil	find_specr%  r   r   rY   ensure_cute_availabler    s3    ~~''	2$>>     # 	//c                     	 t         j                  j                  d      du} | r
t	                | S # t        $ r Y yw xY w)zCheck if NVIDIA Universal GEMM (cutlass_api) is importable; cache the result for reuse.

    Call ensure_nv_universal_gemm_available.cache_clear() after installing cutlass_api
    in the same interpreter to retry the import.
    cutlass_apiNF)r  r  r  r%  _ensure_fp4_dtype_registered)	availables    rY   "ensure_nv_universal_gemm_availabler  
	  sE    NN,,];4G	 $&	  s   !1 	==c                     ddl } 	 | j                  j                  t        j                         y# t
        t        f$ r4 ddl| j                  j                  fd}|| j                  _        Y yw xY w)a$  Patch cutlass_api to handle torch.float4_e2m1fn_x2 -> cutlass.Float4E2M1FN.

    NOTE: cutlass_api doesn't natively map this dtype. We patch the lookup function
    in-place so all callers (including TensorWrapper) pick up the change.
    Remove once cutlass_api adds native FP4 support.
    r   Nc                R    | t         j                  k(  rj                  S  |       S r   )rR   rj  Float4E2M1FN)r   _origr  s    rY   _patchedz._ensure_fp4_dtype_registered.<locals>._patched*	  s'    ...+++<r   )cutlass_api.utilsr  cutlass_type_from_torch_typerR   rj  KeyErrorr
  r  )r  r  r  r  s     @@rY   r  r  	  sh     B66u7M7MNn% 
B!!>>	 
 :B6
Bs   )2 A A54A5c                 d    	 t         j                  j                  d      duS # t        $ r Y yw xY w)aG  Check if nvMatmulHeuristics is importable; cache the result for reuse.

    nvMatmulHeuristics provides performance model-based kernel selection
    for NVIDIA GEMM operations.

    Call ensure_nvmatmul_heuristics_available.cache_clear() after installing
    nvMatmulHeuristics in the same interpreter to retry the import.
    nvMatmulHeuristicsNFr  r   r   rY   $ensure_nvmatmul_heuristics_availabler  2	  s4    ~~''(<=TII r  c                n   t               syt        d      syddlm} t	        |j
                  j                        sy |       syt        j                  g}	t        ||	      syt        j                  st        j                  syt        | ||      syt        d | |fD              ry|r|ry|y||yy)a  
    Returns True if we can use the blackwell kernel for grouped mm.
    Required conditions:
        1. CuTeDSL backend is enabled
        2. CuTeDSL is available
        3. We are on a blackwell arch
        4. The dtype is bf16
        5. Max autotune or max autotune gemm is enabled
        6. A, B, and the output are 16B aligned
        7. We are not using dynamic shapes
        8. A is 2d
        9. B is 3d
        10. Offsets are provided
        11. Bias and Scale are not provided
    FCUTEDSLr9   r  )rc  c              3  2   K   | ]  }t        |        y wr   )
is_dynamicr   rV   s     rY   r   z3use_blackwell_cutedsl_grouped_mm.<locals>.<genexpr>t	  s     
1Q:a=
1r  T)r  rJ  r  r  rB  r   r,  rR   rl  rD  rl   r]  r^  r  r  )
mat_amat_brO  a_is_2db_is_2doffsbiasscale_resultr  r`  s
             rY    use_blackwell_cutedsl_grouped_mmr  B	  s    2 !" +C&--$$%')^^$M 76#;#; ue6:

15%.
11g|<3r   c                <   ddl m} |j                  j                  j	                  ||z  |z  d      }|dk  s|t
        j                  j                  k  ryddlm	} t        j                  j                  ryt        j                  t        j                  t        j                  g}t!        | |      xr/ t
        j"                  xs t
        j$                  xr t'        d      }|r6 |       s/t(        j+                  d	t
        j                  j,                         y|S )
Nr9   r  r   r?  r   F)try_import_cutlassCUTLASSzFailed to import CUTLASS lib. Please check whether _inductor.config.cutlass.cutlass_dir %s is set correctly. Skipping CUTLASS backend for now.)r  r  r  r  rB  rl   r  cutlass_backend_min_gemm_sizecodegen.cutlass.utilsr  rR   ro  ru   r   rl  r{  rD  r]  r^  rJ  r   r  cutlass_dir)	rO  r  r=  kr  	gemm_sizer  r`  r   s	            rY   use_cutlass_templater  	  s      221q519r2JIA~V^^%Q%QQ9 }} ]]ENNEKK@Mfm4 	-  <F$<$<	-!),  !#KK4 **	 Jr   _IntLikec                  
 ddl m t               syt               syt	        d      syddlm
 
j                  ry| j                  j                  dk7  st        j                  j                  ryt        j                  st        j                  sy|||g}||j!                  |       t#        fd|D              ry||g}	||	j!                  |       t#        
fd	|	D              ryy
)a3  
    Return True if we can use the NVIDIA Universal GEMM Template.

    Required conditions:
        1. NVGEMM backend is enabled
        2. cutlass_api is available
        3. We are on a NVIDIA GPU
        4. Max autotune or max autotune gemm is enabled
        5. Not in AOT Inductor mode (requires runtime JIT compilation)
        6. Base pointers are 16-byte aligned
        7. Shape dimensions are not unbacked symbols

    Note:
        - Shape and stride constraints are handled internally by
          cutlass_api.get_kernels() which filters incompatible kernels.
        - GroupedGemm currently only supports TN layout (column-major B).
          Any other layout will act as a noop and fall back to ATen.
        - Dynamic shapes are supported as long as they have hints
          (from example inputs).
    r   )has_free_unbacked_symbolsFNVGEMMr9   r  rJ   c              3  .   K   | ]  } |        y wr   r   )r   dimr  s     rY   r   z1use_nv_universal_gemm_template.<locals>.<genexpr>	  s     
Cc$S)
Cr  c              3  j   K   | ]*  }|j                         j                  j                  v  , y wr   )r  r  rt  )r   tr  s     rY   r   z1use_nv_universal_gemm_template.<locals>.<genexpr>	  s&     
O1::<177444
Os   03T)r  r  r  r  rJ  r  r  aot_compilationr   r,  rR   ro  ru   rl   r]  r^  r  r  )rO  r  r=  r  r  r  r  r  dims_to_checktensors_to_checkr  r  s             @@rY   use_nv_universal_gemm_templater  	  s    < P "-/ *}}V#u}}'8'86#;#;
 1IM}Q

C]
CC u~%

O>N
OOr   c                    t         j                  j                  j                         }|dk(  ry| j                         |j	                  d      D cg c]  }|j                          c}v S c c}w )z8Check if CUTLASS should be used for the given operation.ALLTrd  )rl   r  cutlass_enabled_opsr  r  r  )op_nameenabled_opsrV   s      rY   _use_cutlass_for_opr  	  sU    ..44::<Ke==?+2C2CC2HIQqwwyIIIIs   A,r   c           
        ddl m} t        j                  j                  |z  }|j
                  j                  j                  t        j                  t        j                  ||| z        t        j                  |||z                    xrO |j
                  j                   xr6 |j
                  j                   xr t        j                  j                  dkD  S )Nr   r  )torch._inductor.virtualizedr  rl   r  decompose_k_thresholdr  r  r  r   r  Ger  cpp_wrappernum_decompose_k_splits)r  r=  r  threshold_multipler  r  s         rY   use_decompose_k_choicer  	  s     ."MM??BTT 	
..II1A561A56	
 	5    	5 ###	5 MM0014
r   c           
        t         j                  j                  }ddlm} t        t        j                  j                        xr |j                  j                  j                  t        j                  t        j                  ||| z        t        j                  |||z                    xr0 |j                  j                   xr |j                  j                    S )z
    Check if we should use the contiguous subgraph transform.
    This transform makes the second matrix contiguous before the matmul.
    r   r  )rl   rocmcontiguous_thresholdr  r  rX  rR   ro  ru   r  r  r  r   r  r  r  r  )r  r=  r  r  r  s        rY   use_contiguousr  
  s     ";;;; . 	U]] 	$GG22II01450145
	$    	$ ###
r   c                   t         j                  j                  }g d}t        |t        j
                        r|j                  s|S |dk(  rg S t        | t        j
                        r| j                  r&t        |t        j
                        r|j                  sd}nt        || z  ||z        }d}t	        j                  |      }|D cg c]  }||k  r||k\  r| }}g g g }}
}	|D ]Z  }||z  }|dk  r||dz
  z  dk(  r|dk\  r|	j                  |       0|dz  dk(  r|
j                  |       J|j                  |       \ t         j                  dk(  r|	|
z   |z   S |	|
z   |z   }|d | S c c}w )	N)rv   rz  ry   rw      r   r   r  rw   r9   rz  
EXHAUSTIVE)rl   r  r  r   r   rR  	is_numberr  divisorsr  max_autotune_gemm_search_space)r  r=  r  k_splits_limitdefault_k_splitsmax_k_splitmin_k_splitr  divisorpow_of_2_divisorsmul_of_32_divisorsrest_of_splitsdkPartbest_splitss                  rY   get_k_splitsr  $
  s    ]]99N .!UZZ 	1		1ejj!!++1ejj!!++!q&!q&)K~~a H  k!g&< 	H  =?B>) %Q 3; EAI!#$$Q'RZ1_%%a( !!!$%" ,,< #55FF#&88>IK''=s   
E,c                T    t         j                  j                  |       j                  S r   )rR   rJ   r$  r>  r   s    rY   _rocm_native_device_arch_namer  ]
  s    ::++F3???r   c                     	 dd l } ddlm}m} ddlm} t        j                  j                  | j                        }||||fS # t        $ r dd}dd} G d d      }d }Y %w xY w)	Nr   )gen_ops_librarygen_ops_preselected)CKGemmOperationc                     g S r   r   r   r   rY   r  z*try_import_ck_lib.<locals>.gen_ops_librarys
      Ir   c                     g S r   r   r   r   rY   r  z.try_import_ck_lib.<locals>.gen_ops_preselectedv
  r  r   c                      e Zd Zy)*try_import_ck_lib.<locals>.CKGemmOperationN)r   r   r   r   r   rY   r  r  y
  s    r   r  )r   r  )ck4inductor(ck4inductor.universal_gemm.gen_instancesr  r  ck4inductor.universal_gemm.opr  r  r  dirname__file__r%  )r  r  r  r  package_dirnames        rY   try_import_ck_libr#  b
  sl    	
	
 ''//+*>*>? O-@/QQ  			 	 s   ;A A#"A#c                    t         j                  st         j                  syt        j                  j
                  sy| j                  j                  dk7  ryt        | j                        }t         j                  j                  D ci c]  }|j                  d      d   | c}xs |j                  d      d   |i}|j                         t         j                  j                  z  D cg c]  }||   	 }}|sy| j                  t        j                  t        j                   t        j"                  fvryt%               \  }}}}|st&        j)                  d       y|t         j                  _        yc c}w c c}w )NFrJ   :r   z,Please pip install Composable Kernel packageT)rl   r]  r^  rR   ro  ru   r   r,  r  r  archr  rk  ck_supported_archr   r   rl  rn  r#  r   r  ck_dir)rO  native_archr  requested_archsrequested_supported_archsck_package_dirnamer   s          rY   use_ck_templater-  
  s<   6#;#;==}}V# 0>K39;;3C3CDaqwws|A)D #q!;IO
 !%%'&++*G*GG! 	! ! %||EMM5>>5==II"3"51aBC+FKK+ E!s   E6,E;c                    ddl m} t        d      xr= t        |       xr0 |j                  j
                  j                  ||z  |z  d      dkD  S )Nr9   r  CKr   r  r   r  r  rJ  r-  r  r  rB  rO  r  r=  r  r  s        rY   use_ck_gemm_templater2  
  sR     	d# 	KF#	KGG..q1uqy2.FJr   c                    ddl m} t        d      xr= t        |       xr0 |j                  j
                  j                  ||z  |z  d      dkD  S )Nr9   r  CKTILEr   r  r   r0  r1  s        rY   use_ck_tile_gemm_templater5  
  sR     	h' 	KF#	KGG..q1uqy2.FJr   c                2    t        d      xr t        |       S )Nr/  )rO  r-  rO  s    rY   use_ck_conv_templater8  
  s    %d+G0GGr   c                |    t         j                  xs t         j                  xr | j                  j                  dk(  S r  )rl   r]  r^  r   r,  r7  s    rY   _use_template_for_cpur:  
  s2    7v77&
--


%&r   c                   ddl m} t        |j                  |      sJ |j                  j                  }|j                  j
                  }t        |       xrX |j                         t        j                  k(  xr5 t        |      dk(  xr% t        |      dk(  xr |d   |d   k(  xr |d   dk(  }t        | ||d      xr |j                  j                         xs |S )Nr9   )rC      r  F)require_constant_mat2)r!  rC   r   rO  rH  rW  r:  rs  rR   rn  rT   use_cpp_gemm_templateis_contiguous)rO  mat1mat2rC   	mat1_sizemat1_stridemat1_each_batch_is_contiguouss          rY   use_cpp_bmm_templaterE  
  s     dkk6***
   I++$$Kf% 	"NN-	"^q 	" "	" ^y|+		"
 ^q  " !t5Q !!#D'Dr   c                ~   ddl m} ddlm} ddlm}	 ddlm}
 t        |       rt        d      syt        j                  j                  sy|j                         t        j                  t        j                   fv }t        j"                  t        j$                  t        j&                  t        j                  t        j                   g} |
|||r| j(                  nd ||      \  }}}} }}t+        ||f      ryt-        ||j.                        r|j1                         } |	|j                               \  }} |d	||||j                         |j                         |t3               | |

      }dd}| j(                  |v xr= |d uxr7  ||      xr- t-        ||j4                        xr |j7                         xs | S )Nr9   r   )create_micro_gemm)*get_gemm_template_output_and_compute_dtype)mm_argsCPPF)	out_dtypemat2_transposeduse_4x2_dim
micro_gemm)input_dtypeinput2_dtypeoutput_dtypenum_threadsuse_refq_group_sizec                N    | j                          | j                         d   dk(  S )Nr   r9   )freeze_layoutrr  rV   s    rY   is_last_dim_stride1z2use_cpp_gemm_template.<locals>.is_last_dim_stride1  s"    	||~b!Q&&r   )rV   rB   r   rX  )r"  r!  codegen.cpp_micro_gemmrG  codegen.cpp_utilsrH  kernel.mm_commonrI  r:  rJ  rl   cppweight_prepackrs  rR   r}  rr  rn  rl  halfr   has_free_symbolsr   BaseViewunwrap_viewparallel_num_threadsrJ  is_module_buffer)rO  r@  rA  rL  r=  is_woq_int4rT  r!  rG  rH  rI  	int8_gemmr`  r  r=  r  rQ  r   rN  rX  s                       rY   r>  r>  
  s    9M) (0Ee0L::$$ U[[%**$==I]]ENNEJJUZZXM")"+&,,'#Aq!VT4 A$$!@AQROL!"			NN$^^%!(*!J'
 	% 	Cd"	C%	C tR]]+	C ""$A,A(Ar   c                 b    t         j                  xs t         j                   xs t        d      S )NATEN)rl   r]  r^  rJ  r   r   rY   use_aten_gemm_kernelsrh    s-    7v77 '	v	&'r   c                  T    e Zd ZU  ej                  d      Zded<   ddZddZd	dZ	y)
DebugDirManagerr   r&  prev_debug_namec                @    t        t        j                        | _        y r   )r  rj  counterr   rs  s    rY   rd  zDebugDirManager.__init__&  s    ../r   c                    t         j                  j                  j                  | _        | j                   d| j
                   | _        | j                  t         j                  j                  _        y )N_tmp_)rR   _dynamorl   debug_dir_rootrk  r   new_namers  s    rY   	__enter__zDebugDirManager.__enter__)  sM    $}}33BB//0dggY?.2mm+r   c                    t        j                  | j                         | j                  t        j
                  j                  _        y r   )r)  r*  rr  rk  rR   rp  rl   rq  )r  r   s     rY   __exit__zDebugDirManager.__exit__.  s*    dmm$.2.B.B+r   Nr  )r   r  r   r  )
r   r   r   r  rl  rm  r   rd  rs  ru  r   r   rY   rj  rj  "  s(    iooa G0<
Cr   rj  c                   ddl m} t               dfd}t        j                  j                  |d|      5  t        j                  j                           | |i |}d d d        t              fS # 1 sw Y   xY w)Nr9   r>   c                (    j                  |        y r   )r   codesource_codess    rY   save_output_codez*run_and_get_code.<locals>.save_output_code<  s    r   r{  ry  r&  r   r  )
r  r?   r!   r   patchr  rR   rp  resetr  )r   r   r  r?   r{  r  rz  s         @rY   run_and_get_coder  3  st    
 %$.LL 
		=*<>N	O %T$V$% 4%%%% %s   'A55A>c                   |j                  dd      }t        | g|i |\  }}g }|D ]  }t        j                  rQt        j                  j
                  dur5|j                  t        j                  d|t        j                               n4|j                  t        j                  d|t        j                               |s|D cg c]  }|dd 	 }} ||fS c c}w )Nremove_quoteFTzR"TRITON\((.*?)\)TRITON"z	'''.*?'''r<  )
rU   r  rl   r  r  autotune_at_compile_timerp  r   findallDOTALL)	r   r   r  r  r  rz  kernelsry  r  s	            rY   run_and_get_kernelsr  E  s     ::ne4L+B@@@FLG ;&--"H"HPT"T NN2::&A4STNN2::lD"))DE29:va|:G:; 7? ;s   Cc                &     d fd}t        |      S )Nc                 R            } | j                         j                          | S r   )r?  r0  )r  r   s    rY   run_with_backwardz1run_fw_bw_and_get_code.<locals>.run_with_backwardY  s!    

r   )r   r   )r  )r   r  s   ` rY   run_fw_bw_and_get_coder  X  s    
 -..r   c                X   ddl m} g dfdd	fd}t        j                  j	                  |d|      5  t        j                  j	                  |d      5  t
        j                  j                           | |i |}ddd       ddd       S # 1 sw Y   xY w# 1 sw Y   S xY w)
zLGet the inductor-generated code, but skip any actual compilation or running.r9   r>   c                (    j                  |        y r   r  rx  s    rY   r{  z"get_code.<locals>.save_output_codeg  s    D!r   c                     G d d      }| j                   r| j                         n| j                         \  }} |j                         |r |j                          |       S )Nc                       e Zd ZdZddZddZy)@get_code.<locals>.patched_compile_to_module.<locals>.DummyModulez4This is empty to replace the generated triton modulec                     y r   r   rs  s    rY   rd  zIget_code.<locals>.patched_compile_to_module.<locals>.DummyModule.__init__n  s    r   c                     y r   r   r   s      rY   callzEget_code.<locals>.patched_compile_to_module.<locals>.DummyModule.callq  s    r   Nr  r   r   r  r   r   r  )r   r   r   r   rd  r  r   r   rY   DummyModuler  k  s    Fr   r  )r  codegen_with_cpp_wrappercodegenr   )r  r  wrapper_codekernel_coder{  s       rY   patched_compile_to_modulez+get_code.<locals>.patched_compile_to_modulej  s]    	 	 04/?/?D))+T\\^ 	"k 	++,[../}r   compile_to_moduler{  Nr|  )r  r?   r   r   )r  r?   r   r}  r  rR   rp  r~  )r   r   r  r?   r  r   r{  rz  s         @@rY   get_coder  a  s    $ L", 	

.0I	
  	

-);=MN	  	          s#   "B'BBB	BB)c                |    t        | g|i |}dt        |      cxk  rdk  sn J dt        |              |d   S Nr9   r  z%expected one or two code outputs got r   )r  rT   )r   r   r  rz  s       rY   get_triton_coder    sQ    B000LL!&Q& 
/L0A/BC& ?r   c                    t        | g|i |\  }}dt        |      cxk  rdk  sn J dt        |              |d   S r  )r  rT   )r   r   r  r   rz  s        rY   run_and_get_triton_coder    sW     'r;D;F;OA|L!&Q& 
/L0A/BC& ?r   c                    ddl m ddlm} |j                  g dfd}t
        j                  j                  |d|      5   | |i |}d d d        |fS # 1 sw Y   fS xY w)Nr   r>   rF   c                 ^     | i | | d   }t        |      sJ j                  |       y )Nr  )r   r  )r   r  r  r?   graph_lowerings	real_inits      rY   	fake_initz-run_and_get_graph_lowering.<locals>.fake_init  s7    4"6"Q%///u%r   rd  r  )torch._inductor.graphr?   torch._inductor.output_coderG   rd  r   r}  r  )	r   r   r  rG   r  r  r?   r  r  s	         @@@rY   run_and_get_graph_loweringr    sq     4;((IO& 
		?J		B %T$V$% ?""% ?""s   	AA(c              #     K   ddl m} |j                  |    }	 t        j                  ||      |j                  | <   d ||j                  | <   y# ||j                  | <   w xY ww)z
    Override the lowering of aten_op with override_fn.
    The first argument of override_fn is the original lowering fn.
    r   )loweringN)torch._inductorr  	loweringsr2  partial)aten_opoverride_fnr  orig_fns       rY   override_loweringr    s`      )  )G.&/&7&7W&M7#&-7#g7#s   A$'A  A$A!!A$c                     ddl m} |j                  d fd}t        j                  j
                  j                  |d|      S )zr
    Add hook functions to be called at the beginning and end of Scheduler.__init__.
    Used for unit tests.
    r   )	Schedulerc                B     | |        | |      }r	 | |       |S r   r   )r  rh  outr  post_fnpre_fns      rY   r  z(add_scheduler_init_hook.<locals>.wrapper  s+    y% i'Iu%
r   rd  )r  r   rh  r   r   r   )torch._inductor.schedulerr  rd  unittestr   r}  r  )r  r  r  r  r  s   ``  @rY   add_scheduler_init_hookr    s9     4  G ==%%iWEEr   c                z    t         j                  rt        j                  |        yt        j	                  |        y)z
    Warnings that will be actionable for PyTorch developers, but not
    end users.  Allows us to easily disable them in stable releases but
    keep them on for nightly builds.
    N)rl   developer_warningsr   r  info)msgs    rY   developer_warningr    s$       Cr   c                    	 t         j                  j                  d      } | dz   t        t         j                        k  rTt        t         j                  | dz            dkD  r2t         j                  | dz      d   dk7  rt         j                  | dz      S t         j                  D ]#  }|j                  d      s|t        d      d c S  y# t        $ r Y Bw xY w)a  
    An experimental API used only when config.benchmark_kernel is true.

    The benchmark name is only available at codegen time. So we can not
    directly call it in benchmark_all_kernels which is run after codegen.

    The function assumes the argument after --only is the benchmark name.
    It works for torchbench.py/hugginface.py/timm_models.py. But for ad-hoc
    scripts, this function may return None.

    There are 2 flavors of --only argument we need to handle:
    1. --only model_name
    2. --only=model_name
    z--onlyr9   r   r  z--only=N)r  argvr+  rT   
ValueErrorr  )rv  r  s     rY   get_benchmark_namer    s    	hhnnX&!Gc#((m#CHHS1W%&*q!!$+88C!G$$ xx )>>)$s9~'(()   s   BC 	CCc                &    t        d | D              S )Nc              3  &   K   | ]	  }|d k(    ywr9   Nr   r  s     rY   r   zis_ones.<locals>.<genexpr>       %!qAv%   r   rl  s    rY   is_onesr        %u%%%r   c                &    t        d | D              S )Nc              3  &   K   | ]	  }|d k(    yw)r   Nr   r  s     rY   r   zis_zeros.<locals>.<genexpr>  r  r  r  r  s    rY   is_zerosr    r  r   c                &    t        d | D              S )Nc              3     K   | ]@  }t        |t        j                        r$|j                  t        j                  d       k(   B yw)r)  N)r   rR   r  r   )r   r   s     rY   r   z is_cpu_device.<locals>.<genexpr>  s8      dELL) 	u||E**s   AAr  )inputss    rY   is_cpu_devicer    s       r   c                    t        | t        j                        sJ d       | j                  rt        j
                  S t        j                  S )Nz8only support sympy.Expr as input to get_sympy_Expr_dtype)r   r   rR  r   rR   r|  rp  r  s    rY   get_sympy_Expr_dtyper    s=    c5::& B& ~~{{}}r   c              /     K   | r-t        j                  j                  |i |5 }| d d d        y d  y # 1 sw Y   y xY wwr   )rR   r   r   )should_profiler   r  r   s       rY   maybe_profiler  %  sE     ^^##T4V4 	G	 	 		 	s   "A7AA Ac                 l    t         j                  j                  } | dk  rt        j                         } | S Nr9   )rl   r\  threadsrR   get_num_threads)r  s    rY   rb  rb  .  s+    jj  G{'')Nr   c                     ddl m}   |        }|j                  dt        j                  j
                  rd      S d      S )Nr9   )get_backend_options
num_stagesr  r<  )runtime.triton_helpersr  rj  rR   ro  ru   )r  optionss     rY   get_backend_num_stagesr  5  s2    ;!#G;;|%--*;*;QCCCCr   c                   t        | t        j                  j                  j                  j
                  dk(        }||S t        j                  j                         st        j                  d       yddl	m
}m} t        j                  j                         xr! t        j                  j                         dk\  }| t        j                  t        j                  t        j                  fv sJ t!        j"                  |      j$                  j'                  d      rdd	lm}  |       }| t        j                  t        j                  fv r|r	 || |      S t        j                  j                  j                  j
                  dk(  r |t        j                  |      S  |t        j                  |      S | t        j                  t        j                  fv r
|r ||       S t        j                  j                  j                  j
                  dk(  r |t        j                        S  |t        j                        S )
z
    We don't want to throw errors in this function. First check to see if the device is in device_info.py,
    then fall back to the inaccurate triton estimation.
    tf32)is_tf32zget_device_tflops: no Triton fallback available for non-CUDA devices. Returning 0.0; roofline estimates will use memory bandwidth only.r   r   )get_max_simd_tflopsget_max_tensorcore_tflops)rz   r   
clock_rate)max_clock_rate)r   rR   backendsrJ   matmulfp32_precisionrS   r   r  triton.testingr  r  r+  r   rl  rn  inspect	signature
parametersrj  torch._utils_internalr  )r   ds_topsr  r  SM80OrLaterr  sm_clocks          rY   get_device_tflopsr  =  s    u~~**11@@FJG ::""$P	
 M**))+ 

0P0P0R W 1K
 U]]ENNEMMBBBB,-88<<\J8!#U]]ENN33,UH==>>%%44>,U]]HEE&u}}h??U]]ENN33,U33>>%%44>,U]];;&u}}55r   c                     ddl m}   |        S )Nr   get_dram_gbps)r  r  r  s    rY   get_gpu_dram_gbpsr  o  s    ,?r   c                 x    ddl m}  | j                  j                  j	                  d      j                  dd      S )Nr   r  max_shared_mem)triton.runtimer  r  r  r$  rj  r  s    rY   get_gpu_shared_memoryr   v  s.    %==44Q7;;<LaPPr   c                     t         j                  j                         rht        j                  dt         j                  j	                               } t        j                  |       }|j                  }|j                  }|J ||z  S d}d}||z  S )NrJ   rz  i   )	rR   rJ   rS   r   r.  r   r  warp_size_or_defaultmax_threads_per_block)r   props	warp_sizer  s       rY   get_max_numwarpsr  |  s    zz fejj&?&?&AB ''/..	 % ; ;$000
 !I-- 	 $ I--r   c                $    | j                  d      S )Nwelford)r  reduction_types    rY   is_welford_reductionr    s    $$Y//r   c                &    t        |       ry| dv ryy)Nr<  )argmax_with_valueargmin_with_valueonline_softmax_reducer  r9   )r  r	  s    rY   reduction_num_outputsr    s"    N+	  

 r   c                 0    t        j                         dk(  S )NLinux)platformsystemr   r   rY   is_linuxr    s    ??''r   c                 (    t         j                  dk(  S )Nrn   )r  r  r   r   rY   r(  r(    s    <<7""r   c                &    t        d | D              S )Nc              3  n   K   | ]-  }t        |t        j                        xr |j                    / y wr   )r   r   rR  r  r  s     rY   r   z#has_free_symbols.<locals>.<genexpr>  s)     Jz!UZZ(<_<Js   35r  )itrs    rY   r_  r_    s    JcJJJr   c            	     x   ddl m} | D ]  }t        ||j                  |j                  |j
                  |j                  |j                  f      r=t        |j                         xs d      st        |j                         xs d      s yt        ||j                        st        dt        |              y)Nr9   r   r   Tzunexpected type for is_dynamic F)r"  r!  r   rH  rJ  r`  r  r@   r_  maybe_get_sizemaybe_get_striderB   	TypeErrorr,  )r   r!  r  s      rY   r  r    s     IbmmR[[":K:KRYYW
   0 0 2 8b9=M""$*> Aryy)=d1gYGHHI r   c                      e Zd ZdZdZy)PlaceholderKERNEL_NAMEDESCRIPTIVE_NAMEN)r   r   r   r   r!  r   r   rY   r  r    s      K *r   r  c                ~   ddl m} t        j                  dd      5 }t	        j
                         }t	        j
                         } t        |t        |            j                  |  t        d|j                   |       t        |j                  |       t        j                         }t        ||      5   | |j                         d d d        t        j                         |z
  }	 ||j                         |j                  j                          |j                          t        d	|j                   |       t        |j                  |       |j!                         |j!                         k(  }
t"        j%                  d
||j&                  |
|	       d d d        y # 1 sw Y   xY w# 1 sw Y   y xY w)Nr9   )stable_topological_sortr  zutf-8)modeencoding)r  	fake_modezBefore:
)filezAfter:
zZ%s, save before/after graph to %s, graph before/after are the same = %s, time elapsed = %s)pattern_matcherr#  r"  NamedTemporaryFileior   r`   r\   	propagater  r  r   nowr_   lint	recompilerm  r   r  r   )r  r  inpr  r#  r0  	before_ioafter_io
start_timetime_elapsedr  s              rY   pass_execution_and_saver4    sU    9		$	$
 
 
KKM	;;=C	R#3C#89CCSI	"(($1-bhhY'\\^
#B, 	N	||~
2)


#!,bhhX& H$5$5$77hFF	
+
 
	 	
 
s%   BF3;F'CF3'F0	,F33F<c                ^    ddl m} t        | |j                        xr | j	                         S )zB
    Check if input buffer is a multi-outputs template buffer
    r9   r   )r"  r!  r   r  is_multi_outputs_template	input_bufr!  s     rY   r6  r6    s/      	9b//0 	2//1r   c                    ddl m} t        | |j                        xr2 t	        | j
                        dk(  xr t        | j
                  d         S )zL
    Check if input buffer is a output of multi-outputs template buffer
    r9   r   r   )r"  r!  r   MultiOutputrT   r  r6  r7  s     rY   #is_output_of_multi_outputs_templater;    sL      	9bnn- 	;	  !Q&	;%i&6&6q&9:r   c                &   | yddl m} t        | |j                        xr- t        | |j                         xr |d u xs | j
                  |u xsA t        |       |j                  u xr' t        t        j                  j                  d      xr; | j
                  t        j                  j                  j                  j                  k(  xs t        t        j                  j                  d      xr; | j
                  t        j                  j                  j                  j                  k(  xsa t        t        j                  j                  d      xr; | j
                  t        j                  j                  j                  j                  k(  S )NFr9   r   all_to_all_singleall_gather_into_tensorreduce_scatter_tensor)r"  r!  r   _CollectiveKernel_WaitKernelop_overloadr,  FallbackKernelr#  rR   r$  torchrecr=  defaultr>  r?  r  r  r!  s      rY   is_collectiverG    sG    | 	4--. 	3400	34Z14++r1 	T
b''' 	

 		**,?@ U$$		(:(:(L(L(T(TT
 		**,DE E$$99%%<<DDE 		**,CD Y$$		(:(:(P(P(X(XX/r   c                <    ddl m} t        |       |j                  u S Nr9   r   )r"  r!  r,  rA  )r  r!  s     rY   is_waitrJ  )  s    :''r   c                    ddl m} t        | |      rt        d | j                  D              S t        | j                        xr |d u xs  ||       S )Nr   GroupedSchedulerNodec              3  2   K   | ]  }t        |        y wr   )contains_collectiver  s     rY   r   z&contains_collective.<locals>.<genexpr>6  s     @a&q)@r  )r  rM  r   r  snodesrG  r  )snode	filter_fnrM  s      rY   rO  rO  /  sJ     ?%-.@5<<@@@$P)t*;*Oy?OPr   c                    ddl m} t        | |      rt        d | j                  D              S t        | j                        S )Nr   rL  c              3  2   K   | ]  }t        |        y wr   )contains_waitr  s     rY   r   z contains_wait.<locals>.<genexpr>?  s     :=#:r  )r  rM  r   r  rP  rJ  r  )rQ  rM  s     rY   rU  rU  ;  s4    >%-.:U\\:::uzz""r   c                    ddl m} t        |t        j                  j
                        r|g}t        | |j                        xr | j                  |v S rI  )r"  r!  r   rR   r  r  rC  rB  rF  s      rY   is_fallback_oprW  D  sE     "ejj++,TdB--.I43C3Cr3IIr   c                B    |||    j                   j                            S r   )defining_opr  )buf_namename_to_bufname_to_fused_nodes      rY   buf_name_to_fused_snoder]  O  s#     k(3??HHJKKr   c                     yr  r   rQ  s    rY   r  r  Z  r  r   c                     ||       ry |j                  |        | j                  D ].  }t        |j                  ||      }||v rt	        |||||       0 y )Ncriteria_cb)r   unmet_dependenciesr]  r   find_recursive_deps_of_node)rQ  collected_node_setr[  r\  rb  depdefining_op_for_deps          rY   rd  rd  U  sn     55!'' 
5HHk#5
 "44##	

r   c                     yr  r   r_  s    rY   r  r  s  r  r   c           	     z    ||       ry |j                  |        | j                         D ]  }|j                  D ]}  }|j                  J |j                  j	                         dk(  r/|j                  j	                         |vrL||j                  j	                            }||v rnt        |||||         y )NOUTPUTra  )r   get_outputsr  r  r  find_recursive_users_of_node)rQ  re  r[  r\  rb  or  user_ops           rY   rl  rl  n  s     55!  GG 	D99(((yy!!#x/yy!!#+==(););)=>G,,(""'	r   c                b    t         j                  j                  j                  rdnd}|| z
  |z
  S )zaComputes the number of inputs to the aot fw graph which have fixed addresses (params and buffers)r  r   )rR   
_functorchrl   functionalize_rng_ops)dynamo_gm_num_inputsaot_fw_gm_num_inputsnum_rng_seed_offset_inputss      rY   num_fw_fixed_argumentsru    s6     $$::   "669SSSr   c                    dd}d}g }| j                   j                  D ]0  }|j                  dk(  s ||      r|j                  |       |dz  }2 |t	        t        t        |                  k(  sJ t        |      S )z>
    Infers which inputs are static for a backwards graph
    c                ~    d| j                   vxr. d| j                   vxr d| j                   vxr d| j                   vS )Ntangentsbwd_seedbwd_base_offsetbwd_rng_stater  rW  s    rY   is_saved_tensorz'count_tangents.<locals>.is_saved_tensor  sH    aff$ .!&&(.!/.  qvv-		
r   r   r  r9   )rV   r7   r   rX  )r  rh  r  r  r  r   rT   )fx_gr|  	arg_countstatic_arg_idxsr=  s        rY   count_tangentsr    s    

 IOZZ 44= q!&&y1NI	 d5_)=#>????r   c                    g }t        | j                  j                        D ]D  \  }}|j                  dk7  r |S |j                  j                  d      s4|j                  |       F |S )a  
    Returns indices of backward graph inputs that are always at fixed
    addresses: primals (parameters/buffers/user inputs saved for backward).
    Excludes saved activations which may not be at fixed addresses when
    the forward is partitioned for CUDA graphs.
    r  primals_)rK  r  rh  r  r   r  r  )r}  static_idxsrv  r=  s       rY   get_static_bw_input_idxsr    sl     KDJJ,,- $Q44=   66Z(s#	$
 r   c                  2    e Zd ZU ded<   ddZedd       Zy)	BoxedBoolrX  r   c                    | j                   S r   )r   rs  s    rY   r~  zBoxedBool.__bool__  s    zzr   c                6    t        | t              r	d| _        | S yr  )r   r  r   r  s    rY   disablezBoxedBool.disable  s    c9%CIJr   Nr  )r  r   r   zBoxedBool | bool)r   r   r   r   r~  r  r  r   r   rY   r  r    s     K  r   r  c              #      K   ddl m} |j                  	 	 	 d	 	 	 	 	 	 	 	 	 	 	 	 	 d fd}t        j                  j                  |d|      5  d  d d d        y # 1 sw Y   y xY ww)Nr9   r;   c                @    j                  |        | |||||      S r   r  )r  kernel_namer  ry  gpucpp_definitionkernel_listorig_define_kernels         rY   define_kernelz.collect_defined_kernels.<locals>.define_kernel  s-     	;'!+{Hc>
 	
r   r  )NTN)r  r<   r  r&  r  r&  ry  r  r  rX  r  r  r   r   )codegen.wrapperr<   r  r   r}  r  )r  r<   r  r  s   `  @rY   collect_defined_kernelsr    s     5-;;  $%)
"

 
 	

 
 #
 

 
		/-	P   s   AA*A	A*A'#A*c                    | dz   S )N__original__r   r  s    rY    get_cloned_parameter_buffer_namer    s    .  r   c                    | t         v S r   )rP   r  s    rY   rB  rB    s    Yr   c                 :    t         j                  j                  duS )z,Check if we're running on ROCm/HIP platform.N)rR   ro  ru   r   r   rY   is_rocmr    s    ==D((r   c                    t               sy 	 ddlm} ddlm}  || ||      } ||      j                  i       }t        |dd       }|y t        |      S # t        t        t        t        f$ r Y y w xY w)Nr   )	GPUTarget)make_backendsupported_fp8_dtypes)r$   triton.backends.compilerr  triton.compiler.compilerr  parse_optionsr
  r%  r  r  rQ   r  )backend_namer&  r  r  r  r  r  r  s           rY   _triton_supported_fp8_dtypesr    s     69 <y9v&44R8 #7,BDI#%&& K9= s   -A A10A1)triton_backendtriton_archr  c                  | t         j                  t         j                  t         j                  t         j                  fvryt        |       j                  d      }||yt        j                  |      }|j                  dk7  ryt         j                  j                         sy|j                  }|t         j                  j                         }t         j                  j                  |      }t               r!d}|j                  j!                  d      d   }nd}|j"                  dz  |j$                  z   }t'        |dd       }|t)        d	      |	|dk(  rd
nd}t+        |||      }|y||v S )NTr  rJ   ru   r%  r   r  r  z0triton_arch must be provided with triton_backendry   rz  )rR   rb  rc  rd  rf  r  removeprefixr   r,  rJ   rS   r+  r.  r$  r  r>  r  r  minorrQ   r  r  )	r   r   r  r  r  triton_dtypedevice_index
propertiesr  s	            rY   is_triton_fp8_dtype_supportedr  	  sc    	  E?//4L>f%;;& zz&&(|| ::446LZZ55lC
9"N$0066s;A>K#N$**R/*2B2BBKJT:	KLL(E1Br	7Y #///r   c                &    | dk7  xr t        |       S )NrK   )rB  r  s    rY   device_need_guardr  =  s    U?-vf~-r   c                N   | t         j                  k(  r?t         j                  j                         r!t         j                  j	                         dk  S | t         j                  k(  rt         j
                  j                         ry| t         j                  t         j                  fv S )N)r  r   T)rR   rl  rJ   rS   r+  rL   r|  rX  rS  s    rY   ,needs_fallback_due_to_atomic_add_limitationsr  A  sk    5::#:#:#<zz//1F::	%..	 UYY%;%;%=ejj111r   c                   | j                   t        j                  j                  j                  t        j                  j                  j
                  fv r|y| j                   t        j                  j                  j                  k(  rdnd}|d |fvxs |xr t        |      xr t        |      xs | j                   t        j                  j                  j                  k(  xrW |dk(  xrP |xrL |dk(  xrE t        j                  j                  xr) t        j                  j                  xs t               dk7  xs? ||k(  xr" |t        j                  t        j                  fv xs t        j                          S )NFr   r?  r)  r9   )overloadpacketrR   r$  atenscatter_reduce_scatter_reducescatter_rB  r  rl   r\  fallback_scatter_reduce_sumdynamic_threadsrb  rX  r|  $are_deterministic_algorithms_enabled)rB  r
  
self_dtype	src_dtypesrc_device_typesrc_is_tensor	reduce_tys          rY   use_scatter_fallbackr  J  sZ    	""IINN**EIINN,I,IJ	K" ++uyy~~/F/FFE 
 	tY// 	8 H'H<YG		8 &&%))..*H*HH L%'LL  5(L 

66	L
 ++J/C/E/J	8 i'SJ5::u{{:S,S	8 557!r   c                   ddl m}m} ddlm} t        dt        |        d       t        |       D ]  \  }}t        d|dd       ||u rt        d	       '||u rt        d
       7t        ||      r|j                         }t        |rdnd d       |r:|j                  J t        d|j                  j                  j                          t        d       |j                  j                  D ]  }t        |        t        d       |j                  j                  D ]  }t        |        t!        dt#        |              y)z
    An API that can be used in pdb to dump a node_schedule.
    Right mainly dump the read/write dependencies but can add more as needed.
    r   )DisableReductionEnableReduction)SchedulerNodezNode schedule with z nodesr   3r%  zenable reductionzdisable reductionredpwz scheduler nodeNzoriginal reduction hint zReadDep:z	WriteDep:zUnrecognized node type: )torch._inductor.codegen.simdr  r  r  r  r  rT   rK  r   is_reductionr  rI  reduction_hintr\  r]  r^  r  r,  )r%  r  r  r  rv  r  is_redrf  s           rY   dump_node_scheduler  q  s=   
 O7	M 236
:;}- H	T#al?"$%%%%&m,&&(FfU$/?@yy,,,01N1N0OPQ*''-- c
+''.. c
 !9$t*FGG+Hr   c                z    ddl m}  || j                         t        | j                        z  t
        z  dk(        S )Nr   )r  )r  r  storage_offsetrU  r   GPU_ALIGN_BYTES)r   r  s     rY   tensor_is_alignedr    s:     L 				 >&,,#?	??RVWW r   c                |    t        | j                  j                        syt        j                  xs t        |       S r  )rB  r   r,  rl   assume_aligned_inputsr  )example_inputs    rY   should_assume_input_alignedr    s2     -&&++,''K+<]+KKr   c                 6   t         j                  j                  j                         } | st	        j
                         S | j                  r| j                  j                  st	        j
                         S | j                  j                  }|j                         S r   )	rR   _guardsTracingContexttry_getr  nullcontextr&  r  suppress_guards)tracing_contextr  s     rY   #maybe_get_suppress_shape_guards_ctxr    sv    
 mm22::<O%%'' $$O,E,E,O,O%%''))33I$$&&r   c                   t         j                  j                  j                  t        dd      5  t
        j                  j                          dd l}dd l	} |j                         } |j                  |      }ddlm} |j                  |       |j                  }|j!                  |j"                          | |i |}	|j%                         }
|j!                  |       |j'                  |       d d d        |	|
fS # 1 sw Y   	
fS xY w)Nr   Tr   )output_code_log)r  r   r}  r  rl   rR   rp  r~  r*  r  r   StreamHandlertorch._inductor.codecacher  
addHandlerlevelsetLevelr  rm  removeHandler)r   r   r  r*  r  log_capture_stringchr  
prev_levelr  r   s              rY   run_and_get_cpp_coder    s     
			#	#FGT	: *(R[[]"W""#56=""2&$**
  /T$V$'')  ,%%b)*  19!*  19s   CC>>D
c                   t        |       }||j                  S | D ]  }t        |t        j                        r|j
                  j                  c S t        |t        j                        sP|j                         D ]6  }t        |t        j                        s|j
                  j                  c c S  |j                         D ]6  }t        |t        j                        s|j
                  j                  c c S   y r   )	r\   r  r   rR   r4   r  r  rH  rW  )r  r&  inputrH  rW  s        rY   shape_env_from_inputsr    s     (I """  1eU\\*::''' eU\\*

 /dELL199.../  ,,. 1fell3!;;00011 r   c                <     t              dk(  r S d fd}|S )Nr   c                z    t        |       \  }} |       }t        |      rt        j                  ||       |S r   )copy_misaligned_inputsrT   rR   _foreach_copy_)
new_inputsold_tensorsnew_tensorsr  inputs_to_checkr  mutated_input_idxss       rY   r  z)align_inputs_from_check_idxs.<locals>.run  sE    #9);$
 [ J {  k:
r   )r  list[InputType]r   r   )rT   )r  r  r  r  s   ``` rY   align_inputs_from_check_idxsr    s#    
 ?q  Jr   c                T   d| j                         v rd}n;t        d t        | j                         | j                               D              dz   }t	        j
                  | |fd      j                         }t	        j
                  || j                         | j                               S )Nr   c              3  2   K   | ]  \  }}|d z
  |z    ywr  r   )r   rS  rW  s      rY   r   z)clone_preserve_strides.<locals>.<genexpr>  s     Tf$Tr  r9   r   )rH  r?  r   rW  rR   
as_stridedclone)rV   needed_sizerM  s      rY   clone_preserve_stridesr    s    AFFH} T#affh
:STTWXX 	 a+6<<>FFAFFHahhj99r   c                2   g }g }|du}|D ]  }| |   }t        |t        j                        sJ dt        |              |j	                         t
        z  sMt        |      | |<   |s^||v sc|j                  |       |j                  | |           ||fS )z
    Clones misaligned tensors which we inferred were aligned. Returns a tuple of [old_tensors], [new_tensors] for every
    cloned tensor which is in `return_pair_idxs`.
    Nz Expected tensors only, but got: )r   rR   r  r,  data_ptr	ALIGNMENTr  r  )r  check_inputs_idxsreturn_pair_idxsr  r  ret_pair_definedr   _inps           rY   r  r    s     ')K&(K (t3 
2!}$- 	
.tDzl;	
- ==?Y&248JqMA)9$9""4("":a=1
2 ##r   c                    g }|D ]N  }| |   }t        |t        j                        s#|j                         t        z  dk(  s>|j                  |       P t        |      t        |      k7  r|S |S )z[
    We require all inputs to be aligned, so introduce a copy for any
    that aren't.
    r   )r   rR   r  r  r  r  rT   )r  static_input_idxsaligned_static_input_idxsrv  r  s        rY   remove_unaligned_input_idxsr  1  st     !#  2seU\\*0@90LQR/R%,,S12 $%->)??((r   c                P   ddl m} t        j                  t        j                        j
                  }|j                  j                  j                  }|j                  j                  j                  j                  }t        j                  r'|j                  j                  j                  | |       y|j                  j                  j                  | |k        ry|j                  r)|j                  j                  j                  | dk        ry ||       xr  ||       |k  S )zCheck if an expression fits within 32-bit integer range.

    NOTE: This function intentionally does not install guards. Callers are
    responsible for guarding (e.g. via check_leq) when they decide to use
    32-bit indexing based on this result.
    r9   r  Tg@xDF)r  r  rR   r  r{  r   r  r  guarding_hint_or_throwr  has_guarding_hintrl   assume_32bit_indexing	check_leqr  r  )r   r  int_maxr  r  s        rY   expr_fits_within_32bitr  C  s     kk%++&**GWW--DD((22DD##	""1g. 	ww--a7l; 	 7711!d(;  QH$:1$=$HHr   c                   t         j                  j                  j                         }||j                  t        |j                        dk(  sJ t        |       |j                  J |j                  D ]  }||j                  j                  d        !dt         j                  j                  j                         x}r|j                  dfd|j                  j                  t        fd|D                      y y y )Nr   Fc                f    t        |       S rj                  |       S j                  |       S r   )r   deserialize_symexprevaluate_symexpr)r   fakify_first_callr  s    rY   map_exprz4set_tracing_context_output_strides.<locals>.map_expr  s7     ("1v((<<Q??$55a88r   c              3  .   K   | ]  } |        y wr   r   )r   r   r  s     rY   r   z5set_tracing_context_output_strides.<locals>.<genexpr>  s     5!(1+5r  )r   r   r   z)float | int | SymInt | SymFloat | SymBool)
rR   r  r  r  output_stridesrT   r  r  r  r  )r  compiled_graphrX  rP  r  r  r  r  s        @@@rY   "set_tracing_context_output_stridesr  p  s     mm**224Gw55A7))*a///).9	,,888#22 	E}&&--d3$)!--66>>@@3@(+(=(=%9 &&--5u55		  Br   c                    t         j                  t         j                  S t        j                         syt        j                  j                         ry	 ddlm}  | t        j                  j                  d      k\  S # t        $ r Y yw xY w)NFr   REMOTE_CACHE_VERSIONz.pytorch/remote_cache:fx_graph_memcache_version)
rl   fx_graph_remote_cache	is_fbcoderR   _utils_internalis_fb_unit_testtorch._inductor.fb.remote_cacher  ModuleNotFoundErrorjustknobs_getval_intr  s    rY    should_use_remote_fx_graph_cacher!    s    ##/+++,,.H  5#8#8#M#M8$    s   A> >	B
	B
c                0    t        j                  dd|       S )Nz[^a-zA-Z0-9_]r   )r   subr  s    rY   normalize_namer$    s    66"C..r   ztl.int1ztl.float8e4nvztl.float8e5ztl.float8e4b8ztl.float8e5b16ztl.uint8)ztl.boolztl.float8_e4m3fnztl.float8_e5m2ztl.float8_e4m3fnuzztl.float8_e5m2fnuzztl.float8_e8m0fnuztl.float4_e2m1fn_x2z^.*[.]c                l    t         j                  dt        |             }t        j	                  ||      S )z"Convert torch.dtype to triton typetl.)_triton_type_rer#  r&  _triton_type_mappingrj  )r   triton_type_names     rY   triton_typer*    s.    &**5#e*=##$46FGGr   c                    t         j                  | |       }|j                  dd      }t        t        |      }t        |t        j                        sJ |S )Nr&  r"  )_torch_triton_mappingrj  r  rQ   rR   r   r   )r   adjusted_type	type_namerK  s       rY   triton_type_to_torchr/    sL    )--eU;M%%eR0Iy)Ii---r   c                   | j                    xr | j                         |j                         k(  xr | j                         |j                         k(  xr | j                  |j                  k(  xr{ | j                  |j                  k(  xr` | j                         j                         |j                         j                         k(  xr! | j                         |j                         k(  S r   )	is_mkldnnrH  rW  r   r   untyped_storager  r  rI  r   s     rY   is_same_tensorr4    s    NN 	<IIK5::<'	<KKMU\\^+	< JJ%++%	< KK5<<'		<
   "++-1F1F1H1Q1Q1SS	< !U%9%9%;;r   c                v   | j                   xr | j                         |j                         k(  xr | j                  |j                  k(  xrn | j                  |j                  k(  xrS t        j
                  j                  j                  |       t        j
                  j                  j                  |      k(  S r   )r1  rH  r   r   rR   r$  mkldnnr  r3  s     rY   is_same_mkldnn_tensorr7    s     	PIIK5::<'	PJJ%++%	P KK5<<'	P II%%d+uyy/?/?/H/H/OOr   c                      y)N)r  isnanlogical_notlogical_andsignbitand_leltgegteqner#  xorr   r   r   rY   boolean_opsrE    s    r   c                  "    e Zd ZU ded<   ded<   y)OpDtypeRuler5   type_promotion_kindtorch.dtype | Noneoverride_return_dtypeNrY  r   r   rY   rG  rG    s    88--r   rG  zdict[str, OpDtypeRule]op_dtype_propagation_rulesc                *    t        ||      t        | <   y r   )rG  rK  )r   rH  rJ  s      rY   #register_op_dtype_propagation_rulesrM    s    
 (32(t$r   zOrderedSet[str]op_requires_libdevice_fp64c                .    t         j                  |        y r   )rN  r   r  s    rY   #register_op_requires_libdevice_fp64rP    s    ""4(r   c                   ddl m} | s$|j                  j                         j                  } | dk(  rt
        j                  S | dk(  ry| dk(  rt
        j                  S | dk(  rt
        j                  S t
        j                  S )Nr   r  r)  rK   rL   tpu)
r  r  r  get_current_device_or_throwr,  rl   cpu_backendxpu_backendtpu_backendcuda_backend)r   r  s     rY   get_current_backendrX    ss    -gg99;@@e!!!				!!!		!!!"""r   c                x    | 8| j                   dk(  r)t        j                  j                  |       j                  S y)z+Check if the given device supports float64.rL   T)r,  rR   rL   r$  has_fp64r  s    rY   device_supports_fp64r[    s1    fkkU2yy..v6???r   c                    | t         j                  t         j                  fv r7t        j                  j
                  rt               dk(  rt         j                  S | S )z"Maybe upcast [b]float16 to float32r  )rR   r   rl  rl   r  codegen_upcast_to_fp32rX  rn  rS  s    rY   upcast_compute_typer^     s@     	%--00MM00!X-}}Lr   KeyTypeValTypec                  Z    e Zd ZdZddZddZddZddZdddZddZ	dd	Z
dd
ZddZy)
ScopedDictz
    A dictionary-like object that allows for scoped updates. It maintains
    an original dictionary and a set of new items that can override
    the original items within the scope.  The original dictionary is
    unmodified.
    c                     || _         i | _        y r   original_dict	new_items)r  re  s     rY   rd  zScopedDict.__init__7  s    *13r   c                Z    || j                   v r| j                   |   S | j                  |   S r   rf  re  r  s     rY   r5  zScopedDict.__getitem__;  s.    $.. >>#&&!!#&&r   c                "    || j                   |<   y r   )rf  )r  r  r   s      rY   __setitem__zScopedDict.__setitem__@  s    #sr   c                >    || j                   v xs || j                  v S r   rh  r  s     rY   __contains__zScopedDict.__contains__C  s!    dnn$At/A/A(AAr   Nc                t    || j                   v r| j                   |   S | j                  j                  ||      S r   )rf  re  rj  )r  r  rE  s      rY   rj  zScopedDict.getF  s6    $.. >>#&&!!%%c733r   c                z    t        | j                        }| j                  D ]  }|| j                  vs|dz  } |S r  )rT   re  rf  )r  r=  r  s      rY   r  zScopedDict.__len__K  sC    ""# 	A***Q	 r   c              #     K   | j                   E d {    | j                  D ]  }|| j                   vs|  y 7 )wr   rd  )r  r  s     rY   __iter__zScopedDict.__iter__R  s@     %%%% 	A***	 	&s   ><!>>c                H    t        | j                  xs | j                        S r   )rX  re  rf  rs  s    rY   r~  zScopedDict.__bool__X  s    D&&8$..99r   c                    t         r   r  r  s     rY   __delitem__zScopedDict.__delitem__[  s    !!r   )re  Mapping[KeyType, ValType])r  r_  r   r`  )r  r_  r   r`  r   r  )r  r  r   rX  r   )r  r_  rE  ValType | Noner   ru  r  )r   zIterator[KeyType]r  )r  r_  r   r  )r   r   r   r   rd  r5  rj  rl  rj  r  rp  r~  rs  r   r   rY   rb  rb  /  s5    4'
$B4
:"r   rb  )frozen_defaultc              (    dfd}| |S  ||       S )Nc                4    t        j                  | d      S )NT)kw_onlyr   )dataclasses	dataclass)r   r   s    rY   wrapzir_dataclass.<locals>.wrapa  s    $$S$vFFr   )r   rp   r   rp   r   )r   r   r|  s    ` rY   ir_dataclassr}  _  s    G {9r   c                     t         j                  j                  j                         } | "| j                  r| j                  j
                  S y r   )rR   r  r  r  fw_metadatabw_donated_idxs)r  s    rY   get_donated_idxsr  i  s=    mm22::<O"'B'B**:::r   c                       e Zd ZdZdZdZdZdZy)TritonAttrsDescriptorVersionr   r9   r  r<  r_  N)r   r   r   V0_NO_TRITONV1_COMPILERV2_BACKENDSV3_BACKENDS_TUPLEV4_DICTr   r   rY   r  r  p  s     LKK	  Gr   r  c                 P   t         j                  j                  d      t        j                  S dd l} dd l} t        | j                  j                  d      rt        j                  S t        | j                  j                  d      rt        j                  S t        j                  S )Nr  r   AttrsDescriptor)r  r  r  r  r  r  r  r#  r  compilerr  r  r  )r  s    rY   #get_triton_attrs_descriptor_versionr  z  s{    ~~)1+888##v''):; ,777	))+<	=+777 ,333r   c                 8    t               t        j                  k(  S r   )r  r  r  r   r   rY   triton_version_uses_attrs_dictr    s    .04P4X4XXXr   c                    | j                         }t        | t        j                  j                        r| d| j
                   n|}||fS )Nrt   )r   r   rR   r  r  _overloadname)r  op_overload_packet_nameop_overload_names      rY   get_op_namesr    sR    #%779 b%**//0 #
#1R%5%5$67$ 
 #$444r   c                j   ddl m} | j                  }t        |t        j
                  j                        sy|t        j                  j                  j                  j                  t        j                  j                  j                  j                  t        j                  j                  j                  j                  fv rq ||| j                  | j                  d      }|O|\  }}|d   }|D ]@  }||j                  d   j                   t        j"                  t        j$                  fv s@ y y)a  
    Check if an FX node is cudagraph-unsafe based on its input arguments.

    Some ops are only cudagraph-unsafe depending on their inputs (e.g., index_put
    with boolean indices triggers .nonzero() during capture, but integer indices
    are safe).
    r   )normalize_functionFT)normalize_to_only_use_kwargsrE  r  )torch.fx.operator_schemasr  r  r   rR   r  r  r$  r  	index_putrE  
index_put__unsafe_index_putr   r  r*  r   rX  r}  )r  r  r  
normalizedr   r  rE  rv  s           rY   ,_fx_node_is_input_dependent_cudagraph_unsafer    s     =^^Ffejj334 		  ((		!!))		((00 
 (GLL'..t

 !"IAvY'G  ?sxx'<'<JJKKA (    r   c                   | j                   }t        |      t        v ryt        |t        j
                  j                        r1t        j                  j                  j                  |j                  v ryt        |       ry| j                  j                  d      x}Kt        |t        t        f      s|gn|}|D ]+  }t        |t        j                         s|j"                  s+ y y)a   
    Check if an FX node is cudagraph-unsafe.

    This includes:
    - Ops in FORBIDDEN_CUDAGRAPH_OPS (CPU sync, dynamic alloc, etc.)
    - Ops with the cudagraph_unsafe tag
    - Input-dependent unsafe ops (e.g., index_put with boolean indices)
    - Ops with sparse tensor outputs
    Tr  F)r  r&  FORBIDDEN_CUDAGRAPH_OPSr   rR   r  r  r!  r  cudagraph_unsafer  r  r*  rj  r  r  r  	is_sparse)r  r  r  valsr   s        rY   r  r    s     ^^F 6{-- 	65::001HHLL))V[[8 4G< ||&&3&sT5M:u 	A!U\\*q{{	 r   c                    ddl m} t        | |j                  |j                  f      ryt        | |j
                  |j                  f      syt        | dd      }|t        |      ryy)ah  
    Returns True if the node is an op that is not cudagraphable.
    This includes:
    - Ops in FORBIDDEN_CUDAGRAPH_OPS (CPU sync, dynamic alloc, etc.)
    - Ops with the cudagraph_unsafe tag
    - index_put_ with boolean indices (triggers .nonzero() during capture)
    - Control flow nodes (Conditional, WhileLoop)
    - Ops with sparse tensor outputs
    r9   r   TFr  N)	r"  r!  r   Conditional	WhileLooprC  rA   rQ   r  )r  r!  r  s      rY   is_cudagraph_unsafe_opr    sa      $67dR..@AdIt,G:7Cr   c                    t         j                  j                  dd      } t        j                         rUddlm}  |       }|rFt         j                  j                  |dd      }| r!t         j                  j                  || g      n|} | S )NLD_LIBRARY_PATHr"  r   )get_runtime_pathr  lib)
r  r  rj  rl   r  libfb.py.parutilr  r  r4  pathsep)r  r  runtime_pathlib_paths       rY   get_ld_library_pathr    sg    ::>>+R0D5')ww||L)UCH8<2::??Hd#34(DKr   c                F    ddl m} t        | |      xr | j                  d uS )Nr   )SubgraphPythonWrapperCodegen)torch._inductor.codegen.wrapperr  r   partition_signatures)r  r  s     rY   #is_codegen_graph_partition_subgraphr    s*    L 	789 	5((4r   c                     t         j                  j                  j                  j                  xs t
        j                  d uxr$ t         j                  j                  j                  S r   )rR   r  rl   r  
cudagraphs&_unstable_customized_partition_wrapperr  graph_partitionr   r   rY   is_using_cudagraph_partitionr    sN    %%00 	F199E1 //
 
 
0
01r   c                    ddl m} |j                  j                  j	                  | d      r6|j                  j                  j                  | d      rt        j                  S t        j                  S )Nr9   r  l        i   )	r  r  r  r  statically_known_ltr  rR   r{  r|  )rH  r  s     rY   dtype_from_sizer  "  sP    ww++e
''


/
/h
?{{{{r   )r)  rL   c                h    | dk(  r(t         j                  j                  j                         S d| v ryy)z;
    Returns True if the device supports MKL-DNN BF16.
    r)  rL   TF)rR   r$  r6  _is_mkldnn_bf16_supportedr   s    rY   is_mkldnn_bf16_supportedr  0  3     eyy99;;	+	r   c                h    | dk(  r(t         j                  j                  j                         S d| v ryy)z;
    Returns True if the device supports MKL-DNN FP16.
    r)  rL   TF)rR   r$  r6  _is_mkldnn_fp16_supportedr  s    rY   is_mkldnn_fp16_supportedr  <  r  r   c           
     n   |D cg c]  }t        t        |             }}| D ]R  }t        |      t        |      k(  sJ t        |      D ])  \  }}t        ||   t        t        |                  ||<   + T g }|j	                  dj                  d t        ||      D                     t        |      t        |      dz  z   t        |      dz
  z   }|j	                  d|z         | D ]3  }|j	                  dj                  d t        ||      D                     5 dj                  |      S c c}w )N|c              3  6   K   | ]  \  }}d || dd   ywr   r"  Nr   )r   hr  s      rY   r   ztabulate_2d.<locals>.<genexpr>O  s$     H41aAa0tWA,H   r  r9   r  c              3  6   K   | ]  \  }}d || dd   ywr  r   )r   r   r  s      rY   r   ztabulate_2d.<locals>.<genexpr>T  s$     Htq!!QCp4lHr  re  )rT   r&  rK  r   r  r4  r   r?  )elementsheadersr   widthsrowr   r  total_widths           rY   tabulate_2dr  H  s   #*+ac#a&k+F+ 43x3w<'''cN 	4DAqF1Is3q6{3F1I	44 E	LLH3w3GHHIf+Vq1S[1_EK	LL{"# JSXXHs37GHHIJ99U ,s   D2c              #     K   t        | j                               t        |j                               z  }|D ]3  }| j                  |      }|j                  |      }|||n|||n|f 5 yw)a  
    Zip two dictionaries together, replacing missing keys with default values.

    Args:
        dict1 (dict): The first dictionary.
        dict2 (dict): The second dictionary.
        d1_default (Any): the default value for the first dictionary
        d2_default (Any): the default value for the second dictionary

    Yields:
        tuple: A tuple containing the key, the value from dict1 (or d1_default if missing),
               and the value from dict2 (or d2_default if missing).
    N)r!   rk  rj  )dict1dict2
d1_default
d2_defaultall_keysr  value1value2s           rY   	zip_dictsr  X  sv     ( %**,'*UZZ\*BBH  	
33 (Fj(Fj
 	
	
s   A-A/c                T   	 	 	 	 	 	 	 	 dd}	 	 	 	 	 	 	 	 dd}| j                  dt        j                  j                        }| j	                         } |rm || dd        || dd        || dt
        j                  j                           || dd	        || d
t        j                  j                          || dd       | j                  dt        j                  j                        }| j                  dt        j                  j                        }|dk(  r|rt        d      | S )a6  
    Ensures the configuration is internally consistent for standalone AOTInductor.

    If `aot_inductor_mode.compile_standalone` is set to True in the provided
    `config_patches` (or falls back to the global config), this function ensures
    that the following configs are also enabled:
        - `aot_inductor.package_cpp_only`

    Args:
        config_patches (dict[str, Any]): A dictionary of user-provided config
            overrides for AOTInductor compilation.

    Returns:
        dict[str, Any]: The possibly-updated `config_patches` dictionary.
    c                    | j                  |t        t        |            }||| |<   y |s||k7  rt        d| d| d      y y )NzInvalid config: =z3 when aot_inductor_mode.compile_standalone is True.)rj  rQ   rl   r  config_patchesconfig_nameconfig_valuer   s       rY   patch_configz2maybe_aoti_standalone_config.<locals>.patch_config  s]     "";0LM=*6N;'5L0";-q>qr  1r   c                    | j                  |t        t        |            }||k7  rt        j	                  d||       || |<   y )NzDOverriding: %s=%s when aot_inductor_mode.compile_standalone is True.)rj  rQ   rl   r   r  r  s       rY   force_patch_configz8maybe_aoti_standalone_config.<locals>.force_patch_config  sF     "";0LML KKV
 '3{#r   z$aot_inductor_mode.compile_standalonezaot_inductor.package_cpp_onlyTz aot_inductor.embed_kernel_binaryz#aot_inductor.emit_multi_arch_kernelz+aot_inductor.model_name_for_generated_files
aoti_modelzaot_inductor.link_libtorchzaot_inductor.dynamic_linkageFz"aot_inductor.cross_target_platformz$aot_inductor.package_constants_in_sowindowszconfig.aot_inductor.package_constants_in_so is not supported for windows cross-compilation. Please use config.aot_inductor.package_constants_on_disk_format = binary_blob.)r  dict[str, Any]r  r&  r  r   r   r  )rj  rl   aot_inductor_modecompile_standalonecopyrR   ro  ru   test_configsuse_libtorchaot_inductorcross_target_platformpackage_constants_in_sor  )r  r  r  r  r  r  s         rY   maybe_aoti_standalone_configr  {  sk   "	&	58	HK			
3&
358
3HK
3	
3 (++.  33
 $((*N^%DdK^%GNAu}}GXGXCX	
 	I<	
 	(,,	

 	>+I5Q*..,11
 -00.33
 	).E]
 	

 r   c                   t         j                  j                  r(t         j                  j                  dk(  rt	        d      t         j                  j                  r0t         j                  j
                  dk(  rt	        d      d}d}||fS t         j                  j                  dk(  rd}d}||fS | dk  ryd}t        j                          }||fS )	a  
    Decide whether we should mmap weights, and whether to store the weights with .so.

    If force_mmap_weights or package_constants_on_disk_format == "binary_blob" configs are set, respect the config.

    Returns tuple (use_external_weights, use_mmap_weights).
    binary_blobzconfig.aot_inductor.package_constants_on_disk_format = binary_blob and config.aot_inductor.force_mmap_weights cannot both be True.r  zKwhen cross_target_platform is windows, use_mmap_weights should not be true.TFi 5w)FF)rl   r  force_mmap_weights package_constants_on_disk_formatr  r  r  )consts_sizeuse_mmap_weightsuse_external_weightss      rY   determine_aoti_mmap_flagsr    s     	..@@MQJ
 	

 --44	A]   $#%555;;}L# #%555m# !++--!111r   c                     ddl m}  | j                  j                  }|yt	        |t
              st        d      |dk(  ryt        j                  d|      st        d      y)zL
    Validates if a model name is suitable for use in code generation.

    r   rk   Tz4Invalid AOTI model name: Model name must be a stringr"  z^[a-zA-Z_][a-zA-Z0-9_]*$zVInvalid AOTI model name: Model name can only contain letters, numbers, and underscores)	r  rl   r  model_name_for_generated_filesr   r&  r  r   r   )rl   
model_names     rY   is_valid_aoti_model_namer    sh    
 '$$CCJj#&OPPR 88/<d
 	
 r   c                2    |rt        |       S t        |       S r   )r*   r)   )rV   unbacked_onlys     rY   get_free_symbolsr    s    $Q''Ar   c                    i t         j                  dt         j                  j                  dt         j                  j	                  t
        j                              i} t        j                         rt        j                  d      | d<   | S )zA
    Get a base environment for running Python subprocesses.
    
PYTHONPATHTORCH_CUSTOM_PYTHONPATHrI  
PYTHONHOME)r  r  rj  r  r4  r  r  rl   r  	sysconfigget_path)envs    rY   python_subprocess_envr    sl    

** 	bjjnn%rzzsxx'@
	C  %..v6LJr   c                  &    e Zd ZU dZded<   ded<   y)CUDAGraphWrapperMetadataz
    Metadata for Customized CUDAGraphWrapper.

    Currently assumes there is 1 dynamo graph and will extend to
    multiple graphs in the future.
    r   num_partitionspartition_indexNr   r   r   rY   r  r  6  s      r   r  .c                      e Zd ZU dZded<   y)CUDAGraphWrapperNzCUDAGraphWrapperType | Noner  )r   r   r   r  r   r   r   rY   r  r  M  s    +/G(/r   r  c                    | t         _        y r   )r  r  )r  s    rY   !set_customized_partition_wrappersr  _  s    5<*2r   c                   t        d| j                        }t        |t        j                  j
                  j                        r*|j                  |j                  |j                        \  }}n&g |j                  |j                  }|j                  }|j                  ||      }t        |      }t        |dd       }t        |t        j                  j                        r_|j                   j"                  D cg c]  }|j$                  r|j&                   }}|d t)        |       D ]  }|j+                  |d         t-        j.                  ||f      \  }}	dd}
|D cg c]7  } |
|      r+t        j                  j
                  j1                  |d      n|9 }}d	dd
fd}|D cg c]
  } ||       }}t-        j2                  ||	      \  }}||fS c c}w c c}w c c}w )NrA   rB  c                   t        | t        j                  j                  j                        xrS t        | t        j                  j                  j
                  t        j                  j                  j                  f       S r   )r   rR   r  r!  rB   GeneratorStateOpaqueObjectStaterW  s    rY   _is_tensor_irz(snode_args_kwargs.<locals>._is_tensor_irv  s^    !U__//667 

__..0B0B0T0TUA
 =
 	
r   T)replace_symbols_with_hintsc                2    t        j                  | ||      S )Nr   )rR   r   )rH  r   r   s      rY   _tensorz"snode_args_kwargs.<locals>._tensor  s    {{4uV<<r   c                    t        | t        j                        s| S  | j                         | j                  | j
                        }|S r   )r   rR   r  rH  r   r   )r   r  r  s     rY   to_real_tensorz)snode_args_kwargs.<locals>.to_real_tensor  s7    !U\\*Haffh2
r   r  )r   r  )r   r   r   r   )r   r  r   rR   r  r!  rC  unflatten_argsr  constant_argsr  fill_non_provided_argsr$  rQ   r  r  r  r  
kwarg_onlyr   rT   rU   pytreer"   ir_node_to_tensortree_unflatten)rQ  r  r   r  rB  r  positional_arg_namesarg_name	flat_argsflat_args_pytree_specr  r<  r  r  s                @rY   snode_args_kwargsr$  c  s   

+D$**99:**4;;8J8JKf22t112&&tV4D&\F$t4K+uzz445 + 3 3 = = 
S^^CHH 
  
 -[s4y9 	'HJJx&	''-':':D&>'J$I$
 	   	,,Q4,P	I = -66q"6I6((4IJLD&<? 
  7s   <GG%<G$1G)c                    ddl m} | j                  }|j                  j                  r(|j	                  |j                  j                  dz         }|j                  d      S )Nr9   r  r   )r  r  fwd_rng_stater{  rx  )r  r  r   r  r  r  )rf  r  dep_names      rY   is_nonfreeable_buffersr(    sN    xxH 	ww||(();<I r   c                p    t        ||  dz        5 }|j                         cddd       S # 1 sw Y   yxY w)z,Load a template file and return its content.z	.py.jinjaN)openread)r   template_dirr0  s      rY   load_templater-    s6    	lvY//	0 Avvx  s   ,5c                v   | j                   }t        |t        j                  j                  t        j                  j
                  f      sJ dt        |              t        j                  syt        t        j                  j                  j                  j                  t        j                  j                  j                  j                  g      }||v ryt        t        j                  j                  j                   g      }t        |t        j                  j
                        r||v S t#        |        S )zLDecide whether fallback for a node. This is only used in inductor lite mode.z6Expected OpOverload or HigherOrderOperator, but found F)r  r   rR   r  r  r+  r,  rl   fallback_by_defaultr!   r$  r  _assert_scalarrE  lift_fresh_copyhigher_order triton_kernel_wrapper_functionalr   )r  r  "skip_fallback_due_to_dynamic_shapefallback_hopss       rY   should_fallback_by_defaultr6    s    [[F&&

(F(FG O	?V~NO  %% *4IINN))11IINN**22	
*& 33 				@	@AM &%**889&&&t,,,r   )z-torch.ops._c10d_functional.all_reduce.defaultz.torch.ops._c10d_functional.all_reduce_.defaultz9torch.ops._c10d_functional.all_gather_into_tensor.defaultz8torch.ops._c10d_functional.reduce_scatter_tensor.defaultz4torch.ops._c10d_functional.all_to_all_single.defaultz6torch.ops._c10d_functional_autograd.all_reduce.defaultzBtorch.ops._c10d_functional_autograd.all_gather_into_tensor.defaultzAtorch.ops._c10d_functional_autograd.reduce_scatter_tensor.defaultz=torch.ops._c10d_functional_autograd.all_to_all_single.defaultz(torch.ops._c10d_functional.isend.defaultz(torch.ops._c10d_functional.irecv.defaultz0torch.ops._c10d_functional.batch_p2p_ops.defaultc                    | t         v S )z0Check if an operation is a collective operation.)COLLECTIVE_OPS)r  s    rY   is_collective_opr9    s    n$$r   c                 b    t        j                         r		 ddlm}  | S g S # t        $ r g cY S w xY w)Nr   tlx_only_cuda_options)rl   r  )torch._inductor.fb.tlx_templates.registryr<  r%  r;  s    rY   r<  r<    s9    	W(( 		  	I	s     ..c                    | |z   dz
  |z  |z  S )z(Round x up to the nearest multiple of y.r9   r   )rV   r  s     rY   	_round_upr?    s    UQY1!!r   c                   ddl m}m}  ||d      r|j                  |j                  fS t        |      dk\  r  ||d   | d         r ||d   d      s ||d   d      r' ||d   | d         r|j                  |j                  fS  ||d   | d         r ||d   t        | d   d            s( ||d   | d         r1 ||d   t        | d   d            r|j                  |j                  fS  ||d   t        | d   d            r1 ||d   t        | d   d            r|j                  |j                  fS |t        j                  k(  rdnd}|t        j                  k(  r|t        j                  k(  rt        | d   d      t        t        || d   z  d      d      z  }	t        | d   d      t        t        || d   z  d      d      z  }
 |||	      s	 |||
      r|j                  |j                  fS |t        j                   k(  rt        j"                  j$                  st        | d   d      t        t        || d   z  d      d      z  }	t        | d   d      t        t        || d   z  d      d      z  }
 |||	      s	 |||
      r|j&                  |j                  fS y	t        | d   d      |z  | d   z  }	t        || d   z  d      | d   z  }
 |||	      s	 |||
      r|j&                  |j                  fS y	)
z:
    Core implementation for scale/swizzle inference.
    r   )r8   SwizzleTyper9   r  rw   rv   r_  rz  NN)torch.nn.functionalr8   rA  
TensorWise
NO_SWIZZLErT   RowWiserm   BlockWise1x128BlockWise128x128rR   rj  rb  r?  BlockWise1x16SWIZZLE_32_4_4rh  ro  ru   BlockWise1x32)mat_size
scale_sizescale_numel	mat_dtypescale_dtypeeq_fnr8   rA  K_multiplierexpected_numel_aexpected_numel_bs              rY   _infer_scale_swizzle_implrU    sD    = [!%%{'='=== :!*Q-!-%
1q2I*Q-#jmXa[(I&&(>(>>> *Q-!-jmWXa[#%>?*Q-!-jmWXa[#%>?--{/E/EEE AS 9:uqM78A;4@
 //1G1GGG "U%;%;;1L E***{e>Q>Q/Q$Xa[#6L8A;.3Q:
 
 %Xa[#6L8A;.3Q:
 
 ./5FV3W,,k.H.HHH e***}}  (!c:Yx{2B7>    )!c:Yx{2B7>   ["23u-8 #00+2L2LLL   'x{B7,FRST&|hqk'A2FRST["23u-8 #00+2H2HHHr   c                    t        | j                  d   | j                  d   ft        |j                        |j                         | j                  |j                  d       S )a  
    Infer the scaling type and swizzle mode from matrix and scale tensor shapes/dtypes.

    This function determines how scale factors are laid out relative to the matrix:
    - TensorWise: Single scale for entire tensor
    - RowWise: One scale per row
    - BlockWise1x128/128x128: Block-scaled with float32 scales
    - BlockWise1x32: MXFP8 with float8_e8m0fnu scales (swizzled on NVIDIA)
    - BlockWise1x16: NVFP4 with float8_e4m3fn scales (swizzled)

    Args:
        mat: The matrix tensor (FP8 or FP4)
        scale: The scale factor tensor

    Returns:
        Tuple of (ScalingType, SwizzleType) or (None, None) if unrecognized
    r   r9   c                    | |k(  S r   r   r  s     rY   r  z%infer_scale_swizzle.<locals>.<lambda>`  s
    16 r   rL  rM  rN  rO  rP  rQ  )rU  rS  r  numelr   )matscales     rY   infer_scale_swizzler\  F  sO    ( %))A,		!-%KKM))KK! r   c                \   ddl m | j                         }|j                         }|r
|d   |d   f}|r%t        j                  t
        j                  |d      nd}dfd}t        t        |      dk\  r
|d   |d   fn|d   dft        |      || j                  |j                  |      S )z
    Infer the scaling type and swizzle mode for IR nodes (used during graph lowering).

    This is the IR-compatible version of infer_scale_swizzle, using symbolic
    size comparisons via V.graph.sizevars.statically_known_equals.
    r   r  r9   c                P    j                   j                  j                  | |      S )z5Compare values using symbolic equality when possible.)r  r  r~  )r<  r=  r  s     rY   symbolic_eqz+infer_scale_swizzle_ir.<locals>.symbolic_eq{  s     ww771==r   r  rX  )r<  r   r=  r   r   rX  )r  r  rq  r2  r3  r4  r5  rU  rT   r  r   )rZ  r[  	transposerL  rM  rN  r_  r  s          @rY   infer_scale_swizzle_irra  d  s     .||~H!J QK!- DN)""8<<Q?STK> %/28}/A(1+x{+QRUVGW$))KK r   r  )r}   r   r   r   )r   r   r   rX  )   d   )r   Callable[[], Any]r   r   r   r   r   r  )rb  rc  F)
r   rd  r   r   r   r   r   rX  r   r  )
r
  Iterable[Any]r  re  r   r   r  r]   r   r  r  )r   ztorch.device | None | strr   torch.device)r8  zIterable[sympy.Expr]r   r   )r@  r  rA  r  r   r   )rE  r  rF  r  r   r   )r+  r   rF  r  r   list[sympy.Expr])r8  zIterable[_T]r   zValuesView[_T])rU  r  rV  r  r   r  )r  rI  r   r&  )r  zIterable[int | torch.SymInt]r   rg  )r  int | torch.SymIntr   r  )r   r  r   rh  )r  zIterable[int | sympy.Expr]r   zlist[int | torch.SymInt])r  torch._ops.OpOverloadr   rX  )r  r7   r  z'Callable[[torch._ops.OpOverload], bool]r   rX  )r  r   r   r  r  r  r   z&tuple[GraphModule, list[torch.Tensor]])rJ   )r   r&  r   r  )r9   rJ   )
r  Callable[..., Any]r  Sequence[Any]r   r   r   r&  r   r  )r   r  r  g      ?rJ   )r  rj  r  rk  r   r   r  r   r  r  r   r&  r   r  )r  r   r  r&  r   r  )r  r   r  r   r   r  )r<  r   r=  r   r   r   )rV   zint | Sequence[int]rH  r   r   Sequence[int])rV   ztuple[_T, ...]r   zlist[_T])r   z!Callable[Concatenate[Any, P], RV]r   zCachedMethod[P, RV])r  r&  r   z*Callable[[FN_TYPE[P, RV]], FN_TYPE[P, RV]])r%  *Sequence[BaseSchedulerNode] | ExternKernelr   zOrderedSet[Node])r%  Sequence[BaseSchedulerNode]r5  z8Literal[True, 'torch', 'original_aten', 'inductor_node']r   r&  )r%  rm  r  r<   r   tuple[str, str]r   )r  zIterable[torch.fx.Node]r  zCallable[[Any], bool] | Noner   OrderedSet[torch.fx.Node])r   zSequence[IRNode]r  zdict[str, IRNode]r   rp  r  )r+  r   r   zValueRanges[Any])r  r&  r   rX  )r  rh   rv  r   r   sympy.Symbol)r  rX  r   rX  )r   r&  r   rq  )r  r   r  zdict[sympy.Expr, Any]r   r   )r<  r   r   z&TypeGuard[torch.SymInt | torch.Tensor])r   r   r   rX  )r  torch.fx.GraphModuler   ztorch.fx.Node | None)r  rr  r   r7   )r  rr  r   zOrderedSet[torch.device]r  )r  r   r   r   )r  r&  r   r&  r   r  )NNT)r,  zdict[str, Any] | Noner  r  r-  rX  r   r  )r6  rk  r2  rX  r   	list[int])r  r-   r6  z)Sequence[int | torch.SymInt | sympy.Expr]r2  rX  r   rs  )r   r  r   r   r  r  r  )r   zint | torch.devicer   rX  r  )r9  r   r   rf  r:  
int | Noner   r:   )rv   )r?  r   r   r   )rO  rC   rC  zlist[torch.dtype]r   rX  )rI  r&  r   rX  )
rO  rC   rV  rX  rW  rX  rX  rX  r   rX  )r  rB   rc  r  rd  rX  r   rX  r  )rF  r  rd  rX  r   rX  )r  rB   rc  rC   rd  rX  r   rX  )r  r8   r  r8   r  zlist[ScalingType]r   rX  )r  r   r  r   rO  rC   r  rX  r  rX  r  
Any | Noner  ru  r  ru  r   rX  )
rO  rC   r  r   r=  r   r  r   r   rX  rB  )rO  rC   r  r  r=  r  r  r  r  rB   r  rB   r  zIRNode | Noner  z_IntLike | Noner   rX  )r  r&  r   rX  r   )
r  r  r=  r  r  r  r  r   r   rX  )r  r  r=  r  r  r  r   rX  )r  r  r=  r  r  r  r   rs  )r   r&  r   r&  )r   zNtuple[str | None, Callable[[], list[Any]], Callable[[], list[Any]], type[Any]])rO  rC   r   rX  )rO  rC   r@  zReinterpretView | BufferrA  rB   r   rX  )FTFN)rO  rC   r@  rB   rA  rB   rL  rX  r=  rX  rd  rX  rT  rt  r   rX  )r   Callable[P, _T]r   r  r  r  r   ztuple[_T, list[str]])r   rj  r   ztuple[Any, list[str]])r   rv  r   r  r  r  r   r   )r   rv  r   r  r  r  r   r&  )r   rv  r   r  r  r  r   ztuple[Any, list[GraphLowering]])r  rj  r  rj  r   r  )r  rj  r  zCallable[..., Any] | Noner   r   )r  r&  r   r  r  )rl  rk  r   rX  )r  zSequence[torch.Tensor]r   rX  )r  r   r   r  )r  rX  r   r   r  r   r   zIterator[Any])r   r  r   r  )r
  r&  r   rX  )r
  r&  r   r   )r  re  r   rX  )
r  rj  r  r6   r/  rk  r  r&  r   r  )r8  zBuffer | Operation | Noner   rX  )r  zNode | Operation | Noner  ztorch._ops.OperatorBase | Noner   rX  )r  zIRNode | Operation | Noner   rX  )rQ  rH   rR  z*Callable[[BaseSchedulerNode], bool] | Noner   rX  )rQ  rH   r   rX  )r  zOperation | Noner  z9torch._ops.OpOverload | Collection[torch._ops.OpOverload]r   rX  )rZ  r&  r[  r  r\  r  r   r   )rQ  rH   re  zMutableSet[BaseSchedulerNode]r[  zdict[str, SchedulerBuffer]r\  zdict[str, BaseSchedulerNode]rb  zCallable[[Any], bool]r   r  )rr  r   rs  r   r   r   )r}  rr  r   r   )r}  rr  r   rs  )r  r   r   r  )r   r&  r   r&  )r   r  r   rX  )r  r&  r&  z	int | strr  r   r   ztuple[str, ...] | None)r   r  r   ztorch.device | str | Noner  r  r  zint | str | Noner  rt  r   rX  )r   r&  r   rX  )r   r  r   rX  )rB  ri  r
  r  r  r  r  r  r  r&  r  rX  r   rX  )r%  rn  r   r  )r   r  r   rX  )r  r  r   rX  )r   r  )r   rv  r   r  r  r  r   ztuple[_T, str])r  Sequence[InputType]r   zShapeEnv | None)r  Callable[[list[InputType]], _T]r  rl  r  zOrderedSet[int]r   rx  )rV   r  r   r  )r  r  r  rl  r  zOrderedSet[int] | Noner   z-tuple[list[torch.Tensor], list[torch.Tensor]])r  rw  r  rl  r   rl  )r   r   r   rX  )r  rk  r  rG   r   r  )r   r  r   r&  )r   r&  r   r  )rI  r  r   r  r   rX  )r   ztuple[str, ...])r   r&  rH  r5   rJ  rI  r   r  )r   r&  r   r  )r   r  r   r&  )r   ztorch.device | Noner   rX  )r   r  r   r  )r   ztype[Any] | Noner   rX  r   r   )r   zlist[int] | None)r   r  )r  ztorch._ops.OperatorBaser   ro  )r  torch.fx.Noder   rX  )r  rD   r   rX  )r  r<   r   rX  )rH  r   r   r  )r   r&  r   rX  )r  zSequence[Sequence[T]]r  zSequence[T]r   r&  )
r  rt  r  rt  r  ru  r  ru  r   zEGenerator[tuple[KeyType, ValType | None, ValType | None], None, None])r  r  r   r  )r  r   r   ztuple[bool, bool])rV   r,   r  rX  r   zOrderedSet[sympy.Symbol])r   zdict[str, str])r  CUDAGraphWrapperTyper   r  )rQ  rH   r   z tuple[list[Any], dict[str, Any]])rf  r=   r   rX  )r   r&  r,  r1   r   r&  )r  ry  r   rX  )r   r   )rV   r   r  r   r   r   )rL  ztuple[Any, Any]rM  ztuple[Any, ...]rN  r   rO  r  rP  r  rQ  zCallable[[Any, Any], bool]r   tuple[Any | None, Any | None])rZ  r  r[  r  r   r{  )rZ  r@   r[  r@   r`  rX  r   r{  (  
__future__r   rf  r  rz  enumr2  r  r  r*  r  r  r  r4  r  r  r   r)  r   r  r  r"  r  r  r  collections.abcr   r   r   r   r   r	   r
   r   r   r   typingr   r   r   r   r   r   r   r   r   r   r   typing_extensionsr   r   r   r   r   rR   torch.utils._pytreer  _pytreer  $torch._inductor.analysis.device_infor   torch._inductor.runtime.hintsr   !torch.fx.passes.regional_inductorr   torch.utils._dtype_abbrsr    torch.utils._ordered_setr!   r"   r#   r  r$   OPTIMUS_EXCLUDE_POST_GRAD#torch.fx.experimental._size_hintingr(   r  r)   r*   r+   r,   r-   r.   r/   r0   pathlibr1   r2   r3   r4   torch._prims_commonr5   torch.fxr6   torch.fx.noder7   rC  r8   r5  r:   r  r<   dependenciesr=   r  r?   r!  r@   rA   rB   rC   rD   rE   output_coderG   r  rH   rI   rP   rN   r   rZ   torch._dynamo.device_interfacer[   torch._dynamo.utilsr\   torch.autogradr]   torch.autograd.profiler_utilr^   (torch.fx.passes.graph_transform_observerr_   torch.fx.passes.shape_propr`   torch.utils._sympy.functionsra   rb   rc   rd   re   rf   torch.utils._sympy.symbolrg   rh   torch.utils._sympy.value_rangesri   rj   r"  rl   runtime.runtime_utilsrm   rT  _IS_WINDOWS	getLoggerr   r   r  rp   r$  rR  	VarRangesr  r   	InputTypegetenvXPU_KERNEL_FORMATGPU_KERNEL_BIN_EXTSr  r  ri  r7  r}  rr  r~  rt  r  r{  r|  r   rl  rn  rp  rb  rc  rd  rf  rx   r   r|   r~   r   Functionr   r{  r   r   r   r  r   r'  r0  r9  rB  rI  rL  rO  r  r  r  r  r  r  r  r  r   r  r  r  r  r  r  r  r  r  FN_TYPEr  r  r  r  r&  r:  r  r  r  r  r  r  r  r  r  r  r  r  r  	frozensetro  ru   r  r  r  r  r   r  r  r  r  r  r  r1  clear_on_fresh_inductor_cacheclear_inductor_cachesfresh_inductor_cacher:  rQ  rU  rW  r[  r^  r  r  r  r  r  rj  r  r"  r&  r)  r,  r0  r;  r@  rD  rJ  rO  rR  rU  ra  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r#  r-  r2  r5  r8  r:  rE  r>  rh  rj  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  rb  r  r  r  r   r  r  r  r  r(  r_  r  Enumr  r4  r6  r;  rG  rJ  rO  rU  rW  r]  rd  rl  ru  r  r  r  r  r  rB  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r!  r$  r(  rl  r,  compiler'  r*  r/  r4  r7  rE  rG  rK  rM  rN  rP  rX  r[  r^  r_  r`  rb  r}  r  r  r  r  r  r  r  r  r  r  r  r  SUPPORTED_MKLDNN_DEVICESr  r  r  r  r  r  r  r  r  r  PartitionFnTyperz  r  r  r  r$  r(  r-  r6  r8  r9  r<  r?  rU  r\  ra  )r  r   s   00rY   <module>r     s   "        	     	  	   
              C B    $ $ ? : E 0 / ; 2 ($ 
 <  >>//C$"/,5!$TT,= +	CL
   D 0 % 2 K 0  8 D  = llg%g!  C  T]UZZ'(	LL3-4	 Eibii(I7S 
  !"  	 
 2<

2 . ( {Q'A-+2B XDX XB5
LENN  d#  $"GX #(	 
 !	
 86. 6."6. 6. %	6.
 6.v #(	NNN 
N !	N
 Nb  ;@
	!		 	 +	*%AP+	%++	"/	#//G @OI	I<I 
I0 *8+0' 	!  	
 ( %'!  	
    )'#$  cNTT"
;sAv&*
+E8WQU^ E:)++/+\=,4).4)O4) 	4)n_2=_2!_2 _2H 15*- (G
G$5GG:,^%(	DU	>+- $( ==(  	
 + 6'& 
< !# I "	 " "( +/7(7	7 7 	7 7v !5 $ "  49 ( 	77	27 	7
 7t Q7 7*  , , ,
`' `'FL L^-N -,

 
 @ @ @?' ?  8 J J ) )I  $  	(+<	  #  	
  
< 7;uBB&3BHLB	BL 5:-1	4 BGOO&,O:>O	O2 BGVV&,V:>V	V OOO %O 
	O Q	  	 Q  B0 Q  >>> > 	>
 > > > > 
>BP EEE E 	E
 E E E E 
EPJ EJJ&) & EF!)?B	 (  . 5( 5(p @ @ R R:"JH&2:@	: ""&#==
= = 	=
  = = = 
=@'C C"&&& & 	&$ &2:&/(V		 &	2:		## &#2:#$#* ...@.. .$ FJFF)BFF*	B&&   D D .6 .6b  Q.0
(#K(*$)) * 

 
"- 
4A 
HK 
	 
F	(	" *.#
!#&# 
#L( =A	Q	Q9	Q 
	Q#J
JAJ 
JLL .LDRLL *=

5
 ,
 5	

 '
 

< *=5 , 5	
 ' 
:T 2      ,!)
 T''&'36'' '0 )-10 "&$( 1010%10 	10
 "10 10 
10h.2$&$$ $ 	$
 $ $ 
$NHBL'  &2:2:*" ( %	0	: 04$$$$ -$ 3	$<$ $*IZ!3B	:&/ '#)* $%
  +?*D*D*FG$!QAG  "**Y'H	  & . . .
 68 2 7
8 . 
	 /9l O :)# #" )

)
-" 01 -"` D)D  *499  4 42Y5!H"J4
1 * 		& "&!%	 
$ 
$ 
  
 	 

 K 
FRj&2R66 d#  $ 38$./@ 0 0 *:); &=+\
"-L $%
  "
MMM M 	M
 M &M #M`	*"B "	"" " #	"i Hs   3{