+
    –°³j�‡ ã                  óÌ  a€ 0 t $ ^ RIHt ^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RI	t	^ RI
t
^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RIHtHtHtHtHtHtH t  ^ RI!H!t! ^ RIH"t" ^ RI#H$t$H%t%H&t&H't'H(t(H)t)H*t*H+t+H,t, ^ RI-H.t.H/t/H0t0 ^ RIH1t1 ^ RI2t2^ RI3t3^ RI4H5u H6t7 ^ R	I8H9t9 ^ R
I:H;t; ^ RI<H=t= ^ RI>H?t? ^ RI@HAtAHBtBHCtCHDtD ^ RIEHFtF ^ RI4HGtGHHtH ^ RIIHJtJ RR.tK^ RILHMtM ^ RINHOtOHPtPHQtQHRtRHStS ]*'       di   ^ RIHTtTHUtUHVtV ^ RIWHXtX ^ RI3HYtYHZtZH[t[ ^ RI\H]t] ^ RI^H_t_ ^ RI`Hata ^ RIbHctc ^RIdHete ^RIfHgtg ^RIhHiti ^RIjHktk ^R IlHmtmHntnHotoHptpHqtqHrtr ^R!IsHttt ^R"IuHvtvHwtw . ER0Otx]0! R%4      ty]Pô                  R& R' l4       t{^ R(I|H}t} ^ R)I~Ht ^ R*I€H�t� ^ R+I‚Hƒtƒ ^ R,I„H…t… ^ R-I†H‡t‡ ^ R.IˆH‰t‰HŠtŠH‹t‹HŒtŒH�t�HŽtŽ ^ R/I�H�t�H‘t‘ ^ R0I’H“t“H”t” ^R1I•H–t– ^R2I—H˜t™ ]P                   R38H  tš]EP6                  ! ]œ4      t�R4tž]0! R54      tŸ] ]2EPB                  ]2EPB                  3,          t¢]3EPF                  ]¤,          ]3P¶                  ,          R,          t¥]š'       d   R6M]EPL                  ! R7R84      t§R#R9R:R;R$R<]§ 2/t¨^t©^tª^t«^€t¬]F! ]3EPZ                  ]3EP\                  ]3EP^                  ]3EP`                  ]3EPb                  ]3EPd                  ]3EPf                  ]3EPh                  ]3EPj                  ]3EPl                  ]3EPn                  ]3EPp                  ]3EPr                  ]3EPt                  ]3EPv                  .4      t¼R=]½R>&   ^@t¾]¾]¾^,
          ,          ^ 8X  d   ]¾^8¼  g	   ]¿! R?4      hR@ RA ltÀRB RC ltÁ ! RD RE]2EP„                  4      tÃ]EPˆ                  ! RFRG7       ! RH RI4      4       tÅER1RJ RK lltÆER2RM RN lltÇRO RP ltÈER2RQ RR lltÉ]Pô                  RS RT l4       tÊRU RV ltËRW RX ltÌRY RZ ltÍR[ R\ ltÎR] R^ ltÏR_ R` ltÐRa Rb lt˜Rc Rd ltÑRe Rf ltÒRg Rh ltÓRi Rj ltÔRk Rl ltÕRm Rn ltÖRo 3Rp Rq llt×Rr Rs ltØER3Rt Ru lltÙER4Rv Rw lltÚER5Rx Ry lltÛRz R{ ltÜR| R} ltÝR~ R ltÞR€ R� ltßR‚ Rƒ ltà]/! R„4      tá]0! R…RFR†7      tâ]]&]$]á3,          ]â3,          tã ! R‡ Rˆ])]']á]â3,          4      täR‰ RŠ ltåR‹ RŒ ltæR� RŽ ltçR� R� ltèR‘ R’ ltéR“ R” ltêER6R• R– lltëR— R˜ ltìR™ Rš ltíR› Rœ ltîR� Rž ltïRŸ R  ltðR¡ R¢ ltñR£ R¤ ltòR¥ R¦ ltóR§ R¨ ltôR© Rª ltõR« R¬ ltö]÷! . ER7O]3EPð                  EPò                  e   R­.M. ,           4      túR® R¯ ltûR° R± ltüR² R³ ltý^ RIþtþR´ Rµ ltÿ. Et R¶]½R·&   R¸ R¹ lEtRº R» lEt]EP                  R¼ R½ l4       Et]EP                  ER8R¾ R¿ ll4       EtE]EtE]EtE]EtRÀRL/RÁ RÂ llEt	RÀRL/RÃ RÄ llEt
]PD                  ! ^4      RÅ RÆ l4       Et ! RÇ RÈ]B4      Et ! RÉ RÊ]B4      EtRË RÌ lEt ! RÍ RÎ]B4      Et]EP                  RÏ RÐ l4       Et ! RÑ RÒ]A4      Et]Pô                  ER9RÓ RÔ ll4       Et]PD                  RÕ RÖ l4       Et]PD                  R× RØ l4       EtRÙ RÚ lEtRÛ RÜ lEtER6RÝ RÞ llEtER:Rß Rà llEtRá Râ lEtRã Rä lEtRå Ræ lEtRç Rè lEtRé Rê lEtRëRLRìRLRíRF/Rî Rï llEtRðRRñRL/Rò Ró llEtER;Rô Rõ llEt RñRL/Rö R÷ llEt!RñRL/Rø Rù llEt"Rú Rû lEt#]PD                  ! ^Rü7      Rý Rþ l4       Et$]PD                  ! ^Rü7      Rÿ ER  l4       Et%ER Et&]PD                  ! ^Rü7      ER ER l4       Et'ER ER lEt(ER ER lEt)ER<ER ER	 llEt*ER
 ER lEt+]¤]2EPB                  ,          Et,ER]½ER&   ]Pô                  ER=ER ER ll4       Et-]Pô                  ER ER l4       Et.]Pô                  ER ER l4       Et/]Pô                  ER ER l4       Et0]PD                  ER ER l4       Et1]Pô                  ER ER l4       Et2ER ER lEt3ER ER lEt4ER ER lEt5ER  ER! lEt6ER" ER# lEt7ER$ ER% lEt8ER>ER& ER' llEt9ER( ER) lEt: ! ER* ER+4      Et;ER, ER- lEt<ER. ER/ lEt=ER0 ER1 lEt>ER2 ER3 lEt?ER4 ER5 lEt@ER6 ER7 lEtAER8 ER9 lEtB]EP                  ER: ER; l4       EtCER6ER< ER= llEtDER> ER? lEtEER@ ERA lEtFERB ERC lEtGERD ERE lEtHERF ERG lEtIERH ERI lEtJ]EP                  ERJ ERK l4       EtKERL ERM lEtL]Pô                  ERN ERO l4       EtM]Pô                  ERP ERQ l4       EtN]Pô                  ERR ERS l4       EtOERT ERU lEtPERV ERW lEtQERX ERY lEtRERZ ER[ lEtSER\ ER] lEtTER^ ER_ lEtUER` ERa lEtVERb ERc lEtW ! ERd ERe]EP°                  4      EtYERf ERg lEtZERh ERi lEt[ERj ERk lEt\ER6ERl ERm llEt]ERn ERo lEt^ER6ERp ERq llEt_ERr ERs lEt`ERt ERu lEtaERv ERw lEtbERx 3ERy ERz llEtcER{ 3ER| ER} llEtdER~ ER lEteER€ ER� lEtfER‚ ERƒ lEtg]EPˆ                   ! ER„ ER…4      4       Eth]EP                  ER† ER‡ l4       EtiERˆ ER‰ lEtjERŠ ER‹ lEtkERŒ ER� lEtl]PD                  ! R4      ERŽ ER� l4       EtmER6ER�RER‘RER’R/ER“ ER” lllEtnER• ER– lEtoER— ER˜ lEtpER™ ERš lEtqER› ERœ lEtrER� ERž lEtsERŸ ER  lEttER¡ ER¢ lEtuER£ ER¤ lEtvER¥ ER¦ lEtwER§ ER¨ lEtxER© ERª lEtyER6ER« ER¬ llEtzER­ ER® lEt{ER¯ ER° lEt|ER± ER² lEt}ER³ ER´ lEt~ERµ ER¶ lEtER·ER¸ER¹ERºER»ER¼ER½ER¾ER¿ERÀERÁERÂERÃERÂ/Et€E]€EP                  4        U Uu/ uF  w  rWbK	  	  upp Et‚]EP                  ! ERÄ4      Et„ERÅ ERÆ lEt…ERÇ ERÈ lEt†ERÉ ERÊ lEt‡ERË ERÌ lEtˆ]Pô                  ERÍ ERÎ l4       Et‰]EPˆ                   ! ERÏ ERÐ4      4       EtŠ/ Et‹ERÑ]½ERÒ&   ERÓ ERÔ lEtŒ]F! 4       Et�ERÕ]½ERÖ&   ER× ERØ lEtŽER6ERÙ ERÚ llEt�ERÛ ERÜ lEt�ERÝ ERÞ lEt‘]0! ERß4      Et’]0! ERà4      Et“ ! ERá ERâ]E]’E]“3,          4      Et”].! RFERã7      ER6ERäRF/ERå ERæ lll4       Et•ERç ERè lEt– ! ERé ERê]EP°                  4      Et—]Pô                  ERë ERì l4       Et˜ERí ERî lEt™ERï ERð lEtšERñ ERò lEt›ERó ERô lEtœERõ ERö lEt�ER÷ ERø lEtžERù ERú lEtŸERû ERü lEt ERý ERþ lEt¡ER?Et¢ERÿ ER  lEt£ER ER lEt¤ER ER lEt¥ER<ER ER llEt¦ER ER lEt§ER	 ER
 lEt¨ER ER lEt©ER ER lEtªER ER lEt«ER ER lEt¬]EP                  ER ER l4       Et­]EPˆ                  ! RFRG7       ! ER ER4      4       Et®]ER]$3,          Et¯]E]¯E]®.E]¯3,          Et° ! ER ER4      Et±E]±! 4       Et²ER ER lEt³ER ER lEt´ER ER lEtµER  ER! lEt¶ER" ER# lEt·]F! . ER@O4      Et¸ER$ ER% lEt¹]"ER& ER' l4       EtºER( ER) lEt»ER* ER+ lEt¼ER, ER- lEt½ER;ER. ER/ llEt¾R# u upp i (A  é    )ÚannotationsN)ÚCallableÚ
CollectionÚ	GeneratorÚIteratorÚMappingÚMutableMappingÚ
MutableSet)Údatetime)Ú	lru_cache)	ÚAnyÚcastÚConcatenateÚGenericÚLiteralÚProtocolÚTYPE_CHECKINGÚ	TypeAliasÚ	TypeGuard)Údataclass_transformÚ	ParamSpecÚTypeVar)Úmock)Údatasheet_tops)ÚDeviceProperties)Ú_needs_inductor_compile)Údtype_abbrs)ÚDeferredLineBaseÚIndentedBufferÚLineContextÚValueWithLineMap)Ú
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ÚxpuÚTc               ó   € V ^8„  d   QhRR/# ©é   ÚreturnÚstr© )Úformats   "ÚV/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torch/_inductor/utils.pyÚ__annotate__rU   r   s   € ÷ ñ �cñ ó    c                 ó"  € \          U u. uF+  p \        \        V 4      P                  4       '       g   K)  V NK-  	  pp \	        V4      ^8:  g   \        R\	        V4       RV 24      h\	        V4      ^ 8X  d   RpV# VP                  4       pV# u up i )é   z+Expected at most 1 available GPU type, got ú: rJ   )Ú	GPU_TYPESÚgetattrÚtorchÚis_availableÚlenÚAssertionErrorÚpop)ÚxÚ
avail_gpusÚgpu_types      rT   Úget_gpu_typerd   q   sˆ   € å&ÓK™Y˜¬'´%¸Ó*;×*HÑ*H×*J—!�!™Y€JÐKÜˆz‹?˜aÔÜØ9¼#¸j»/Ð9JÈ"ÈZÈLÐYó
ð 	
ô ˜Z›¨AÔ-ˆv€HØ€Oð 4>·>±>Ó3C€HØ€Oùò Ls
   ‰&B´B)Ú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.cubinÚhipz.hsacoÚ.zOrderedSet[torch.dtype]Ú_TMA_SUPPORTED_DTYPESzmust be power of 2c               ó    € V ^8„  d   QhRRRR/# )rO   ÚnbytesÚintrP   rR   )rS   s   "rT   rU   rU   É   s   € ÷ 5ñ 5�3ð 5˜3ñ 5rV   c                óB   € V \         ,           ^,
          \         ) ,          # )z/Round up to the nearest multiple of ALIGN_BYTES)ÚALIGN_BYTES)r‚   s   &rT   Ú_alignr†   É   s   € à”[Õ  1Õ$¬¨Õ4Ð4rV   c               ó    € V ^8„  d   QhRRRR/# )rO   Úvú
sympy.ExprrP   ÚboolrR   )rS   s   "rT   rU   rU   Î   s   € ÷ Lñ L�:ð L $ñ LrV   c                ó&  € \        V \        P                  \        P                  \        34      '       d$   \	        \        \        V P                  4      4      # \        V \        4      ;'       g#    \        P                  ! V \        4      \        8H  # )z:v can be statically proven to be a multiple of ALIGN_BYTES)Ú
isinstanceÚsympyÚAddro   ÚallÚmapÚ_is_alignedÚargsÚalignÚgcdr…   )rˆ   s   &rT   r‘   r‘   Î   sY   € ä�!”e—i‘i¤§¡¬CÐ0×1Ò1Ü”3”{ A§F¡FÓ+Ó,Ð,Ü�aœÓ×KÐK¤5§9¢9¨Q´Ó#<ÄÑ#KÐKrV   c                  ó8   € ] tR t^ÕtRtRtRt]R R l4       tRt	R# )r“   z<Symbolically round up to the nearest multiple of ALIGN_BYTESTc               ó    € V ^8„  d   QhRRRR/# )rO   Úvaluer‰   rP   zsympy.Expr | NonerR   )rS   s   "rT   rU   Úalign.__annotate__Ü   s   € ÷ ñ ˜ð Ð(9ñ rV   c                	ó    € \        V\        \        P                  34      '       d   \	        \        V4      4      # \        V4      '       d   V# R # ©N)rŒ   rƒ   r�   ÚIntegerr†   r‘   )Úclsr—   s   &&rT   ÚevalÚ
align.evalÛ   s<   € ä�eœc¤5§=¡=Ð1×2Ò2Üœ#˜e›*Ó%Ð%Ü�u×ÒØˆLñ rV   rR   N©rX   )
Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__ÚnargsÚ
is_integerÚclassmethodr�   Ú__static_attributes__rR   rV   rT   r“   r“   Õ   s!   † ÙFà€EØ€Jàôó ôrV   r“   T)Úfrozenc                  óD   € ] tR t^ãt$ RtR]R&   R]R&   R]R&   R]R	&   R
tR# )ÚGraphPartitionMapzH
Mapping from the partition info (e.g., input/output) to the graph info
rƒ   Úidzlist[int | None]Úinput_index_mappingzlist[list[int]]Úoutput_index_mappingú	list[str]Úconstant_namesrR   N©r    r¡   r¢   r£   r¤   Ú__annotations__r¨   rR   rV   rT   r«   r«   ã   s&   ‡ ñð
 	ƒGð *Ó)ð *Ó)ð ×rV   r«   c               ó(   € V ^8„  d   QhRRRRRRRR/# )rO   ÚfnúCallable[[], Any]Úwarmuprƒ   ÚreprP   ÚfloatrR   )rS   s   "rT   rU   rU   ø   s/   € ÷ Gñ GÐ#ð G¨Sð G¸Cð GÈ%ñ GrV   c                ó2	  € V ! 4        \         P                  P                  4        \         P                  ! \	        R4      \         P
                  RR7      p\         P                  P                  RR7      p\         P                  P                  RR7      pVP                  4        \        ^4       F  pVP                  4        V ! 4        K  	  VP                  4        \         P                  P                  4        VP                  V4      ^,          p\        ^\	        W,          4      4      p\        ^\	        W',          4      4      p	\        V4       F
  pV ! 4        K  	  \        V	4       Uu. uF#  p\         P                  P                  RR7      NK%  	  pp\        V	4       Uu. uF#  p\         P                  P                  RR7      NK%  	  pp\         P                  P                  \         P                  P                  P                  .R7      ;_uu_ 4       p
\         P                  P                  4        \        V	4       F€  pVP                  4        WK,          P                  4        \         P                  P                   P                  R4      ;_uu_ 4        V ! 4        RRR4       W[,          P                  4        K‚  	  \         P                  P                  4        \         P"                  ! \%        WE4       UUu. uF  w  rÍVP                  V4      NK  	  upp4      pRRR4       \         P&                  ! X4      P)                  4       p\*        P-                  R4       \*        P-                  X
P/                  4       P1                  R	RR
7      4       \3        V
P5                  4        Uu. uFK  pVP6                  \8        P                  8X  g   K$  \:        P<                  ! RVP>                  4      f   KI  VNKM  	  up4      pV'       d,   V\@        P&                  ! R V 4       4      R,          ,          p\*        P-                  RV4       V# u upi u upi   + '       g   i     ELÄ; iu uppi   + '       g   i     EL\; iu upi )a:  
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.
rJ   ©ÚdtypeÚdeviceT©Úenable_timing©Ú
activitiesÚRunCudaModuleNú
raw eventsÚself_device_time_total©Úsort_byÚ	row_limitzfused_abs_max_\dc              3  ó8   "  € T F  qP                   x € K  	  R # 5irš   )Údevice_time_total)Ú.0Úevents   & rT   Ú	<genexpr>Úfp8_bench.<locals>.<genexpr>:  s   é € ÐQÁ¸×3Ö3Ãùó   ‚ç     @�@úprofiling results: %s msç    €„ŽAéÿÿÿÿ)!r\   rJ   ÚsynchronizeÚemptyrƒ   Úfloat16ÚEventÚrecordÚrangeÚzero_Úelapsed_timeÚmaxÚprofilerÚprofileÚProfilerActivityÚCUDAÚnvtxÚtensorÚzipÚmeanÚitemÚlogÚdebugÚkey_averagesÚtablerh   ÚeventsÚdevice_typerg   ÚreÚmatchÚnameÚ
statistics)r´   r¶   r·   ÚcacheÚstart_eventÚ	end_eventÚ_Úestimate_msÚn_warmupÚn_repeatÚpÚiÚsÚeÚtimesÚresrÊ   Úfiltered_eventss   &&&               rT   Ú	fp8_benchrü   ø   sP  € ñ „DÜ	‡J�J×ÑÔÜ�KŠKœ˜J›¬u¯}©}ÀVÔL€Eô —*‘*×"Ñ"°Ð"Ó6€KÜ—
‘
× Ñ ¨tÐ Ó4€IØ×ÑÔÜ�1ŽXˆØ�‰ŒÙ
Žñ ð ×ÑÔÜ	‡J�J×ÑÔØ×*Ñ*¨9Ó5¸Õ9€Kô �1”c˜&Õ.Ó/Ó0€HÜ�1”c˜#Õ+Ó,Ó-€Hô �8Ž_ˆÙ
Žñ ô BGÀxÄÓQÁ¸A”5—:‘:×#Ñ#°$Ð#Ö7Á€KÐQÜ?DÀX¼ÓO¹¸!”—‘×!Ñ!°Ð!Ö5¹€IÐOÜ	�‰×	Ñ	ä�N‰N×+Ñ+×0Ñ0ð
ð 
 ÷ 
ô 
ð 
Ü�
‰
×ÑÔ Ü�x–ˆAØ�K‰KŒMØ�N×!Ñ!Ô#Ü—‘—‘×&Ñ& ×7Õ7Ù”÷ 8à�L×ÑÖ!ñ !ô 	�
‰
×ÑÔ Ü—’Ü+.¨{Ô+FÔGÑ+F¡4 1ˆQ�^‰^˜AÖÑ+FÒGó
ˆ÷
ô" �*Š*�UÓ
×
 Ñ
 Ó
"€CÜ‡I�IˆlÔÜ‡I�Iˆa�n‰nÓ×$Ñ$Ð-EÐQSÐ$ÓTÔUÜð Ÿ™œó	
á#�à×!Ñ!¤Z§_¡_Ñ4ô ô —H’HÐ0°%·*±*Ó=÷	 ˆEÙ#ñ	
ó	€O÷ ØÜ�OŠOÑQÁÓQÓQØõõ	
ˆô
 ‡I�IÐ(¨#Ô.Ø€JùòO RùÚO÷ 8×7Ð7üó
 H÷
÷ 
ð 
üò*	
sP   Å1)QÆ))Q!ÈBR Ê#Q&Ê+AR Ì	Q:Ì%R Î9!RÏ RÐRÑ&Q7Ñ1R Ò R	Fc          
     ó,   € V ^8„  d   QhRRRRRRRRRR	/# ©
rO   r´   rµ   r¶   rƒ   r·   Úis_vetted_benchmarkingrŠ   rP   r¸   rR   )rS   s   "rT   rU   rU   B  s:   € ÷ ñ Øðàðð 
ðð !ð	ð
 ñrV   c                óH   € ^ RI HpHp V! \        4      pV! V4      ! WW#4      # )r   )Úgpu_benchmark_lockÚmay_distort_benchmarking_result)Ú$torch._inductor.runtime.benchmarkingr  r  Ú_do_bench_using_profiling)r´   r¶   r·   rÿ   r  r  Úlocked_benchs   &&&&   rT   Údo_bench_using_profilingr  B  s+   € ÷"ñ
 &Ô&?Ó@€LÙ*¨<Ô8Ø
�Cóð rV   c          
     ó,   € V ^8„  d   QhRRRRRRRRRR	/# )
rO   Úkineto_eventsúIterable[Any]Úprofiler_eventsrô   rƒ   Úexpected_device_typerg   rP   r¸   rR   )rS   s   "rT   rU   rU   ^  s:   € ÷ 6.ñ 6.Ø ð6.à"ð6.ð ð6.ð %ð	6.ð
 ñ6.rV   c           
     ó  a
a€ \        4       o
R  V
V3R lloV Uu. uF=  pVP                  \        8X  g   K  VP                  \        P
                  8X  g   K;  VNK?  	  pp\        V4      V8w  d#   \        RV R\         R\        V4       R24      hV F  pS! V4       K  	  RpV  Fµ  pVP                  4       pVP                  4       pVP                  4       p	VP                  4       V8X  g   KJ  V	R8w  g   KS  VS
9   g   V^ 8X  g   Kc  VS
9   g   Kl  VP                  4       R8w  g   Kƒ  WdP                  4       VP                  4       ,
          R	,          ,          pK·  	  V^ 8:  d   \        R
\         R24      hVR	,          V,          # u upi )c               ó    € V ^8„  d   QhRRRR/# )rO   rÊ   r   rP   ÚNonerR   )rS   s   "rT   rU   Ú2_get_do_bench_profile_result.<locals>.__annotate__f  s   € ÷ )ñ ) Sð )¨Tñ )rV   c                ó¶   <€ V P                   \        P                  8w  d   R # SP                  V P                  4       V P
                   F  pS! V4       K  	  R # rš   )ré   rg   ÚCPUÚaddr¬   Úcpu_children)rÊ   ÚchildÚbenchmark_event_idsÚcollect_cpu_event_idss   & €€rT   r  Ú;_get_do_bench_profile_result.<locals>.collect_cpu_event_idsf  sB   ø€ Ø×Ñ¤
§¡Ô.Ùà×Ñ §¡Ô)Ø×'Ô'ˆEÙ! %Ö(ó (rV   z	Expected Ú z profiling events. Found z events.ç        Úgpu_user_annotationzContext SyncrÎ   z$Failed to capture device events for r   )r"   rì   Ú_DO_BENCH_PROFILE_EVENT_NAMEré   rg   r  r^   ÚRuntimeErrorÚlinked_correlation_idÚcorrelation_idÚactivity_typeÚend_nsÚstart_ns)r  r
  rô   r  rÊ   Úbenchmark_eventsÚdevice_time_usr  r  r  r  r  s   &&&&      @@rT   Ú_get_do_bench_profile_resultr$  ^  s—  ù€ ô ,6«<Ð÷)ñ )ñ %óá$ˆEØ�:‰:Ô5Ñ5ô 	ð ×Ñ¤§¡Ñ/÷ 	ˆÙ$ð ð ô ÐÓ Ô(ÜØ˜�z Ô#?Ð"@ð AÜÐ)Ó*Ð+¨8ð5ó
ð 	
ó
 "ˆÙ˜eÖ$ñ "ð €NÛˆØ %× ;Ñ ;Ó =ÐØ×-Ñ-Ó/ˆØ×+Ñ+Ó-ˆà×ÑÓÐ#7Ö7ØÐ!6Ö6à%Ð)<Ô<à)¨QÖ.°>ÐEXÖ3Xð —
‘
“ Ö.àŸ|™|›~°·±Ó0@Õ@ÀFÕJÕJŠNñ ð" ˜ÔÜØ2Ô3OÐ2PÐPQÐRó
ð 	
ð ˜FÕ" XÕ-Ð-ùòMs   ›F
·F
ÁF
c          
     ó,   € V ^8„  d   QhRRRRRRRRRR	/# rþ   rR   )rS   s   "rT   rU   rU   —  sA   € ÷ Nñ NØðNàðNð 
ðNð !ð	Nð
 ñNrV   c           
     óà  € V'       g   ^ RI Hp V! 4        \        4       pVP                  4       p\	        V4      pV ! 4        VP                  4        \        P                  ! \        R4      \        P                  VR7      pVP                  RR7      p	VP                  RR7      p
V	P                  4        \        ^4       F  pVP                  4        V ! 4        K  	  V
P                  4        VP                  4        V	P                  V
4      ^,          p\        ^\        W,          4      4      p\        ^\        W,,          4      4      p\        V4       F
  pV ! 4        K  	  VP                  4        \        \        P                   P"                  V4      p\        P                   P%                  \        P                   P"                  P&                  V.R7      ;_uu_ 4       p\        V4       FN  pVP                  4        \        P                   P)                  \*        4      ;_uu_ 4        V ! 4        RRR4       KP  	  VP                  4        RRR4       \,        P/                  R4       \,        P1                  \2        P4                  4      '       d5   \,        P/                  XP7                  4       P9                  RRR	7      4       \;        XP                   P<                  P?                  4       VP?                  4       V\        \@        V4      4      p\,        P/                  R
4       \,        P/                  RV4       V#   + '       g   i     EKg  ; i  + '       g   i     EL; i)a*  
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.
)Úmay_ban_benchmarkingrº   Tr½   r¿   NrÂ   rÃ   rÄ   zprofiling time breakdownrÏ   rÐ   rÑ   )!r  r'  rd   Úupperre   rÒ   r\   rÓ   rƒ   rÕ   rÖ   r×   rØ   rÙ   rÚ   r[   rÛ   rÝ   rÜ   r  Úrecord_functionr  rä   rå   ÚisEnabledForÚloggingÚDEBUGræ   rç   r$  Úkineto_resultsrè   rg   )r´   r¶   r·   rÿ   r'  ré   Údevice_type_upperÚdevice_interfacerî   rï   rð   rñ   rò   ró   rô   Úprofile_activityrõ   Úresults   &&&&              rT   r  r  —  s�  € ÷ "ÝMáÔä“.€KØ#×)Ñ)Ó+ÐÜ/°Ó<ÐÙ„DØ× Ñ Ô"Ü�KŠKœ˜J›¬u¯y©yÀÔM€Eð #×(Ñ(°tÐ(Ó<€KØ ×&Ñ&°TÐ&Ó:€IØ×ÑÔÜ�1ŽXˆØ�‰ŒÙ
Žñ ð ×ÑÔØ× Ñ Ô"Ø×*Ñ*¨9Ó5¸Õ9€Kô �1”c˜&Õ.Ó/Ó0€HÜ�1”c˜#Õ+Ó,Ó-€Hô �8Ž_ˆÙ
Žñ ð × Ñ Ô"ÜœuŸ~™~×>Ñ>Ð@QÓRÐÜ	�‰×	Ñ	ä�N‰N×+Ñ+×/Ñ/Øð
ð 
 ÷ 
ô 
ð
 
ä�x–ˆAà�K‰KŒMä—‘×/Ñ/Ô0L×MÕMÙ”÷ NÑMñ	 !ð 	×$Ñ$Ô&÷
ô  ‡I�IˆlÔÜ
×ÑœŸ™×&Ò&Ü�	‰	Ø�N‰NÓ×"Ñ"Ð+CÈrÐ"ÓRô	
ô *Ø	�
‰
×!Ñ!×(Ñ(Ó*Ø	�‰‹
ØÜ”
Ð-Ó.ó	€Fô ‡I�IÐ(Ô)Ü‡I�IÐ(¨&Ô1Ø€M÷) N×MÑMú÷
÷ 
ð 
ús%   Ç&A
MÈ0MÈ8MÍMÍ
MÍM-	c               ó   € V ^8„  d   QhRR/# ©rO   rP   rŠ   rR   )rS   s   "rT   rU   rU   é  s   € ÷ ñ  4ñ rV   c                 ó<  €  ^ RI Hp  \        P                  P	                  RR4       V RJ;'       d&    \        \        \        P                  RR4      R4      #   \         d     R# \         d+   pR\        T4      9  d   \        R	T 24      Th Rp?R# Rp?ii ; i)
r   )Ú	roi_alignztorchvision::nmsÚMetaNÚtorchvisionr5  Fztorchvision::nms does not existzUnexpected RuntimeError: )Útorchvision.opsr5  r\   Ú_CÚ%_dispatch_has_kernel_for_dispatch_keyÚhasattrr[   ÚopsÚImportErrorr  rQ   r_   )r5  rø   s     rT   Úhas_torchvision_roi_alignr>  è  s�   € ðÝ-ä�‰×6Ñ6Ð7IÈ6ÔRØ Ð$÷ 
ð 
¬Ü”E—I‘I˜}¨dÓ3°[ó*
ð 	
øô ô ÚÜô Ø,´C¸³FÔ:Ü Ð#<¸Q¸CÐ!@ÓAÀqÐHÝûðús(   ‚0A ³%A ÁBÁ'BÁ0BÁ1BÂBc               ó    € V ^8„  d   QhRRRR/# )rO   r¼   ztorch.device | None | strrP   útorch.devicerR   )rS   s   "rT   rU   rU   ù  s   € ÷ ñ Ð3ð ¸ñ rV   c                ó‚  € V f!   \         P                  ! R4      P                  # \        V \        4      '       d   \         P                  ! V 4      p V P
                  R9  d^   V P                  fP   \        V P
                  4      p\         P                  ! V P
                  VP                  P                  4       R7      # V # )Nr  )Úindex)ÚcpuÚmeta)
r\   rà   r¼   rŒ   rQ   ÚtyperB  re   ÚWorkerÚcurrent_device©r¼   r/  s   & rT   Údecode_devicerI  ù  s…   € Ø‚~Ü�|Š|˜CÓ ×'Ñ'Ð'Ü�&œ#×ÒÜ—’˜fÓ%ˆØ‡{�{˜/Ô)¨f¯l©lÒ.BÜ3°F·K±KÓ@ÐÜ�|Š|˜FŸK™KÐ/?×/FÑ/F×/UÑ/UÓ/WÔXÐXØ€MrV   c               ó    € V ^8„  d   QhRRRR/# )rO   ÚitzIterable[sympy.Expr]rP   r‰   rR   )rS   s   "rT   rU   rU     s   € ÷ ;ñ ;Ð*ð ;¨zñ ;rV   c                ó~   € \         P                  ! \        P                  V \        P
                  P                  4      # rš   )Ú	functoolsÚreduceÚoperatorÚmulr�   ÚSÚOne)rK  s   &rT   Úsympy_productrS    s#   € Ü×ÒœHŸL™L¨"¬e¯g©g¯k©kÓ:Ð:rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   Úseq1úSequence[sympy.Expr]Úseq2rP   r‰   rR   )rS   s   "rT   rU   rU     s(   € ÷ @ñ @Ð(ð @Ð0Dð @Èñ @rV   c           	     óØ   € \        V 4      \        V4      8X  g$   \        R \        V 4       R\        V4       24      h\        P                  ! \	        R \        W4       4       4      4      # )zLength mismatch: len(seq1)=z, len(seq2)=c              3  ó6   "  € T F  w  rW,          x € K  	  R # 5irš   rR   )rÉ   ÚaÚbs   &  rT   rË   Úsympy_dot.<locals>.<genexpr>  s   é € Ð>©o¡d a˜AŸEšE«oùs   ‚)r^   r_   r�   ÚexpandÚsumrá   )rU  rW  s   &&rT   Ú	sympy_dotr_    sW   € Üˆt‹9œ˜D›	Ô!ÜØ)¬#¨d«)¨°LÄÀTÃÀÐLó
ð 	
ô �<Š<œÑ>¬c°$¬oÓ>Ó>Ó?Ð?rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   ÚindicesrV  ÚsizesrP   r‰   rR   )rS   s   "rT   rU   rU     s&   € ÷ ñ Ø!ðàðð ñrV   c                óø   € \        V 4      \        V4      8w  d$   \        R\        V 4       R\        V4       24      h\        P                  P                  p\        W4       F  w  r4W$,          V,           pK  	  V# )z7Row-major flatten: per-dimension indices -> flat index.z.indices and sizes must have equal length, got z and )r^   r_   r�   rQ  ÚZerorá   )ra  rb  ÚflatrB  Úsizes   &&   rT   Úflatten_indexrg    sm   € ô
 ˆ7ƒ|”s˜5“zÔ!ÜØ<Ü�7‹|ˆn˜E¤# e£* ð/ó
ð 	
ô �7‰7�<‰<€DÜ˜7Ö*‰ˆØ�{˜UÕ"Šñ +à€KrV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   rB  r‰   rb  rV  rP   úlist[sympy.Expr]rR   )rS   s   "rT   rU   rU      s&   € ÷ ñ Øðàðð ñrV   c                ó†   € \        V4       UUu. uF$  w  r#\        V \        W^,           R 4      V4      NK&  	  upp# u uppi )z=Row-major decomposition: flat index -> per-dimension indices.N)Ú	enumeraterp   rS  )rB  rb  rö   rf  s   &&  rT   Údecompose_indexrl     sF   € ô ! Ô'ôá'‰GˆAô 	˜œ}¨U°qµ5°7¨^Ó<¸dÖCÙ'òð ùó s   �*=c               ó    € V ^8„  d   QhRRRR/# )rO   rK  zIterable[_T]rP   zValuesView[_T]rR   )rS   s   "rT   rU   rU   +  s   € ÷ +ñ +ˆ|ð + ñ +rV   c                ó`   € V  Uu/ uF  p\        V4      VbK  	  upP                  4       # u upi rš   )r¬   Úvalues)rK  ra   s   & rT   Úuniquerp  +  s+   € Ù Ó!™b˜ŒBˆq‹E�1ŠH™bÑ!×(Ñ(Ó*Ð*ùÒ!s   …+c               ó$   € V ^8„  d   QhRRRRRR/# )rO   Únumberúint | sympy.ExprÚdenomrP   rR   )rS   s   "rT   rU   rU   /  s$   € ÷ *ñ *Ð$ð *Ð-=ð *ÐBRñ *rV   c                ó¨  € \        V \        P                  4      '       g!   \        V\        P                  4      '       d5   \        \        P                  ! V 4      \        P                  ! V4      4      # \        V \
        4      '       d   \        V\
        4      '       g)   \        V  R \        V 4       RV R \        V4       24      h\        W4      # )rY   ú, )	rŒ   r�   ÚExprrk   Úsympifyrƒ   r_   rE  Úruntime_ceildiv)rr  rt  s   &&rT   rw   rw   /  s’   € Ü�&œ%Ÿ*™*×%Ò%¬°E¼5¿:¹:×)FÒ)FÜ”u—}’} VÓ,¬e¯mªm¸EÓ.BÓCÐCô �vœs×#Ò#¬
°5¼#×(>Ò(>Ü ˜x r¬$¨v«,¨°r¸%¸ÀÄ4ÈÃ;À-ÐPÓQÐQÜ˜6Ó)Ð)rV   c               ó    € V ^8„  d   QhRRRR/# )rO   Úkeyútorch.dtype | NonerP   rQ   rR   )rS   s   "rT   rU   rU   :  s   € ÷ %Añ %AÐ$ð %A¨ñ %ArV   c                ó–  € V f   R# \        V 4      P                  R4      R.,          p/ RRbRRbRRbR	R
bRRbRRbRRbRRbRRbRRbRRbRRbRRbRRbRRbRRbR R!bR"R#R$R%R&RR'R(R)R*R+R,/CpTP                  \        VP	                  4       4       Uu/ uF  q3VbK  	  up4       \        V \         4      '       d   V # R-W!,           2# u upi )/Nz*i8r   rŠ   Ú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Úint32Úi32Úint64Úi64Úuint8Úuint16Úu16Úuint32Úu32Úuint64Úu64Ú*rÑ   )rQ   ÚsplitÚupdateÚlistro  rŒ   )r{  Ú	dtype_strÚtysrˆ   s   &   rT   Ú_type_ofr¬  :  si  € ð
 ‚{ÙÜ�C“—‘˜sÓ# BÕ'€IðØ�ðà�iðð 	�Gðð 	�zð	ð
 	˜ðð 	˜ðð 	�wðð 	˜9ðð 	˜:ðð 	˜$ðð 	˜Dðð 	�6ðð 	�Fðð  	�6ð!ð" 	�6ð#ð$ 	�ð%ð& 	�ð'ð( 	�Ø�Ø�Ø�%Ø�%Ø�%ñ3€Cð8 ‡J�Jœd 3§:¡:£<Ô0Ó1Ñ0˜�1’Ñ0Ñ1Ô2Ü˜S¤#×&Ò&ˆ3Ð@¨a°µÐ/?Ð,@Ð@ùò 2s   ÂCc               ó    € V ^8„  d   QhRRRR/# )rO   ÚlstzIterable[int | torch.SymInt]rP   ri  rR   )rS   s   "rT   rU   rU   b  s   € ÷ +ñ +Ø	%ð+àñ+rV   c                óX   € V  Uu. uF  p\         P                  ! V4      NK  	  up# u upi )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�   rx  ©r®  rö   s   & rT   Úconvert_shape_to_inductorr±  b  s%   € ñ '*Ó*¡c ŒE�MŠM˜!Ö¡cÑ*Ð*ùÒ*s   …'c               ó    € V ^8„  d   QhRRRR/# )rO   Úvalúint | torch.SymIntrP   rs  rR   )rS   s   "rT   rU   rU   m  s   € ÷ 	ñ 	Ð 2ð 	Ð7Gñ 	rV   c                ór   € \        V \        P                  4      '       d   V P                  P                  # V # )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Œ   r\   r5   ÚnodeÚexpr©r³  s   &rT   Úconvert_symint_to_exprr¹  m  s(   € ô �#”u—|‘|×$Ò$Ø�x‰x�}‰}ÐØ€JrV   c               ó    € V ^8„  d   QhRRRR/# )rO   rö   rs  rP   r´  rR   )rS   s   "rT   rU   rU   y  s   € ÷ ñ Ð)ð Ð.@ñ rV   c                óö   € ^RI Hp \        V \        4      '       d   V # \        V \        P
                  4      '       d   \        V 4      # VP                  P                  P                  P                  V RR7      # )zD
Like convert_shape_to_symint, but operates on a single expression.
©ÚVN)Úhint)
Úvirtualizedr½  rŒ   rƒ   r�   r›   ÚgraphÚsizevarsÚ	shape_envÚcreate_symintnode)rö   r½  s   & rT   Úconvert_to_symintrÄ  y  sk   € õ ô �aœ×Òð 	
ðô
 ˜!œUŸ]™]×+Ò+ô �‹Fð	ð —‘×!Ñ!×+Ñ+×=Ñ=¸aÀdÐ=ÓKðrV   c               ó    € V ^8„  d   QhRRRR/# )rO   r®  zIterable[int | sympy.Expr]rP   zlist[int | torch.SymInt]rR   )rS   s   "rT   rU   rU   Š  s   € ÷ /ñ /Ø	#ð/àñ/rV   c                óB   € V  Uu. uF  p\        V4      NK  	  up# u upi )zn
Takes a list of shapes from Inductor and converts them into symints (or just
ints if all shapes are static).
)rÄ  r°  s   & rT   Úconvert_shape_to_symintrÇ  Š  s"   € ñ +.Ó.©# QÔ˜aÖ ©#Ñ.Ð.ùÒ.s   …c               ó    € V ^8„  d   QhRRRR/# )rO   Úopútorch._ops.OpOverloadrP   rŠ   rR   )rS   s   "rT   rU   rU   ”  s   € ÷ Gñ GÐ%ð G¨$ñ GrV   c                óº   € \         ;QJ d0    R V P                  P                   4       F  '       g   K   R# 	  R# ! R V P                  P                   4       4      # )z%
Does this op overload have aliasing
c              3  ó<   "  € T F  qP                   R Jx € K  	  R # 5irš   )Ú
alias_info©rÉ   rZ  s   & rT   rË   Úis_view.<locals>.<genexpr>˜  s   é € ÐFÑ1E¨A�|‰| 4Õ'Ó1Eùs   ‚TF)ÚanyÚ_schemaÚ	arguments©rÉ  s   &rT   Úis_viewrÔ  ”  sA   € ÷ ‹3ÑF°·±×1EÒ1EÓF�3Œ3ÐFŠ3ÐFˆ3ÑF°·±×1EÒ1EÓFÓFÐFrV   c                ó   € R # ©FrR   )rñ   s   &rT   Ú<lambda>r×  �  s   € ÉrV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   Úuser8   Úis_pointwise_fnz'Callable[[torch._ops.OpOverload], bool]rP   rŠ   rR   )rS   s   "rT   rU   rU   ›  s+   € ÷ Iñ IØ	ðIà<ðIð 
ñIrV   c                ó‚  a€ V P                   R8w  d   R# \        V P                  \        P                  P
                  4      '       g!   V P                  \        P                  J g   R# \        \        P                  P
                  V P                  4      pV\        P                  J g   \        V4      '       dO   \        ;QJ d)    V3R lV P                   4       F  '       d   K   R# 	  R# ! V3R lV P                   4       4      # \        P                  P                  VP                  9   ;'       g	    S! V4      # )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  ó<   <"  € T F  p\        VS4      x € K  	  R # 5irš   )Úis_pointwise_use)rÉ   ÚurÚ  s   & €rT   rË   Ú#is_pointwise_use.<locals>.<genexpr>®  s   øé € ÐKÁ¸AÔ# A ×7Ð7Ãùs   ƒT)rÉ  rŒ   Útargetr\   Ú_opsÚ
OpOverloadrO  Úgetitemr   rÔ  r�   ÚusersÚTagÚ	pointwiseÚtags)rÙ  rÚ  rá  s   &f rT   rÞ  rÞ  ›  sÊ   ø€ ð ‡v�v�Ô Ùä�3—:‘:œuŸz™z×4Ñ4×5Ò5¸¿¹Äx×GWÑGWÓ9Wáä”%—*‘*×'Ñ'¨¯©Ó4€FØ”×!Ñ!Ó!¤W¨V§_¢_ß‹sÔKÀÇÂÓK�sŒsÐKŠsÐKˆsÔKÀÇÂÓKÓKÐKä�9‰9×Ñ &§+¡+Ñ-×HÐH±ÀÓ1HÐHrV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )	rO   rá  r   r’   ú	list[Any]Úkwargsúdict[str, Any]rP   z&tuple[GraphModule, list[torch.Tensor]]rR   )rS   s   "rT   rU   rU   ³  s,   € ÷ ñ ØðØ ðØ*8ðà+ñrV   c           	     óØ  aa€ \         P                  P                  4       o. oR  VV3R llpSP                  ! V .\	        \         P
                  W1V34      O5!  p\        V P                  P                  4      ^8X  d9   \        V P                  P                  ^ ,          P                  4      R8X  d   V3pSP                  V4       \         P                  P                  / S4      pVS3# )c               ó    € V ^8„  d   QhRRRR/# )rO   Úargútorch.TensorrP   r8   rR   )rS   s   "rT   rU   Ú'gen_gm_and_inputs.<locals>.__annotate__¹  s   € ÷ 6ñ 6˜Lð 6¨Tñ 6rV   c                ó`   <€ SP                  V 4       SP                  R \        S4       24      # )rï  )ÚappendÚplaceholderr^   )rï  ÚgÚ
graph_argss   &€€rT   Úadd_tensor_argÚ)gen_gm_and_inputs.<locals>.add_tensor_arg¹  s,   ø€ Ø×Ñ˜#ÔØ�}‰}˜s¤3 z£?Ð"3Ð4Ó5Ð5rV   ÚTensor)r\   ÚfxÚGraphrÜ  r$   rù  r^   rÑ  ÚreturnsrQ   rE  Úoutputr7   )rá  r’   rë  r÷  r¶  Úgmrõ  rö  s   &&&   @@rT   Úgen_gm_and_inputsrÿ  ³  s»   ù€ ô 	�‰�‰Ó€AØ%'€J÷6ñ 6ð �?Š?ØðÜœuŸ|™|¨^ÀF¸^ÓLó€Dô 	ˆF�N‰N×"Ñ"Ó# qÔ(Ü�—‘×&Ñ& qÕ)×.Ñ.Ó/°8Ô;àˆwˆØ‡H�HˆT„Nä	�‰×	Ñ	˜b !Ó	$€BØˆzˆ>ÐrV   c               ó    € V ^8„  d   QhRRRR/# )rO   r¼   rQ   rP   r  rR   )rS   s   "rT   rU   rU   Ë  s   € ÷ 'ñ '˜ð '¨ñ 'rV   c                ó~   € V R 8X  d   R# \        V 4      pVP                  4       '       d   VP                  4        R# R# ©rC  N)re   r]   rÒ   rH  s   & rT   rÒ   rÒ   Ë  s7   € Ø�„ÙÜ/°Ó7ÐØ×$Ñ$×&Ò&Ø×$Ñ$Ö&ñ 'rV   c          
     ó,   € V ^8„  d   QhRRRRRRRRR	R
/# )rO   ÚmodelúCallable[..., Any]Úexample_inputsúSequence[Any]rù   rƒ   r¼   rQ   rP   r¸   rR   )rS   s   "rT   rU   rU   Ó  s:   € ÷ ñ Øðà!ðð ðð ð	ð
 ñrV   c                ó  € \        V4       \        P                  ! R 4       \        P                  ! 4       p\        V4       F  pV ! V!  p\        V4       K  	  \        P                  ! 4       pXf   \        R4      hWt,
          # )i9  z4Expected result to be not None after timed execution)rÒ   r\   Úmanual_seedÚtimeÚperf_counterr×   r_   )r  r  rù   r¼   Út0rñ   r1  Út1s   &&&&    rT   Útimedr  Ó  sp   € ô �ÔÜ	×Ò�dÔÜ	×	Ò	Ó	€BÜ�5Ž\ˆÙ˜Ñ'ˆÜ�FÖñ ô 
×	Ò	Ó	€Bà‚~ÜÐSÓTÐTØ�7€NrV   c               ó4   € V ^8„  d   QhRRRRRRRRRR	R
RRR	/# )rO   r  r  r  r  rù   rƒ   ÚrepeatÚbaseliner¸   r¼   rQ   rP   rR   )rS   s   "rT   rU   rU   æ  sN   € ÷ ñ Øðà!ðð ðð ð	ð
 ðð ðð ñrV   c                óþ   € \         P                  ! \        V4       Uu. uF  p\        WW%4      NK  	  up4      p\         P                  ! V4      V,          p\        W„,          R  4       VP                  4       # u upi )z.6f)r\   rà   r×   r  ÚmedianÚprintrã   )	r  r  rù   r  r  r¼   rñ   ÚtimingsÚtooks	   &&&&&&   rT   Úprint_performancer  æ  se   € ô �lŠlÜ>CÀF¼mÓL¹m¸Œˆu eÖ	4¹mÑLó€Gô �<Š<˜Ó  5Õ(€DÜ	ˆT�_˜SÐ!Ô#Ø�9‰9‹;Ðùò	 	Ms   žA:c               ó$   € V ^8„  d   QhRRRRRR/# )rO   Úobjr   ÚmethodrQ   rP   r  rR   )rS   s   "rT   rU   rU   ö  s!   € ÷ )ñ )˜3ð )¨ð )°ñ )rV   c                óH   a€ \        W4      ! 4       o\        WV3R l4       R# )zKReplace obj.method() with a new method that returns a precomputed constant.c                 ó   <€ S # rš   rR   )r1  s   €rT   r×  Ú#precompute_method.<locals>.<lambda>ù  s   ø€ ¡rV   N)r[   Úsetattr)r  r  r1  s   &&@rT   Úprecompute_methodr  ö  s   ø€ ä�SÔ!Ó#€FÜˆCœÖ(rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r  r   Úmethodsr¯   rP   r  rR   )rS   s   "rT   rU   rU   ü  s!   € ÷ 'ñ '˜Cð '¨)ð '¸ñ 'rV   c                ó.   € V F  p\        W4       K  	  R# )zFReplace methods with new methods that returns a precomputed constants.N)r  )r  r!  r  s   && rT   Úprecompute_methodsr#  ü  s   € ãˆÜ˜#Ö&ó rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   rZ  rƒ   r[  rP   rR   )rS   s   "rT   rU   rU     s!   € ÷ #ñ #ˆ3ð #�3ð #˜3ñ #rV   c                ó@   € \        W8„  4      \        W8  4      ,
          # rš   )rƒ   ©rZ  r[  s   &&rT   Úcmpr'    s   € Üˆq‰u‹:œ˜A™E›
Õ"Ð"rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   ra   zint | Sequence[int]rf  rƒ   rP   úSequence[int]rR   )rS   s   "rT   rU   rU     s"   € ÷ ñ Ð'ð ¨sð °}ñ rV   c                ó¦   € \        V \        4      '       d   V .V,          # \        V 4      ^8X  d!   \        V 4      ! V ^ ,          .4      V,          # V # rŸ   )rŒ   rƒ   r^   rE  )ra   rf  s   &&rT   Úpad_listliker+    sD   € Ü�!”S×ÒØˆs�T�zÐÜ
ˆ1ƒv�„{Ü�AŒw˜˜!��v‹ Õ%Ð%Ø€HrV   c               ó    € V ^8„  d   QhRRRR/# )rO   ra   ztuple[_T, ...]rP   zlist[_T]rR   )rS   s   "rT   rU   rU     s   € ÷ $ñ $�Nð $ xñ $rV   c                óJ   € \        V 4      ^ 8X  d   . # R R lp\        WR7      # )r   c               ó    € V ^8„  d   QhRRRR/# )rO   Úelemrz   rP   rQ   rR   )rS   s   "rT   rU   Ú"tuple_sorted.<locals>.__annotate__  s   € ÷ ñ ˜ð ˜sñ rV   c                ó®   € \        V \        4      '       d   V # ^RIHp \        W4      '       g   \	        R\        V 4       24      hV P                  4       # )rX   )rH   z Expected BaseSchedulerNode, got )rŒ   rQ   Ú	schedulerrH   r_   rE  Úget_name)r/  rH   s   & rT   Ú	sort_funcÚtuple_sorted.<locals>.sort_func  sE   € Ü�dœC× Ò ØˆKå0ä˜$×2Ò2Ü Ð#CÄDÈÃJÀ<Ð!PÓQÐQØ�}‰}‹ÐrV   ©r{  )r^   Úsorted)ra   r4  s   & rT   Útuple_sortedr8    s$   € Ü
ˆ1ƒv�„{Øˆ	õô �!Ô#Ð#rV   ÚPÚRV)Ú	covariantc                  ó8   € ] tR tRt]R R l4       tR R ltRtR# )ÚCachedMethodi%  c               ó    € V ^8„  d   QhRRRR/# )rO   rî   r   rP   r  rR   )rS   s   "rT   rU   ÚCachedMethod.__annotate__'  s   € ×,Ñ,˜3Ð, 4Ñ,rV   c                	ó   € R # rš   rR   )rî   s   &rT   Úclear_cacheÚCachedMethod.clear_cache&  s   € Ù),rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r’   úP.argsrë  úP.kwargsrP   r:  rR   )rS   s   "rT   rU   r?  )  s   € ×DÑD˜fÐD°ÐD¸RÑDrV   c                	ó   € R # rš   rR   ©Úselfr’   rë  s   &*,rT   Ú__call__ÚCachedMethod.__call__)  s   € ÁrV   rR   N)r    r¡   r¢   r£   ÚstaticmethodrA  rI  r¨   rR   rV   rT   r=  r=  %  s   † ØÜ,ó Ø,çDÑDrV   r=  c               ó    € V ^8„  d   QhRRRR/# ©rO   r´   z!Callable[Concatenate[Any, P], RV]rP   zCachedMethod[P, RV]rR   )rS   s   "rT   rU   rU   -  s   € ÷ ñ Ð7ð Ð<Oñ rV   c           	     óä   a€ V P                   pR V R2oRV /p\        RV RS RS R2P                  4       V4       \        P                  ! V 4      ! W! R2,          4      pR V3R	 llpWCn        V# )
Ú__Ú_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               ó    € V ^8„  d   QhRRRR/# ©rO   rH  r   rP   r  rR   )rS   s   "rT   rU   Ú#cache_on_self.<locals>.__annotate__B  s   € ÷ ñ ˜#ð  $ñ rV   c                óH   <€ \        V S4      '       d   \        V S4       R # R # rš   ©r;  Údelattr©rH  r{  s   &€rT   rA  Ú"cache_on_self.<locals>.clear_cacheB  s   ø€ Ü�4˜×ÒÜ�D˜#Öñ rV   ©r    ÚexecÚlstriprM  ÚwrapsrA  )r´   rì   ÚctxÚwrapperrA  r{  s   &    @rT   Úcache_on_selfr`  -  s™   ø€ Ø�;‰;€DØˆtˆf�FÐ
€Cð �ˆ*€CÜðØˆFð à ˜Eð "'ð (+ eð ,	ð		÷ ‰F‹HØôô �oŠo˜bÔ! #¨¨nÐ&=Õ">Ó?€G÷ð ð &ÔØ€NrV   c               ó    € V ^8„  d   QhRRRR/# rM  rR   )rS   s   "rT   rU   rU   J  s   € ÷ ñ Ø)ðàñrV   c                ó   € \        V 4      # )zU
Variant of cache_on_self for properties. The only difference is the type signature.
)r`  )r´   s   &rT   Úcache_property_on_selfrc  J  s   € ô ˜ÓÐrV   c               ó    € V ^8„  d   QhRRRR/# )rO   Ú
class_namerQ   rP   z*Callable[[FN_TYPE[P, RV]], FN_TYPE[P, RV]]rR   )rS   s   "rT   rU   rU   T  s   € ÷ +ñ +Øð+à/ñ+rV   c                ó   a € R  V 3R llpV# )c               ó    € V ^8„  d   QhRRRR/# )rO   r´   zFN_TYPE[P, RV]rP   rR   )rS   s   "rT   rU   Ú,cache_on_self_and_args.<locals>.__annotate__Y  s   € ÷ $ñ $Øð$à	ñ$rV   c           	     óä   <a€ R S RV P                    R2oRV /p\        RS RS RS R2P                  4       V4       \        P                  ! V 4      ! VR,          4      pR	 V3R
 llpW2n        V# )rO  rñ   rP  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               ó    € V ^8„  d   QhRRRR/# rS  rR   )rS   s   "rT   rU   Ú=cache_on_self_and_args.<locals>.wrapper.<locals>.__annotate__x  s   € ÷ 	#ñ 	#˜cð 	# dñ 	#rV   c                óH   <€ \        V S4      '       d   \        V S4       R # R # rš   rV  rX  s   &€rT   rA  Ú<cache_on_self_and_args.<locals>.wrapper.<locals>.clear_cachex  s   ø€ Ü�t˜S×!Ò!Ü˜˜cÖ"ñ "rV   rZ  )r´   r^  rj  rA  r{  re  s   &   @€rT   r_  Ú'cache_on_self_and_args.<locals>.wrapperY  s–   ù€ ð �:�,˜a §¡˜}¨FÐ3ˆð �RˆjˆÜð'ð (+ eð ,/Ø/2¨eð 4à!˜Uð #ð÷$ ‘“Øô)	
ô, —’ Ô# C¨¥LÓ1ˆ÷	#ð 	#ð (ÔØˆrV   rR   )re  r_  s   f rT   Úcache_on_self_and_argsrp  T  s   ø€ ÷
$ð $ðL €NrV   c               ó    € V ^8„  d   QhRRRR/# )rO   Únode_scheduleú*Sequence[BaseSchedulerNode] | ExternKernelrP   zOrderedSet[Node]rR   )rS   s   "rT   rU   rU   ‚  s   € ÷ ñ Ø=ðàñrV   c           
     óš  € ^RI Hp \        V \        4      '       d{   \        P
                  ! \        P                  V  Uu. uFA  p\        VR4      '       g   K  VP                  '       g   K+  VP                  P                  NKC  	  up\        4       4      # \        WP                  4      '       d   V P                  # \        4       # u upi )rX   ©Úirr¶  )Ú rv  rŒ   r©  rM  rN  rO  Úor_r;  r¶  Úoriginsr"   rA   )rr  rv  r¶  s   &  rT   Úaggregate_originsrz  ‚  s    € õ ä�-¤×&Ò&Ü×ÒÜ�L‰Lñ *óñ *�DÜ˜4 ×(ô "à-1¯Y­Yô "�—	‘	×!Ô!Ù)ñô ‹Ló	
ð 		
ô 
�M§?¡?×	3Ò	3Ø×$Ñ$Ð$ä‹|Ðùòs   Á C
ÁC
Á-C
c               ó$   € V ^8„  d   QhRRRRRR/# )rO   rr  úSequence[BaseSchedulerNode]Údescriptive_namesz8Literal[True, 'torch', 'original_aten', 'inductor_node']rP   rQ   rR   )rS   s   "rT   rU   rU   ˜  s&   € ÷ <ñ <Ø.ð<àOð<ð 	ñ<rV   c                óÀ  € \        V 4      pVR 8X  dn   R pV Uu. uFH  pVP                  R8X  g   K  R VP                  9   g   K)  VP                  R ,          f   K@  V! V4      NKJ  	  pp\        \	        V4      4      pEMDVR8X  Ed   . pV Fä  pVP                  R8X  g   K  RpRpRVP                  9   d   VP                  R,          R,          pM-RVP                  9   d   VP                  R,          R,          pRpV'       g   K}  \        V^,          \        4      '       d"   VP                  V^,          V,           4       K»  VP                  V^,          P                  V,           4       Kæ  	  \        \	        V4      4      pM;VR	8X  d/   V Uu. uF!  qDP                  R8X  g   K  VP                  NK#  	  ppM\        hR
P                  R.V,           4      p\        4       '       dh   \        V4      ^28”  dX   \        P                  ! VP!                  R4      4      P#                  4       R,          p	VR,          P%                  R
4       R
V	 2pV# u upi u upi )Úoriginal_atenc                ó:  € V P                   R ,          pRp\        V\        P                  P                  4      '       d   VP
                  P                  pV# \        V\        P                  P                  4      '       d   \        VP                  4       4      pV# )r  rw  )
rD  rŒ   r\   râ  rã  Ú_overloadpacketr    ÚHigherOrderOperatorrQ   rì   )Úoriginr  r{  s   &  rT   Úget_origin_meta_strÚ2get_fused_kernel_name.<locals>.get_origin_meta_strŸ  su   € Ø"ŸK™K¨Õ8ˆMØˆCÜ˜-¬¯©×)>Ñ)>×?Ò?Ø#×3Ñ3×<Ñ<�ð ˆJô ˜M¬5¯:©:×+IÑ+I×JÒJÜ˜-×,Ñ,Ó.Ó/�ØˆJrV   rÜ  Nr\   rw  Úsource_fn_stackÚfwd_source_fn_stackÚbackwardÚinductor_noderñ   Úfusedúutf-8:Né   N:Né)   NrÑ   )rz  rÉ  rD  r7  r"   rŒ   rQ   ró  r    rì   ÚNotImplementedErrorÚjoinÚ
is_windowsr^   ÚhashlibÚsha256ÚencodeÚ	hexdigestÚrstrip)
rr  r}  Úall_originsr„  rƒ  ÚsourcesÚ	source_fnÚsuffixrì   Úhs
   &&        rT   Úget_fused_kernel_namer›  ˜  s  € ô $ MÓ2€KØ˜OÔ+ò	ñ &ó
á%�Ø�y‰y˜OÑ+ô (ð   6§;¡;Ñ.ô (ð —‘˜OÕ,ô	 (Ñ Ö'Ù%ð 	ð 
ô œ GÓ,Ó-ŠØ	˜gÕ	%àˆÛ!ˆFØ�y‰y˜OÖ+Ø �	Ø�Ø$¨¯©Ô3Ø &§¡Ð,=Õ >¸rÕ B‘IØ*¨f¯k©kÔ9à &§¡Ð,AÕ BÀ2Õ F�IØ'�Fß ÙÜ˜i¨�l¬C×0Ò0Ø—N‘N 9¨Q¥<°&Õ#8Ö9à—N‘N 9¨Q¥<×#8Ñ#8¸6Õ#AÖBñ "ô" œ GÓ,Ó-‰Ø	˜oÔ	-á&1ó
Ù&1˜F·Y±YÀ/Ñ5QŒKˆF�KŒK¡kð 	ð 
ˆô "Ð!Ø�8‰8�W�I Õ'Ó(€Dô
 ‡|‚|œ˜D›	 BœÜ�NŠN˜4Ÿ;™; wÓ/Ó0×:Ñ:Ó<¸RÕ@ˆØ�s•)×"Ñ" 3Ó'Ð(¨¨!¨Ð-ˆØ€KùòW
ùò<
s"   šI²IÁIÁIÆIÆ+Ic               ó$   € V ^8„  d   QhRRRRRR/# )rO   rr  rs  r_  r<   rP   útuple[str, str]rR   )rS   s   "rT   rU   rU   ×  s+   € ÷ h2ñ h2Ø=ðh2à!ðh2ð ñh2rV   c                óÞ  aa a!a"€ \        V 4      pV Uu. uF  q3P                  R8X  g   K  VNK  	  pp\        P                  ! \        4      p\        P                  ! \        4      pRo!V'       d�   \        R V 4       4      p\        V4      ^8X  dk   V^ ,          P                  o!\        S!R4      '       g0   \        S!P                  4       UU	u/ uF  w  r‰W˜bK	  	  p
pp	V
S!n        VP                  V!3R lR7       V EF’  pRVP                  9   d×   VP                  R,          eÂ   VP                  R,          pRp\        V\        P                   P"                  4      '       d   \%        VP&                  4      pMC\        V\        P                   P(                  4      '       d   \%        VP+                  4       4      pV'       d"   Wm,          P-                  VP*                  4       RVP                  9   dI   VP                  R,          ^ ,          P*                  pW],          P-                  VP*                  4       EKD  VP                  P/                  R	4      R
8X  g   EKg  W[P*                  ,          P-                  VP*                  4       EK•  	  S!e   RMRpSP0                   RV RRP3                  VP5                  4       4       RRP3                  VP5                  4       4       R2pSP0                   R2.p\7        VP9                  4       4       FA  w  ppVP-                  SP0                   RV RRP3                  \7        V4      4       24       KC  	  S!Eeï   ^RIHo  VP-                  SP0                   R24       \        4       p. p\        V S P>                  4      '       Egé   ^RI H!p R V 3R llpR R lo"R V"3R llpV  EFÃ  p	\        V	R4      '       d   V	PD                  f   K&  \        V	PD                  R4      '       dÒ   V	PD                  PF                  eº   V	PD                  PF                   FŸ  pVP*                  V9   d   K  VPI                  VP*                  4       VP                  PK                  VP*                  4      pVf   K\  V! VVP*                  4      w  ppVP-                  SP0                   R V R!V! V4       R"V R24       K¡  	  \        V	PD                  R#4      '       g   EK2  V	PD                  PL                  f   EKM  V	PD                  PL                   F\  pVP                  PK                  VP*                  4      pVf   K.  V! VVP*                  4      w  ppVP-                  R$V,           4       K^  	  EKÆ  	  V F„  pVPO                  R%R&7      pVeL   \        PP                  PR                  '       d,   VPU                  V3R' lVPW                  4        4       4       Kd  VP-                  SP0                   RV 24       K†  	  VP-                  SP0                   R(R)P3                  V4       24       VR*P3                  V4      3# u upi u up	pi )+a  
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  ó8   "  € T F  qP                   x € K  	  R # 5irš   )rÀ  )rÉ   Úns   & rT   rË   Ú&get_kernel_metadata.<locals>.<genexpr>ô  s   é € Ð"C±N¨q§7¦7³NùrÍ   Ú)_inductor_kernel_metadata_node_to_idx_mapc                ó*   <€ SP                   V ,          # rš   )r¢  )r   Úsingle_graphs   &€rT   r×  Ú%get_kernel_metadata.<locals>.<lambda>ü  s   ø€ ˜l×TÑTÐUVÖWrV   r6  r  Ú	from_nodeÚpartitioner_tagÚis_backwardzTopologically SortedÚUnsortedr  z Source Nodes: [rv  z], Original ATen: [Ú]z" Source node to ATen node mapping:ú   z => ru  z Graph fragment:r¼  c               ó$   € V ^8„  d   QhRRRRRR/# )rO   Úbufferz-ir.TensorBox | ir.Buffer | ir.TorchBindObjectÚrw_namerQ   rP   ztuple[str, ir.Layout | None]rR   )rS   s   "rT   rU   Ú)get_kernel_metadata.<locals>.__annotate__&  s$   € ÷ $ñ $ØEð$ØPSð$à-ñ$rV   c                óZ  <€ \        V SP                  4      '       dH   \        V P                  SP                  4      '       d"   V P                  P                  P                  pMV P                  pVf   TpMVP
                  p V P                  4       pW43#   \         d    R p Y43# i ; irš   )rŒ   Ú	TensorBoxÚdataÚ
StorageBoxÚorigin_noderì   Ú
get_layoutrŽ  )r­  r®  r´  rì   Úlayoutrv  s   &&   €rT   Úget_buffer_infoÚ,get_kernel_metadata.<locals>.get_buffer_info&  s¢   ø€ ô ˜f b§l¡l×3Ò3¼
Ø—K‘K §¡÷9ò 9ð #)§+¡+×"2Ñ"2×">Ñ">‘Kà"(×"4Ñ"4�KØÒ&à"‘Dà&×+Ñ+�Dð"Ø#×.Ñ.Ó0�Fð �|Ð#øô +ô "Ø!‘FØ�|Ð#ð"ús   ÂB ÂB*Â)B*c               ó    € V ^8„  d   QhRRRR/# )rO   ÚshapezIterable[int]rP   rQ   rR   )rS   s   "rT   rU   r¯  :  s   € ÷ Añ A }ð A¸ñ ArV   c           	     óh   € R RP                  V  Uu. uF  p\        V4      NK  	  up4       R2# u upi )Ú[rv  rª  )r�  rQ   )rº  ra   s   & rT   Ústringify_shapeÚ,get_kernel_metadata.<locals>.stringify_shape:  s1   € Ø˜4Ÿ9™9±eÓ%<±e°¤c¨!¦f±eÑ%<Ó=Ð>¸aÐ@Ð@ùÒ%<s   ‘/
c               ó    € V ^8„  d   QhRRRR/# )rO   r¶  zir.Layout | NonerP   rQ   rR   )rS   s   "rT   rU   r¯  =  s   € ÷ ñ Ð(8ð ¸Sñ rV   c                ó  <€ V f   R# S! V P                   4       pS! V P                  4       p^ RIHp V P                  pV! 4       '       d   Ve   VP
                  MT pR\        V P                  ,           V V V R2# )Nrw  )Ú_coor_enabledÚ")rf  ÚstrideÚ"torch.fx.experimental.proxy_tensorrÁ  r¼   rE  r   r»   )r¶  Úshape_annotationÚstride_annotationrÁ  r¼   Údevice_annotationr½  s   &     €rT   Ústringfy_layoutÚ,get_kernel_metadata.<locals>.stringfy_layout=  s�   ø€ Ø’>ÙÙ&5°f·k±kÓ&BÐ%CÐ Ù'6°v·}±}Ó'EÐ&FÐ!õ MàŸ™�ñ &Ÿš¨FÒ,>ð —K’Kà"˜8ð "ð œ F§L¡LÕ1Ð2Ð3CÐ2DØ(Ð)Ð*;Ð)<¸Að?ðrV   Úread_writesÚreadsz   %z
 : Tensor z = PlaceHolder[target=ÚwritesÚ%T)Úinclude_tensor_metadatac              3  óF   <"  € T F  pSP                    R V 2x € K  	  R# 5i)r«  N)Úcomment)rÉ   Úliner_  s   & €rT   rË   r¡  v  s(   øé € ð )á ;˜ð —‘Ð' s¨4¨&Õ1Û ;ùs   ƒ!z
   return Ú,Ú
),rz  rÉ  ÚcollectionsÚdefaultdictr©  r"   r^   rÀ  r;  rk  Únodesr¢  ÚsortrD  rŒ   r\   râ  rã  rQ   r�  r‚  rì   ró  ÚgetrÐ  r�  Úkeysr7  Úitemsrw  rv  rA   r¿  r½  rÊ  rË  r  Útry_get_bufferrÌ  Úformat_nodeÚversionr~   ÚextendÚ
splitlines)#rr  r_  r–  rƒ  Úinductor_nodesÚfrom_node_dictÚoriginal_aten_dictÚunique_graphsÚidxr   Únode_to_idx_mapr¶  r  r{  Úsort_strÚmetadataÚdetailed_metadataÚoriginal_noderÖ  Ú	all_readsÚ
all_writesr½  r·  rÈ  Úrr­  Ú
input_namer¶  ÚwÚoutput_namerñ   Úformatted_noderv  r¤  r½  s#   &f                              @@@rT   Úget_kernel_metadatarñ  ×  s>  û€ ô$ $ MÓ2€KÙ+6ÓW©; ¿)¹)ÀÑ:V—f�f©;€NÐWä ×,Ò,¬TÓ2€NÜ$×0Ò0´Ó6Ðð
 €LßÜ"Ñ"C±NÓ"CÓCˆÜˆ}Ó Ô"Ø)¨!Õ,×2Ñ2ˆLä˜<Ð)T×UÒUÜ8AÀ,×BTÑBTÔ8UÔ"VÑ8U©f¨c 1¢6Ñ8U�Ñ"VØIX�ÔFØ×ÑÜWð  ô ô ˆØ˜dŸi™iÔ'¨D¯I©I°oÕ,FÒ,RØ ŸI™I oÕ6ˆMØˆCÜ˜-¬¯©×)>Ñ)>×?Ò?Ü˜-×7Ñ7Ó8‘Ü˜M¬5¯:©:×+IÑ+I×JÒJÜ˜-×,Ñ,Ó.Ó/�ßØ"Õ'×.Ñ.¨t¯y©yÔ9Ø˜$Ÿ)™)Ô#Ø—)‘)˜KÕ(¨Õ+×0Ñ0ˆCØÕ×&Ñ& t§y¡y×1Ø�Y‰Y�]‰]Ð,Ó-°×>àŸ9™9Õ%×,Ñ,¨T¯Y©Y×7ñ ð  *6Ò)AÑ%Àz€Hà�?‰?Ð
˜1˜X˜JÐ&6°t·y±yÀ×ATÑATÓAVÓ7WÐ6Xð YØŸ9™9Ð%7×%<Ñ%<Ó%>Ó?Ð@Àð	Cð ð $ŸO™OÐ,Ð,NÐOÐPÐÜ & ~×';Ñ';Ó'=Ö >Ñˆ�uØ× Ñ Ø�‰Ð˜s = /°°d·i±iÄÀuÃÓ6NÐ5OÐPö	
ñ !?ð ÓÝà× Ñ  G§O¡OÐ#4Ð4DÐ!EÔFÜ%/£\ˆ	Ø "ˆ
Ü˜-¨¯©×9Ó9Ý&÷$ð $õ(A÷ð ô* #�Ü˜q -×0Ò0°A·M±MÒ4IÙÜ˜1Ÿ=™=¨'×2Ò2°q·}±}×7JÑ7JÒ7VØŸ]™]×0Ô0˜àŸ6™6 YÔ.Ù$Ø!Ÿ™ a§f¡fÔ-Ø!"§¡×!7Ñ!7¸¿¹Ó!?˜Ø!š>Ù$Ù-<¸VÀQÇVÁVÓ-LÑ*˜
 FØ)×0Ñ0Ø&Ÿ™Ð/¨t°J°<¸zÙ.¨vÓ6Ð7Ð7MÈjÈ\ÐYZð\öñ 1ô ˜AŸM™M¨8×4Õ4ØŸ™×,Ñ,Õ8àŸ]™]×1Ô1˜Ø!"§¡×!7Ñ!7¸¿¹Ó!?˜Ø!š>Ù$Ù)8¸ÀÇÁÓ)H™˜ Qà"×)Ñ)¨#°Õ*;Ö<ô 2ñ- #ó< #ˆDØ!×-Ñ-ÀdÐ-ÓKˆNØÒ)¬e¯m©m×.?×.?Ð.?ð "×(Ñ(ô )à .× 9Ñ 9Ô ;ó)ö ð
 "×(Ñ(¨G¯O©OÐ+<¸CÀÐ?OÐ)PÖQñ #ð 	× Ñ  G§O¡OÐ#4°J¸s¿x¹xÈ
Ó?SÐ>TÐ!UÔVà�T—Y‘YÐ0Ó1Ð1Ð1ùòk Xùó #Ws   ”Y$«Y$ÃY)c               ó$   € V ^8„  d   QhRRRRRR/# )rO   Úinitial_queuezIterable[torch.fx.Node]Úskip_filterzCallable[[Any], bool] | NonerP   úOrderedSet[torch.fx.Node]rR   )rS   s   "rT   rU   rU   ‚  s&   € ÷ ñ Ø*ðà-ðð ñrV   c                ó  € \        V 4      p \        V 4      pV '       dh   V P                  4       pVP                   FE  pV'       d   V! V4      '       d   K  WB9  g   K#  VP	                  V4       V P                  V4       KG  	  Ko  V# )zJReturns the set of nodes whose values depend on those within initial_queue)r©  r"   r`   rå  r  ró  )ró  rô  Údominated_setr¶  Úusers   &&   rT   Údominated_nodesrù  ‚  sq   € ô
 ˜Ó'€MÜ˜}Ó-€Mç
Ø× Ñ Ó"ˆØ—J”JˆDß™{¨4×0Ò0ÙØÖ(Ø×!Ñ! $Ô'Ø×$Ñ$ TÖ*ó ð ÐrV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r’   zSequence[IRNode]rë  zdict[str, IRNode]rP   rõ  rR   )rS   s   "rT   rU   rU   –  s)   € ÷ Gñ GØ
ðGØ$5ðGàñGrV   c                ó^  aa	€ ^RI Ho R VV	3R llo	\        V4      w  r#V Uu. uF  pS	! V4      '       g   K  VP                  NK!  	  pp\        V 4      w  rcV Uu. uF  pS	! V4      '       g   K  VP                  NK!  	  pp\	        \
        P                  ! . VOVO5!  4      # u upi u upi )rX   ru  c               ó    € V ^8„  d   QhRRRR/# )rO   r   rB   rP   rŠ   rR   )rS   s   "rT   rU   Ú$gather_origins.<locals>.__annotate__›  s   € ÷ 
ñ 
˜fð 
¨ñ 
rV   c                ór  <€ \        V SP                  4      '       d   S! V P                  4      # \        V SP                  4      '       d   S! V P                  4      # \        V SP                  4      ;'       d>    \        V SP
                  SP                  SP                  SP                  34      '       * # rš   )	rŒ   r±  r²  r³  rB   ÚComputedBufferÚInputsKernelÚInputBufferÚTemplateBuffer)r   rv  Úis_unrealized_nodes   &€€rT   r  Ú*gather_origins.<locals>.is_unrealized_node›  s”   ø€ Ü�a˜Ÿ™×&Ò&Ù% a§f¡fÓ-Ð-Ü�a˜Ÿ™×'Ò'Ù% a§f¡fÓ-Ð-Ü˜!˜RŸY™YÓ'÷ 
ð 
´
Øà×!Ñ!Ø—‘Ø—‘Ø×!Ñ!ð	ó1
ô -
ð 	
rV   )rw  rv  r#   ry  r"   Ú	itertoolsÚchain)
r’   rë  Úkwargs_flattenrñ   r³  Úkwargs_originsÚargs_flattenÚargs_originsrv  r  s
   &&      @@rT   Úgather_originsr  –  s—   ù€ õ ÷
ñ 
ô" % VÓ,Ñ€NÙ-;ÓW©^ cÑ?QÐRU×?V”k�c—k”k©^€NÐWÜ" 4Ó(�O€LÙ+7ÓS©< CÑ;MÈc×;R”K�C—K”K©<€LÐSÜ”i—o’oÐE |ÐE°nÓEÓFÐFùò XùâSs   ¤B%¹B%ÁB*Á1B*c               ó    € V ^8„  d   QhRRRR/# ©rO   r·  r‰   rP   rQ   rR   )rS   s   "rT   rU   rU   ³  s   € ÷ ,ñ ,�Jð , 3ñ ,rV   c                ó`   aaaa€ R R loR VV3R lloR VV3R lloR V3R lloS! V 4      # )	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               ó    € V ^8„  d   QhRRRR/# )rO   r·  r‰   rP   rŠ   rR   )rS   s   "rT   rU   Úsympy_str.<locals>.__annotate__º  s   € ÷ 
ñ 
˜*ð 
¨ñ 
rV   c                ó²   € \        V \        P                  4      ;'       d7    \        V P                  4      ^8H  ;'       d    V P                  ^ ,          R8H  # )rO   rÑ   )rŒ   r�   ÚMulr^   r’   )r·  s   &rT   Úis_neg_leadÚsympy_str.<locals>.is_neg_leadº  s@   € ä�tœUŸY™YÓ'×VÐV¬C°·	±	«N¸aÑ,?×VÐVÀDÇIÁIÈaÅLÐTVÑDVð	
rV   c               ó    € V ^8„  d   QhRRRR/# r  rR   )rS   s   "rT   rU   r  ¿  s   € ÷ 	'ñ 	'˜Jð 	'¨3ñ 	'rV   c                óœ  <€ \        V \        P                  4      '       d¥   \        V P                  4      ^8X  df   S! V P                  ^,          4      '       dG   S! V P                  ^ ,          4       RS! V P                  ^,          P                  ^,          4       2# RP                  \        SV P                  4      4      # S! V 4      # )rO   z - z + )rŒ   r�   rŽ   r^   r’   r�  r�   )r·  r  Úsympy_str_muls   &€€rT   Úsympy_str_addÚ sympy_str.<locals>.sympy_str_add¿  s—   ø€ Ü�dœEŸI™I×&Ò&ô �4—9‘9‹~ Ô"¡{°4·9±9¸Qµ<×'@Ò'@Ù'¨¯	©	°!­Ó5Ð6°c¹-ÈÏ	É	ÐRSÍ×HYÑHYÐZ[ÕH\Ó:]Ð9^Ð_Ð_à—z‘z¤# m°T·Y±YÓ"?Ó@Ð@á  Ó&Ð&rV   c               ó    € V ^8„  d   QhRRRR/# r  rR   )rS   s   "rT   rU   r  Ê  s   € ÷ 	(ñ 	(˜Jð 	(¨3ñ 	(rV   c                óò   <€ \        V \        P                  4      '       dP   S! V 4      '       d   R S! V P                  ^,          4       2# RP	                  \        SV P                  4      4      # S! V 4      # )Ú-z * )rŒ   r�   r  r’   r�  r�   )r·  r  Úsympy_str_atoms   &€€rT   r  Ú sympy_str.<locals>.sympy_str_mulÊ  sa   ø€ Ü�dœEŸI™I×&Ò&Ù˜4× Ò ð ™>¨$¯)©)°A­,Ó7Ð8Ð9Ð9à—z‘z¤# n°d·i±iÓ"@ÓAÐAá! $Ó'Ð'rV   c               ó    € V ^8„  d   QhRRRR/# r  rR   )rS   s   "rT   rU   r  Õ  s   € ÷ ñ ˜Zð ¨Cñ rV   c                ó¼  <€ \        V \        P                  4      '       d   V P                  # \        V \        P                  \        P
                  34      '       d   R S! V 4       R2# \        V \        \        \        \        34      '       dD   V P                  P                   R RP                  \        \        V P                  4      4       R2# \!        V 4      # )Ú(Ú)rv  )rŒ   r�   ÚSymbolrì   rŽ   r  rp   rl   rm   rn   Úfuncr    r�  r�   Ú	sympy_strr’   rQ   )r·  r  s   &€rT   r  Ú!sympy_str.<locals>.sympy_str_atomÕ  s¡   ø€ Ü�dœEŸL™L×)Ò)Ø—9‘9ÐÜ˜œuŸy™y¬%¯)©)Ð4×5Ò5Ø‘} TÓ*Ð+¨1Ð-Ð-Ü˜œ´¼(ÄHÐM×NÒNØ—i‘i×(Ñ(Ð)¨¨4¯9©9´S¼ÀDÇIÁIÓ5NÓ+OÐ*PÐPQÐRÐRä�t“9ÐrV   rR   )r·  r  r  r  r  s   &@@@@rT   r%  r%  ³  s3   û€ õ
÷
	'ñ 	'÷	(ñ 	(÷ð ñ ˜ÓÐrV   c               ó    € V ^8„  d   QhRRRR/# )rO   rB  r‰   rP   zValueRanges[Any]rR   )rS   s   "rT   rU   rU   â  s   € ÷ %ñ % ð %Ð0@ñ %rV   c                óÚ   € ^RI Hp \        P                  '       d<   \	        VP
                  RR4      ;p'       d   VP                  R8w  d   \        V 4      # \        P                  ! 4       # )rX   r¼  Úcurrent_nodeNÚ
index_expr)
r¿  r½  rv   Úcompute_all_boundsr[   Úinterpreterrá  rs   rt   Úunknown)rB  r½  Úfx_nodes   &  rT   Úget_bounds_index_exprr/  â  sQ   € Ýô 	×!×!Ð!Ü §¡¨~¸tÓDÐDˆWÖDØ�N‰N˜lÔ*ä˜5Ó!Ð!ä×"Ò"Ó$Ð$rV   c               ó    € V ^8„  d   QhRRRR/# ©rO   ÚprefixrQ   rP   rŠ   rR   )rS   s   "rT   rU   rU   ð  s   € ÷ ñ  ð ¨ñ rV   c                ó   € V ^ ,          R8H  # )r   rì  rR   ©r2  s   &rT   Úprefix_is_reductionr5  ð  s   € Ø�!�9˜ÑÐrV   c               ó    € V ^8„  d   QhRRRR/# r1  rR   )rS   s   "rT   rU   rU   ô  s   € ÷ (ñ ( ð (¨ñ (rV   c                ó   € V ^ ,          R9   # )r   )ra   ÚyÚzrR   r4  s   &rT   Úprefix_is_pointwiser:  ô  s   € Ø�!�9˜Ñ'Ð'rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r2  rr   rä  rƒ   rP   úsympy.SymbolrR   )rS   s   "rT   rU   rU   ø  s&   € ÷ 
Dñ 
D¨4ð 
D°cð 
D¸lñ 
DrV   c                ód   € V \         P                  8X  d   \        RV  24      h\        WRRR7      # )ú1
Used to generate an integer-nonnegative symbol.
z"prefix must not be SymT.SIZE, got T©ÚintegerÚnonnegative)rr   ÚSIZEr_   rq   )r2  rä  s   &&rT   Úsympy_index_symbol_with_prefixrC  ø  s4   € ð ”—‘ÔÜÐAÀ&ÀÐJÓKÐKô �v¨D¸dÔCÐCrV   c               ó    € V ^8„  d   QhRRRR/# )rO   ÚcheckrŠ   rP   rR   )rS   s   "rT   rU   rU     s   € ÷ Uñ U˜4ð U Dñ UrV   c                óf   € T ;'       g    \         P                  ;'       d    \         P                  # rš   )rv   Údebug_index_assertsÚassert_indirect_indexing)rE  s   &rT   Úgenerate_assertrI    s$   € Ø×/Ð/”V×/Ñ/×TÐT´V×5TÑ5TÐTrV   c               ó    € V ^8„  d   QhRRRR/# )rO   rì   rQ   rP   r<  rR   )rS   s   "rT   rU   rU   	  s   € ÷ 
>ñ 
>˜Sð 
> \ñ 
>rV   c                ón   € V ^ ,          R8X  d   \        RV : 24      h\        P                  ! V RRR7      # )r>  r÷   z)Symbol name must not start with 's', got Tr?  )r_   r�   r#  ©rì   s   &rT   Úsympy_index_symbolrM  	  s7   € ð ˆA…w�#„~ÜÐHÈÉÐQÓRÐRô �<Š<˜ d¸Ô=Ð=rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r·  r‰   Úreplacementszdict[sympy.Expr, Any]rP   rR   )rS   s   "rT   rU   rU     s"   € ÷ +ñ +�Zð +Ð/Dð +Èñ +rV   c                ó   € \        W4      # )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·  rO  s   &&rT   Ú
sympy_subsrQ    s   € ô
 �tÓ*Ð*rV   c               ó    € V ^8„  d   QhRRRR/# )rO   rZ  r   rP   z&TypeGuard[torch.SymInt | torch.Tensor]rR   )rS   s   "rT   rU   rU     s   € ÷ ñ �3ð ÐAñ rV   c                ó¢   € \        V \        P                  4      ;'       g/    \        V \        P                  4      ;'       d    V P                  # rš   )rŒ   r\   r5   rù  Ú_has_symbolic_sizes_strides)rZ  s   &rT   Úis_symbolicrU    s;   € Ü�aœŸ™Ó&÷ ð Ü�1”e—l‘lÓ#×EÐE¨×(EÑ(EðrV   c               ó    € V ^8„  d   QhRRRR/# ©rO   r’   r   rP   rŠ   rR   )rS   s   "rT   rU   rU   $  s   € ÷ -ñ -˜3ð - 4ñ -rV   c                 ój   € \         ;QJ d    R  V  4       F  '       g   K   R# 	  R# ! R  V  4       4      # )c              3  ó8   "  € T F  p\        V4      x € K  	  R # 5irš   )rU  rÎ  s   & rT   rË   Ú"any_is_symbolic.<locals>.<genexpr>%  s   é € Ð,¡t !Œ{˜1�~ˆ~£tùrÍ   TF©rÐ  )r’   s   *rT   Úany_is_symbolicr\  $  s+   € ß‹3Ñ,¡tÓ,�3Œ3Ð,Š3Ð,ˆ3Ñ,¡tÓ,Ó,Ð,rV   zaten.topk.defaultc               ó    € V ^8„  d   QhRRRR/# )rO   rþ  útorch.fx.GraphModulerP   ztorch.fx.Node | NonerR   )rS   s   "rT   rU   rU   G  s   € ÷ ñ ØðàñrV   c                óÜ   € ^ RI Hp V P                  P                   FK  p\	        V4      '       d   Vu # VP
                  P                  R4      ;pf   K9  V! V4      '       g   KI  Vu # 	  R# )r   )r+   r³  N)Ú%torch.fx.experimental.symbolic_shapesr+   rÀ  rÖ  Úis_cudagraph_unsafe_fx_noderD  rØ  )rþ  r+   r¶  r³  s   &   rT   Ú%get_first_incompatible_cudagraph_noderb  G  sW   € õ Là—‘—”ˆÜ& t×,Ò,ØŠKà—9‘9—=‘= Ó'Ð'ˆCÔ4Ñ9NÈs×9SÔ9SØŠKñ ñ rV   c               ó    € V ^8„  d   QhRRRR/# )rO   rþ  r^  rP   r8   rR   )rS   s   "rT   rU   rU   V  s   € ÷ ñ Ð(ð ¨Tñ rV   c                ó¼   € \        \        \        V P                  P                  4      4      4      pVP
                  R8w  d   \        RVP
                  : 24      hV# )z$Get the output node from an FX graphrý  z*Expected last node op to be 'output', got )ÚnextÚiterÚreversedrÀ  rÖ  rÉ  r_   )rþ  Ú	last_nodes   & rT   Úoutput_noderi  V  sN   € ä”Tœ( 2§8¡8§>¡>Ó2Ó3Ó4€IØ‡|�|�xÔÜØ8¸¿¹Ñ8HÐIó
ð 	
ð ÐrV   c               ó    € V ^8„  d   QhRRRR/# )rO   rþ  r^  rP   zOrderedSet[torch.device]rR   )rS   s   "rT   rU   rU   `  s   € ÷ 'ñ 'Ð,ð 'Ð1Iñ 'rV   c                ó   € V P                   P                  R R7      p\        R V 4       4      p\        V 4      P                  ^ ,          p\        V\        4      '       d   TMV3p\        R V 4       4      pW%,          # )rô  rÓ  c              3  óÒ   "  € T F]  p\        VP                  P                  R 4      \        P                  4      '       g   K>  VP                  R ,          P
                  x € K_  	  R# 5i©r³  N)rŒ   rD  rØ  r\   rù  r¼   )rÉ   r¶  s   & rT   rË   Ú"get_all_devices.<locals>.<genexpr>b  sC   é € ð 9á%ˆDÜ�d—i‘i—m‘m EÓ*¬E¯L©L×9ô 	 ˆ�	‰	�%Õ×ÖÛ%ùs   ‚;A'Á%A'c              3  ó*  "  € T F‰  p\        V\        P                  P                  4      '       g   K/  \        VP                  P                  R 4      \        P                  4      '       g   Kj  VP                  R ,          P                  x € K‹  	  R# 5irm  )rŒ   r\   rú  r8   rD  rØ  rù  r¼   )rÉ   rï  s   & rT   rË   rn  j  s[   é € ð 7áˆCÜ�cœ5Ÿ8™8Ÿ=™=×)ô 	ô �s—x‘x—|‘| EÓ*¬E¯L©L×9ô 	ˆ�‰��×ÖÛùs   ‚,B³6BÁ.%B)rÀ  Ú
find_nodesr"   ri  r’   rŒ   Útuple)rþ  Úplaceholder_nodesÚinput_devicesÚout_argÚout_argsÚout_devicess   &     rT   Úget_all_devicesrw  `  s~   € ØŸ™×+Ñ+¨}Ð+Ó=ÐÜ.8ñ 9á%ó9ó /€Mô ˜"‹o×"Ñ" 1Õ%€GÜ$ W¬e×4Ò4‰w¸7¸*€HÜ,6ñ 7áó7ó -€Kð Õ&Ð&rV   c               ó   € V ^8„  d   QhRR/# ©rO   rP   r  rR   )rS   s   "rT   rU   rU   v  s   € ÷ "ñ " ñ "rV   c                 ó®  € \        \        P                  P                  4       4       EFT  p V P	                  R 4      '       g   K  \        P                  V ,          pVP
                   Fÿ  pVP	                  R4      '       g   K  \        W4      p\        V\        P                  P                  P                  P                  4      '       g   Kg  VP                   F‡  p\        V\        P                  P                  P                  P                  4      '       g   KC  VP                  P                   p\#        VR4      '       g   Km  VP$                  P'                  4        K‰  	  EK  	  \        P                  V  EKW  	  R\        P                  9   d«   \        P                  R,          p\#        VP(                  P*                  R4      '       dn   \-        VP(                  P*                  P.                  4      p\#        VR4      '       d   V=M\#        VR4      '       d   RVn        VP(                  P*                  =\4        P6                  ! 4        R# )z&torch._inductor.runtime.compile_tasks.Útriton_Úmodztriton.runtime.driverÚutilsÚinstanceÚ	_instanceN)r©  ÚsysÚmodulesrÙ  Ú
startswithÚ__dict__r[   rŒ   r\   Ú	_inductorÚruntimeÚtriton_heuristicsÚCachingAutotunerÚcompile_resultsÚTritonCompileResultÚkernelÚrunr;  r|  Ú__del__ÚdriverÚactiverE  r}  r~  r  ÚgcÚcollect)Úmodule_nameÚmÚ	attr_namerŠ  r1  r‹  Ú
driver_modÚ	utils_clss           rT   Úunload_xpu_triton_pydsr–  v  sœ  € äœCŸK™K×,Ñ,Ó.×/ˆØ×%Ñ%Ð&N×OÒOÙÜ�K‰K˜Õ$ˆØŸœˆIØ×#Ñ# I×.Ô.Ü  Ó.�ÜØœEŸO™O×3Ñ3×EÑE×VÑV÷ô ð #)×"8Ô"8˜Ü%Ø"Ü!ŸO™O×3Ñ3×EÑE×YÑY÷ô ð
 #)§-¡-×"3Ñ"3˜CÜ& s¨E×2Ô2Ø #§¡§¡Ö 1ô #9ñ $ô �K‰K˜Ó$ñ' 0ð, ¤#§+¡+Ô-Ü—[‘[Ð!8Õ9ˆ
Ü�:×$Ñ$×+Ñ+¨W×5Ò5Ü˜Z×.Ñ.×5Ñ5×;Ñ;Ó<ˆIÜ�y *×-Ò-ØÑ&Ü˜ K×0Ò0Ø&*�	Ô#Ø×!Ñ!×(Ñ(Ð.ä‡J‚J†LrV   rê  Ú_registered_cachesc               ó    € V ^8„  d   QhRRRR/# )rO   r  r   rP   rR   )rS   s   "rT   rU   rU   ž  s   € ÷ 	ñ 	˜cð 	 cñ 	rV   c                ó¦   € \        V R4      '       d   \        V P                  4      '       g   \        V  R24      h\        P                  V 4       V # )z\
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   &rT   Úclear_on_fresh_cacherž  ž  sE   € ô
 �3˜×&Ò&¬h°s·±×.GÒ.GÜ ˜uÐ$GÐHÓIÐIä×Ñ˜cÔ"Ø€JrV   c               ó   € V ^8„  d   QhRR/# ry  rR   )rS   s   "rT   rU   rU   ª  s   € ÷ ñ �dñ rV   c                 ó@   € \          F  p V P                  4        K  	  R# )z
Clear all registered caches.
N)r—  rš  r�  s    rT   Úclear_cachesr¡  ª  s   € ÷ "ˆØ�‰Öó "rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r{  rQ   r—   rP   úIterator[None]rR   )rS   s   "rT   rU   rU   ³  s!   € ÷ "ñ "�#ð "˜cð " nñ "rV   c              #  óh  "  € \         P                  P                  V 4      p V\         P                  V &   Rx € Vf#   \         P                  P                  V R4       R# V\         P                  V &   R#   Tf"   \         P                  P                  T R4       i T\         P                  T &   i ; i5i)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ÚenvironrØ  r`   )r{  r—   Úolds   && rT   Ú_set_envr¨  ²  sy   é € ô �*‰*�.‰.˜Ó
€Cð"ØŒ�
‰
�3‰ÛàŠ;Ü�J‰J�N‰N˜3 Ö%à!ŒB�J‰J�s‹Oøð Š;Ü�J‰J�N‰N˜3 Õ%à!ŒB�J‰J�sŠOüs   ‚ B2£A5 º;B2Á5:B/Â/B2c               ó(   € V ^8„  d   QhRRRRRRRR/# )	rO   Úcache_entrieszdict[str, Any] | NoneÚdirú
str | NoneÚdeleterŠ   rP   r£  rR   )rS   s   "rT   rU   rU   È  s0   € ÷ 8ñ 8Ø(ð8à	ð8ð ð8ð ñ	8rV   c              #  ó¬  a"  € \        4        ^ RIHp V! \        P                  ! VR7      4      o \        RS4      ;_uu_ 4        \        P                  RS4       V! \        P                  P                  SR4      4      p\        RV4      ;_uu_ 4        Rx € \        V \        4      '       d¼   \        V 4      ^ 8w  d   \        R4      h\        P                  P                  V4      '       d|   \        P                   ! V4      pT P#                  V Uu/ uFI  pR	V9  g   K  V\        P                  P%                  \        P                  P                  WF4      4      bKK  	  up4       RRR4       RRR4       V'       dd   \'        4       '       d/   \(        P*                  P-                  4       '       d   \/        4        \0        P2                  ! S\'        4       V3R
 lR7       \        4        R# u upi   + '       g   i     L”; i  + '       g   i     LŸ; i  \4         d    \        P7                  RS4       h i ; i  \        4        i ; i5i)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.
)Únormalize_path_separator)r«  ÚTORCHINDUCTOR_CACHE_DIRzUsing inductor cache dir %sÚtritonÚTRITON_CACHE_DIRNz!expected empty cache_entries dictz.lockc                ó4   <€ \         P                  R SVR7      # )z*Failed to remove temporary cache dir at %s)Úexc_info)rä   Úwarning)r$  Úpathr´  Úinductor_cache_dirs   &&&€rT   r×  Úfresh_cache.<locals>.<lambda>ö  s   ø€ ´S·[±[Ø@Ø&Ø%ð 6Aô 6rV   )Úignore_errorsÚonerrorz(on error, temporary cache dir kept at %s)r¡  Útorch._inductor.cpp_builderr¯  ÚtempfileÚmkdtempr¨  rä   rå   r¥  r¶  r�  rŒ   Údictr^   r_   ÚexistsÚlistdirr¨  Úgetsizer�  r\   rK   r]   r–  ÚshutilÚrmtreeÚ	Exceptionrµ  )rª  r«  r­  r¯  Útriton_cache_dirÚfilesÚfr·  s   &&&    @rT   Úfresh_cacherÈ  Ç  s±  øé € ô „NåDá1´(×2BÒ2BÀsÔ2KÓLÐð(ÜÐ/Ð1C×DÕDÜ�I‰IÐ3Ð5GÔHÙ7Ü—‘—‘Ð/°Ó:ó Ðô Ð,Ð.>×?Õ?ÛÜ˜m¬T×2Ò2Ü˜=Ó)¨QÔ.Ü,Ð-PÓQÐQÜ—w‘w—~‘~Ð&6×7Ò7Ü "§
¢
Ð+;Ó <˜Ø%×,Ñ,ñ */óá). AØ#*°!Ñ#3ô !V ¤2§7¡7§?¡?´2·7±7·<±<Ð@PÓ3TÓ#UÒ UÙ).ñô÷ @÷ E÷& Ü�|Š|¤§	¡	× 6Ñ 6× 8Ò 8Ü&Ô(ä�MŠMØ"ô )›lôõô  	Žùò5÷ @×?ú÷ E×DûôF ô Ü�‰Ð>Ð@RÔSØðûô 	�üs…   ƒ.I²H ÁAHÂAG9	Ã2%G9	Ä	G4Ä%AG4Å&G9	Å-HÅ5H Æ3H Æ9/H Ç(IÇ4G9	Ç9H	ÈHÈH	ÈH È#IÉI ÉIÉIÚreversec               ó$   € V ^8„  d   QhRRRRRR/# )rO   Úseqr  rÉ  rŠ   rP   ú	list[int]rR   )rS   s   "rT   rU   rU   	  s!   € ÷ ñ �ð ¨Dð ¸Yñ rV   c               óª   € V P                   p\        \        V 4      4      p\        \	        W2R R7      4      pV'       g   \        \        V4      4      # V# )T©r{  rÉ  )Ú__getitem__r×   r^   r©  r7  rg  )rË  rÉ  ÚgetterÚa_rÚsort_idxs   &$   rT   ÚargsortrÓ  	  sC   € Ø�_‰_€FÜ
”�C“‹/€Cô ”F˜3°DÔ9Ó:€HßÜ”H˜XÓ&Ó'Ð'Ø€OrV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )	rO   rÂ  r.   rË  z)Sequence[int | torch.SymInt | sympy.Expr]rÉ  rŠ   rP   rÌ  rR   )rS   s   "rT   rU   rU     s0   € ÷ 7ñ 7Øð7à	2ð7ð ð	7ð
 ñ7rV   c          	     óL  a € R V 3R llp\        V4       UUu. uF?  w  rET\        V\        P                  4      '       d   VP                  P
                  MT3NKA  	  ppp\        V\        P                  ! V4      VR7      pV UUu. uF  w  rGVNK	  	  pppV# u uppi u uppi )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               ó$   € V ^8„  d   QhRRRRRR/# )rO   rZ  ztuple[int, sympy.Expr]r[  rP   rƒ   rR   )rS   s   "rT   rU   Ú!argsort_sym.<locals>.__annotate__&  s#   € ÷ !ñ !Ð%ð !Ð*@ð !ÀSñ !rV   c                óÎ   <aaa€ V w  poVw  poR  V3R llpR V3R lloV! SS8  VVV3R l4      '       d   R# V! SS8„  VVV3R l4      '       d   ^# W#8  d   ^# W#8”  d   R# ^ # )c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r·  z bool | torch.SymInt | sympy.ExprÚfallbackzCallable[[], bool]rP   rŠ   rR   )rS   s   "rT   rU   Ú.argsort_sym.<locals>.cmp.<locals>.__annotate__*  s&   € ÷ 		"ñ 		"Ø2ð		"à(ð		"ð ñ		"rV   c                óˆ   <€ \        V \        4      '       d   V #  SP                  V 4      #   \         d    T! 4       u # i ; irš   )rŒ   rŠ   Úevaluate_exprr,   )r·  rÚ  rÂ  s   &&€rT   ÚevaluateÚ*argsort_sym.<locals>.cmp.<locals>.evaluate*  sD   ø€ ô ˜$¤×%Ò%Ø�ð"Ø ×.Ñ.¨tÓ4Ð4øÜ.ô "Ù“zÒ!ð"ús   ›, ¬AÁ Ac               ó$   € V ^8„  d   QhRRRRRR/# )rO   Úlhsr‰   ÚrhsrP   rŠ   rR   )rS   s   "rT   rU   rÛ  5  s&   € ÷ 	Wñ 	W˜ð 	W¨*ð 	W¸ñ 	WrV   c                óJ   <€ SP                  V 4      SP                  V4      8  # rš   )Úoptimization_hint)rá  râ  rÂ  s   &&€rT   Úhint_ltÚ)argsort_sym.<locals>.cmp.<locals>.hint_lt5  s'   ø€ ð ×.Ñ.¨sÓ3°i×6QÑ6QÐRUÓ6VÑVÐVrV   c                 ó   <€ S! S S4      # rš   rR   ©Úa_valÚb_valrå  s   €€€rT   r×  Ú*argsort_sym.<locals>.cmp.<locals>.<lambda>;  ó   ø€ ©7°5¸%Ô+@rV   c                 ó   <€ S! SS 4      # rš   rR   rè  s   €€€rT   r×  rë  =  rì  rV   rÑ   rR   )	rZ  r[  Úa_idxÚb_idxrÞ  ré  rê  rå  rÂ  s	   &&   @@@€rT   r'  Úargsort_sym.<locals>.cmp&  sp   û€ Ø‰ˆˆuØ‰ˆˆu÷		"ð 		"÷	Wð 	Wñ �E˜E‘MÖ#@×AÒAØˆIÙ�E˜E‘MÖ#@×AÒAÙð
 Œ=ÙØŒ=ØˆIÙrV   rÎ  )	rk  rŒ   r\   r5   r¶  r·  r7  rM  Ú
cmp_to_key)	rÂ  rË  rÉ  r'  rä  r÷   Úexprsrñ   r1  s	   f&$      rT   Úargsort_symró    s•   ø€ ÷!ð !ôL   ”nôá$‰FˆCð 
œZ¨¬5¯<©<×8Ò8ˆa�f‰f�kŠk¸aÓ@Ù$ð 
ñ ô �5œi×2Ò2°3Ó7ÀÔI€EÙ %Ô&¡‘f�c‹c¡€FÑ&Ø€Mùóùó
 's   ™ABÂB c               ó    € V ^8„  d   QhRRRR/# )rO   r»   útorch.dtyperP   rƒ   rR   )rS   s   "rT   rU   rU   T  s   € ÷ 7ñ 7˜+ð 7¨#ñ 7rV   c                ó|   € V \         P                  8X  d   ^# \         P                  ! RV R7      P                  4       # )rŒ  ©r»   rR   )r\   r¤  rÓ   Úelement_sizer÷  s   &rT   Úget_dtype_sizerù  S  s-   € ð ”—‘ÔÙÜ�;Š;�r Ô'×4Ñ4Ó6Ð6rV   c                  óâ   a € ] tR tRtRtRR V 3R lllt]R R l4       tR tR R	 lt	R
 R lt
R tRR R lltRR R lltRR R lltRR R lltRR R lltRR R lltR V 3R lltRtV ;t# )ÚDualIndentedBufferi\  zïIndentedBuffer 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               ó    € V ^8„  d   QhRRRR/# )rO   Úinitial_indentrƒ   rP   r  rR   )rS   s   "rT   rU   ÚDualIndentedBuffer.__annotate__d  s   € ÷ 2ñ 2 sð 2°4ñ 2rV   c                	óF   <€ \         SV `  V4       \        V4      V n        R # rš   )ÚsuperÚ__init__r   Úaot)rH  rý  Ú	__class__s   &&€rT   r  ÚDualIndentedBuffer.__init__d  s   ø€ Ü‰Ñ˜Ô(Ü! .Ó1ˆŽrV   c               ó   € V ^8„  d   QhRR/# ©rO   rP   r   rR   )rS   s   "rT   rU   rþ  i  s   € ÷ ñ �^ñ rV   c                ó   € V # )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.
rR   ©rH  s   &rT   ÚjitÚDualIndentedBuffer.jith  s	   € ð ˆrV   c                	óh   € \         P                  ! W4       V P                  P                  V4       R # rš   )r   Ú	writeliner  ©rH  rÑ  s   &&rT   r  ÚDualIndentedBuffer.writelineq  s"   € Ü× Ò  Ô,Ø�‰×Ñ˜4Ö rV   c               ó   € V ^8„  d   QhRR/# ry  rR   )rS   s   "rT   rU   rþ  u  s   € ÷ -ñ - Tñ -rV   c                	ó2   € \         P                  ! W4       R # rš   )r   r  r  s   &&rT   Úwriteline_jitÚ DualIndentedBuffer.writeline_jitu  s   € Ü× Ò  Ö,rV   c               ó   € V ^8„  d   QhRR/# ry  rR   )rS   s   "rT   rU   rþ  x  s   € ÷ !ñ ! Tñ !rV   c                	ó<   € V P                   P                  V4       R # rš   )r  r  r  s   &&rT   Úwriteline_aotÚ DualIndentedBuffer.writeline_aotx  s   € Ø�‰×Ñ˜4Ö rV   c                	ó:   € V F  pV P                  V4       K  	  R # rš   ©r  )rH  ÚlinesrÑ  s   && rT   Ú
writelinesÚDualIndentedBuffer.writelines{  s   € ÛˆDØ�N‰N˜4Ö ó rV   c               ó    € V ^8„  d   QhRRRR/# ©rO   ÚstriprŠ   rP   r  rR   )rS   s   "rT   rU   rþ    s   € ÷ 0ñ 0¨ð 0¸ñ 0rV   c                	ó¶   € \         P                  ! WVR 7       \        V\        4      '       d   VP                  MTpV P                  P                  W2R 7       R# ©)r  N)r   ÚsplicerŒ   rû  r  ©rH  Ú
other_coder  Ú	aot_others   &&& rT   r!  ÚDualIndentedBuffer.splice  sC   € ô 	×Ò˜d°eÕ<ä(¨Ô5G×HÒHˆJ�NŠNÈjð 	ð 	�‰�‰˜	ˆÖ/rV   c               ó    € V ^8„  d   QhRRRR/# r  rR   )rS   s   "rT   rU   rþ  ˆ  s   € ÷ =ñ =¨Dð =¸Tñ =rV   c                	ó6   € \         P                  ! WVR 7       R# r   )r   r!  ©rH  r#  r  s   &&&rT   Ú
splice_jitÚDualIndentedBuffer.splice_jitˆ  s   € Ü×Ò˜d°e×<rV   c               ó    € V ^8„  d   QhRRRR/# r  rR   )rS   s   "rT   rU   rþ  ‹  s   € ÷ 0ñ 0¨Dð 0¸Tñ 0rV   c                	ó†   € \        V\        4      '       d   VP                  MTpV P                  P                  W2R 7       R# r   )rŒ   rû  r  r!  r"  s   &&& rT   Ú
splice_aotÚDualIndentedBuffer.splice_aot‹  s1   € ä(¨Ô5G×HÒHˆJ�NŠNÈjð 	ð 	�‰�‰˜	ˆÖ/rV   c               ó   € V ^8„  d   QhRR/# )rO   Úoffsetrƒ   rR   )rS   s   "rT   rU   rþ  ‘  s   € ÷ ñ ˜Sñ rV   c                	óH   a a€ \         P                  VV 3R  l4       pV! 4       # )c               3  óŒ  <"  € S;P                   S ,          un         SP                  ;P                   S ,          un          R x € S;P                   S ,          un         SP                  ;P                   S ,          un         R #   S;P                   S ,          un         SP                  ;P                   S ,          un         i ; i5irš   ©Ú_indentr  )r0  rH  s   €€rT   r^  Ú&DualIndentedBuffer.indent.<locals>.ctx’  sz   øé € à�LŠL˜FÕ"�LØ�H‰H×Ò Õ&Õð+Ûà—’ Õ&•Ø—‘× Ò  FÕ*× øð —’ Õ&•Ø—‘× Ò  FÕ*Ö üs   ƒ=CÁB Á>CÂ>CÃC)Ú
contextlibÚcontextmanager)rH  r0  r^  s   ff rT   ÚindentÚDualIndentedBuffer.indent‘  s$   ù€ Ü	×	"Ñ	"õ	+ó 
#ð	+ñ ‹uˆrV   c               ó    € V ^8„  d   QhRRRR/# ©rO   r0  rƒ   rP   r  rR   )rS   s   "rT   rU   rþ  ž  s   € ÷ #ñ # ð #¨Dñ #rV   c                	ó~   € V ;P                   V,          un         V P                  ;P                   V,          un         R # rš   r3  ©rH  r0  s   &&rT   Ú	do_indentÚDualIndentedBuffer.do_indentž  ó%   € Ø�Š˜Õ�Ø�‰×Ò˜FÕ"×rV   c               ó    € V ^8„  d   QhRRRR/# r;  rR   )rS   s   "rT   rU   rþ  ¢  s   € ÷ #ñ # #ð #¨dñ #rV   c                	ó~   € V ;P                   V,          un         V P                  ;P                   V,          un         R # rš   r3  r=  s   &&rT   Údo_unindentÚDualIndentedBuffer.do_unindent¢  r@  rV   c               ó   € V ^8„  d   QhRR/# ry  rR   )rS   s   "rT   rU   rþ  ¦  s   € ÷ ñ �tñ rV   c                	óX   <€ \         SV `  4        V P                  P                  4        R # rš   )r   Úclearr  ©rH  r  s   &€rT   rG  ÚDualIndentedBuffer.clear¦  s   ø€ Ü‰‰ŒØ�‰�‰ÖrV   )r  ©r   rÖ  rŸ   )r    r¡   r¢   r£   r¤   r  Úpropertyr	  r  r  r  r  r!  r)  r-  r8  r>  rC  rG  r¨   Ú__classcell__©r  s   @rT   rû  rû  \  sb   ø† ñ÷2ò 2ð ôó ðò!õ-õ!ò!÷0÷=÷0÷÷#÷#÷ö rV   rû  c                  óR   € ] tR tRtRtR R ltR R ltRR R lltRR	 R
 lltRt	R# )ÚAotOnlyBufferi«  a  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               ó   € V ^8„  d   QhRR/# ry  rR   )rS   s   "rT   rU   ÚAotOnlyBuffer.__annotate__´  s   € ÷ ñ  Tñ rV   c                	ó   € R # rš   rR   r  s   &&rT   r  ÚAotOnlyBuffer.writeline_jit´  ó   € ÙrV   c               ó   € V ^8„  d   QhRR/# ry  rR   )rS   s   "rT   rU   rQ  ·  s   € ÷ ñ  Tñ rV   c                	ó(   € V P                  V4       R # rš   r  r  s   &&rT   r  ÚAotOnlyBuffer.writeline_aot·  s   € Ø�‰�tÖrV   c               ó    € V ^8„  d   QhRRRR/# r  rR   )rS   s   "rT   rU   rQ  º  s   € ÷ ñ ¨Dð ¸Tñ rV   c                	ó   € R # rš   rR   r(  s   &&&rT   r)  ÚAotOnlyBuffer.splice_jitº  rT  rV   c               ó    € V ^8„  d   QhRRRR/# r  rR   )rS   s   "rT   rU   rQ  ½  s   € ÷ -ñ -¨Dð -¸Tñ -rV   c                	ó*   € V P                  WR 7       R# r   )r!  r(  s   &&&rT   r-  ÚAotOnlyBuffer.splice_aot½  s   € Ø�‰�JˆÖ,rV   rR   NrÖ  )
r    r¡   r¢   r£   r¤   r  r  r)  r-  r¨   rR   rV   rT   rO  rO  «  s    † ñõõ÷÷-ó -rV   rO  c               ó   € V ^8„  d   QhRR/# r  rR   )rS   s   "rT   rU   rU   Á  s   € ÷ ñ ˜^ñ rV   c                 óº   € ^RI Hp  V P                  P                  '       d   \	        4       # V P                  P
                  '       d   \        4       # \        4       # )a   Construct the IndentedBuffer subclass matching the current codegen mode.

Dual-wrapper mode -> DualIndentedBuffer (JIT and AOTI both active).
Pure AOTI -> AotOnlyBuffer (writeline_aot writes; writeline_jit drops).
Pure JIT  -> IndentedBuffer  (writeline_jit writes; writeline_aot drops).
r¼  )r¿  r½  rÀ  Úis_dual_wrapper_moderû  Úaot_moderO  r   r¼  s    rT   Úmake_codegen_bufferrb  Á  sA   € õ à‡w�w×#×#Ð#Ü!Ó#Ð#Ø‡w�w××ÐÜ‹ÐÜÓÐrV   c                  ó:   a € ] tR tRtR V 3R lltR R ltRtV ;t# )ÚFakeIndentedBufferiÑ  c               ó   € V ^8„  d   QhRR/# ry  rR   )rS   s   "rT   rU   ÚFakeIndentedBuffer.__annotate__Ò  s   € ÷ ñ ˜$ñ rV   c                	ó$   <€ \         SV `  4        R # rš   )r   r  rH  s   &€rT   r  ÚFakeIndentedBuffer.__init__Ò  s   ø€ Ü‰ÑÖrV   c               ó    € V ^8„  d   QhRRRR/# )rO   rì   rQ   rP   r   rR   )rS   s   "rT   rU   rf  Õ  s   € ÷ 
ñ 
 Sð 
¨Sñ 
rV   c                	óX   € VR 8X  d   \         P                  W4      # \        RV R24      h)r  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  )rH  rì   s   &&rT   rl  Ú#FakeIndentedBuffer.__getattribute__Õ  s9   € Ø�;ÔÜ×*Ñ*¨4Ó6Ð6ÜØ! $ ð (=ð =ó
ð 	
rV   rR   )r    r¡   r¢   r£   r  rl  r¨   rL  rM  s   @rT   rd  rd  Ñ  s   ø† ÷ð ÷
ó 
rV   rd  c               ó   € V ^8„  d   QhRR/# )rO   rP   r£  rR   )rS   s   "rT   rU   rU   á  s   € ÷ @ñ @˜~ñ @rV   c               #  ó¸   "  € \         P                  \         P                  r R x € Wu\         n        \         n        R #   Yu\         n        \         n        i ; i5irš   )r€  ÚstdoutÚstderr)Úinitial_stdoutÚinitial_stderrs     rT   Úrestore_stdout_stderrrt  à  s9   é € ä%(§Z¡Z´·±�Nð@Ûà!/ÐŒŒ
”C–Jø ÐŒŒ
”C•Jüs   ‚ A£? §A¿AÁAc                  óJ   a € ] tR tRtRtR V 3R lltR R ltR R ltR	tV ;t	# )
ÚDelayReplaceLineié  z6At end of codegen call `line.replace(key, value_fn())`c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r{  rQ   Úvalue_fnzCallable[[], str]rÑ  rR   )rS   s   "rT   rU   ÚDelayReplaceLine.__annotate__ì  s"   € ÷ !ñ !˜Cð !Ð+<ð !ÀCñ !rV   c                	ó>   <€ \         SV `  V4       Wn        W n        R # rš   )r   r  r{  rx  )rH  r{  rx  rÑ  r  s   &&&&€rT   r  ÚDelayReplaceLine.__init__ì  s   ø€ Ü‰Ñ˜ÔØŒØ ŽrV   c               ó   € V ^8„  d   QhRR/# rN   rR   )rS   s   "rT   rU   ry  ñ  s   € ÷ <ñ <˜#ñ <rV   c                	ój   € V P                   P                  V P                  V P                  4       4      # rš   )rÑ  Úreplacer{  rx  r  s   &rT   rI  ÚDelayReplaceLine.__call__ñ  s#   € Ø�y‰y× Ñ  §¡¨4¯=©=«?Ó;Ð;rV   c               ó    € V ^8„  d   QhRRRR/# )rO   rÑ  rQ   rP   rv  rR   )rS   s   "rT   rU   ry  ô  s   € ÷ ?ñ ?˜cð ?Ð&6ñ ?rV   c                	óD   € \        V P                  V P                  V4      # rš   )rv  r{  rx  r  s   &&rT   Ú	_new_lineÚDelayReplaceLine._new_lineô  s   € Ü §¡¨$¯-©-¸Ó>Ð>rV   )r{  rx  )
r    r¡   r¢   r£   r¤   r  rI  r‚  r¨   rL  rM  s   @rT   rv  rv  é  s   ø† Ù@÷!ð !õ
<÷?ó ?rV   rv  c               ó    € V ^8„  d   QhRRRR/# )rO   Úindex_or_devicezint | torch.devicerP   rŠ   rR   )rS   s   "rT   rU   rU   ù  s   € ÷ ñ Ð 2ð ¸4ñ rV   c                ó&  € \        V \        P                  4      '       d   T pM\        P                  ! \        4       V 4      p\        P
                  ! V4      p\        P                  P                  '       dU   VP                  f   \        R4      hVP                  ^	8  g   VP                  ^
8X  d   \        P                  R4       R# R# VP                  R8X  d   ^M^DpVP                  pWC8  d   \        P                  RRVRV/R	7       R# R# )
Nz+DeviceProperties.major is None for ROCm GPUz6GPU arch does not support max_autotune_gemm mode usageFTrK   z,Not enough SMs to use max_autotune_gemm modeÚmin_smsÚ	avail_sms)Úextra)rŒ   r\   r¼   rd   r   ÚcreaterÝ  r~   Úmajorr_   rä   rµ  rE  Úmulti_processor_count)r…  r¼   Úpropr‡  rˆ  s   &    rT   Ú
is_big_gpurŽ  ø  sÖ   € ä�/¤5§<¡<×0Ò0Ø ‰ä—’œl›n¨oÓ>ˆä×"Ò" 6Ó*€Dô ‡}�}××ÐØ�:‰:ÒÜ Ð!NÓOÐOØ�:‰:˜Œ>˜TŸZ™Z¨2Ô-Ü�K‰KÐPÔQÙÙà—K‘K 5Ô(‰b¨b€GØ×*Ñ*€IØÔÜ�‰Ø:Ø˜g {°IÐ>ð 	ô 	
ñ ÙrV   c               ó   € V ^8„  d   QhRR/# ©rO   rP   rƒ   rR   )rS   s   "rT   rU   rU     s   € ÷ Jñ J˜ñ JrV   c                 óì   € \         P                  P                  4       '       d)   \         P                  P                  4       P                  # \         P
                  P                  R 4      P                  # ©rJ   )r\   rK   r]   Úget_device_propertiesÚgpu_subslice_countrJ   rŒ  rR   rV   rT   Úget_max_num_smsr•    sI   € ä‡y�y×Ñ×ÒÜ�y‰y×.Ñ.Ó0×CÑCÐCÜ�:‰:×+Ñ+¨FÓ3×IÑIÐIrV   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU     s   € ÷ )ñ )�Dñ )rV   c                 óâ   € \         P                  P                  4       '       g   R# \         P                  P                  \         P                  P	                  4       4      p V P
                  ^
8H  # )zEReturns true if the device is a NVIDIA B200, otherwise returns false.F)r\   rJ   r]   r“  rG  r‹  )Údevice_propertiess    rT   Ú
using_b200r™    sM   € ô �:‰:×"Ñ"×$Ò$ÙäŸ
™
×8Ñ8¼¿¹×9RÑ9RÓ9TÓUÐØ×"Ñ" bÑ(Ð(rV   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU   '  s   € ÷ ñ  $ñ rV   c                 óØ   € \         P                  P                  4       ;'       dF    \         P                  P                  RJ ;'       d"    \         P                  P                  4       R8¬  # )zCReturn true for NVIDIA CUDA devices with compute capability SM100+.N)é
   r   )r\   rJ   r]   rÝ  r~   Úget_device_capabilityrR   rV   rT   Úis_nvidia_sm100_or_laterrž  '  sT   € ô 	�
‰
×ÑÓ!÷ 	:ð 	:Ü�M‰M×Ñ Ð%÷	:ð 	:ä�J‰J×,Ñ,Ó.°'Ñ9ðrV   c               ó   € V ^8„  d   QhRR/# r�  rR   )rS   s   "rT   rU   rU   0  s   € ÷ Iñ I�Sñ IrV   c                 óÔ   € \         P                  P                  4       '       d   \        4       # \         P                  P                  4       p \        4       V e	   V ,
          # ^ ,
          # )zFHandle experimental carveout if set otherwise return hardware SM count)r\   rK   r]   r•  r9  Ú_get_sm_carveout_experimental)Úcarveouts    rT   Úget_num_smsr£  0  sM   € ô ‡y�y×Ñ×ÒÜÓ Ð Ü�x‰x×5Ñ5Ó7€HÜÓ¨HÒ,@ ÕHÐHÀaÕHÐHrV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )	rO   Únum_tma_descriptorsrƒ   r¼   r@  Únum_programsú
int | NonerP   r:   rR   )rS   s   "rT   rU   rU   9  s0   € ÷ ñ Øðàðð ðð ñ	rV   c                ó¬   € ^RI HpHp Vf   \        4       pVP	                  R4      pW ,          \
        ,          pV! VVVVP                  ! 4       R7      # )zKBuilds and returns a WorkspaceArg for the device side TMA workspace buffer.)r:   ÚWorkspaceZeroModeF)ÚcountÚ	zero_moder¼   Ú
outer_name)Úcodegen.commonr:   r©  r£  Ú	from_boolÚTMA_DESCRIPTOR_SIZEÚunique_name)r¥  r¼   r¦  r:   r©  r«  rf  s   &&&    rT   Úget_tma_workspace_argr±  9  sU   € ÷ @àÒÜ"“}ˆØ!×+Ñ+¨EÓ2€IØÕ-Ô0CÕC€DÙØØØØ×+Ò+Ó-ô	ð rV   c               ó    € V ^8„  d   QhRRRR/# )rO   Úblock_krƒ   rP   rR   )rS   s   "rT   rU   rU   M  s   € ÷ ñ ˜sð ¨Cñ rV   c                ó¸   € \         P                  P                  '       g   ^ # R\         P                  P	                  ^ 4      P
                  9   d
   V ^8:  d   ^# ^# )r   Úgfx942)r\   rÝ  r~   rJ   r“  ÚgcnArchName)r³  s   &rT   Úget_default_kpackr·  M  s?   € Ü�=‰=××ÐÙØ”5—:‘:×3Ñ3°AÓ6×BÑBÔBÀwÐRTÄ}ÙÙrV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r¶  rC   Úallowed_layout_dtypeszlist[torch.dtype]rP   rŠ   rR   )rS   s   "rT   rU   rU   U  s$   € ÷ ñ ØðØ+<ðà	ñrV   c                ó  € V P                   V9  d"   \        P                  R V P                   V4       \        V P                  P
                  4      ;'       d-    V P                   V9   ;'       d    \        V P                  4      # )zDNot using template since dtype %s is not in allowed layout dtypes %s)r»   rä   rå   Úis_gpur¼   rE  rŽ  )r¶  r¹  s   &&rT   Ú_use_template_for_gpur¼  U  sp   € ð ‡|�|Ð0Ô0Ü�	‰	ØRØ�L‰LØ!ô	
ô 	ˆv�}‰}×!Ñ!Ó"÷ 	&ð 	&Ø�L‰LÐ1Ñ1÷	&ð 	&ä�v—}‘}Ó%ðrV   c               ó    € V ^8„  d   QhRRRR/# ©rO   ÚbackendrQ   rP   rŠ   rR   )rS   s   "rT   rU   rU   e  s   € ÷ ñ  3ð ¨4ñ rV   c                óÂ   € V P                  4       \        P                  P                  4       P                  R 4       Uu. uF  qP	                  4       NK  	  up9   # u upi ©rÒ  )r(  rv   Úmax_autotune_gemm_backendsr§  r  ©r¿  ra   s   & rT   Ú_use_autotune_backendrÄ  e  óP   € Ø�=‰=‹?Ü!×<Ñ<×BÑBÓD×JÑJÈ3ÔOóÙO�a�‰Ž	ÑOññ ð ùò ó   ¿Ac               ó    € V ^8„  d   QhRRRR/# r¾  rR   )rS   s   "rT   rU   rU   k  s   € ÷ ñ ¨ð °ñ rV   c                óÂ   € V P                  4       \        P                  P                  4       P                  R 4       Uu. uF  qP	                  4       NK  	  up9   # u upi rÁ  )r(  rv   Úmax_autotune_conv_backendsr§  r  rÃ  s   & rT   Ú_use_conv_autotune_backendrÊ  k  rÅ  rÆ  c               ó    € V ^8„  d   QhRRRR/# r¾  rR   )rS   s   "rT   rU   rU   q  s   € ÷ ñ °3ð ¸4ñ rV   c                óÄ   € V P                  4       \        P                  P                  4       P                  R 4       Uu. uF  pVP	                  4       NK  	  up9   # u upi rÁ  )r(  rv   Ú%max_autotune_conv_bwd_weight_backendsr§  r  rÃ  s   & rT   Ú%_use_conv_bwd_weight_autotune_backendrÎ  q  sU   € Ø�=‰=‹?ä×=Ñ=×CÑCÓE×KÑKÈCÔPóáPˆAð 	
�‰Ž	ÙPññ ð ùò ó   ¿Ac               ó    € V ^8„  d   QhRRRR/# r¾  rR   )rS   s   "rT   rU   rU   x  s   € ÷ ñ °#ð ¸$ñ rV   c                óÄ   € V P                  4       \        P                  P                  4       P                  R 4       Uu. uF  pVP	                  4       NK  	  up9   # u upi rÁ  )r(  rv   Ú$max_autotune_conv_bwd_input_backendsr§  r  rÃ  s   & rT   Ú$_use_conv_bwd_input_autotune_backendrÓ  x  sU   € Ø�=‰=‹?ä×<Ñ<×BÑBÓD×JÑJÈ3ÔOóáOˆAð 	
�‰Ž	ÙOññ ð ùò rÏ  Úenable_int32Úenable_float8Úcheck_max_autotunec          
     ó,   € V ^8„  d   QhRRRRRRRRRR/# )rO   r¶  rC   rÔ  rŠ   rÕ  rÖ  rP   rR   )rS   s   "rT   rU   rU     s<   € ÷ ñ Øðð ðð ð	ð
 ðð 
ñrV   c               ó,  € ^RI HpHp \        P                  \        P
                  \        P                  .pV'       d?   \        P                  \        P
                  \        P                  \        P                  .pV'       d0   VP                  \        P                  \        P                  .4       \        V P                  P                  4      ;'       d    \        W4      ;'       g9    V P                  P                  R8H  ;'       d‡    V P                  V9   ;'       dp    \         P"                  ;'       g(    \         P$                  ;'       g    V'       * ;'       d1    \'        R4      ;'       d    V! V P                  VP(                  4      # )rX   )ÚBackendFeatureÚhas_backend_featurerC  ÚTRITON)r­  rÙ  rÚ  r\   rÔ   r‘  r“  r›  rÞ  r‡  rˆ  r»  r¼   rE  r¼  r»   rv   Úmax_autotuneÚmax_autotune_gemmrÄ  ÚTRITON_TEMPLATES)r¶  rÔ  rÕ  rÖ  rÙ  rÚ  Úlayout_dtypess   &$$$   rT   Úuse_triton_templaterà    s0  € ÷ Dä—]‘]¤E§N¡N´E·M±MÐB€MßÜŸ™¬¯©¼¿¹ÄuÇ{Á{ÐSˆßØ×Ñœe×1Ñ1´5×3DÑ3DÐEÔFô �v—}‘}×)Ñ)Ó*÷ Að AÜ)¨&Ó@÷Oð Oð —‘×"Ñ" eÑ+÷
	Pð
 N°·±ÀÑ0M÷
	Pð 
	Pô × Ñ ×VÐV¤F×$<Ñ$<×VÐVÐDVÔ@V÷
	Pð 
	Pô " (Ó+÷
	Pð 
	Pñ   §¡¨~×/NÑ/NÓOðrV   Úoutput_layoutÚ
add_guardsc               ó(   € V ^8„  d   QhRRRRRRRR/# )rO   ÚmatricesrB   rá  úLayout | Nonerâ  rŠ   rP   rR   )rS   s   "rT   rU   rU   œ  s2   € ÷ Bñ BØðBØ&3ðBØHLðBà	ñBrV   c                ó<  aaaaaa	€ ^ RI Hp ^RIHo R V3R lloR VV3R llpR VVV	3R lloR	 VVV3R
 lloR V3R llo	V! 4       ;'       dK    \        ;QJ d    V3R lV 4       F  '       d   K   RM	  RM! V3R lV 4       4      ;'       d	    V! V 4      # )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
)Úhas_triton_tma_devicer¼  c               ó    € V ^8„  d   QhRRRR/# )rO   Ú
expr_bytesrs  rP   rŠ   rR   )rS   s   "rT   rU   Ú!can_use_tma.<locals>.__annotate__±  s   € ÷ Xñ XÐ-ð X°$ñ XrV   c                óX   <€ SP                   P                  P                  V \        4      # rš   )rÀ  rÁ  Ústatically_known_multiple_ofÚTMA_ALIGNMENT)ré  r½  s   &€rT   Ú_alignedÚcan_use_tma.<locals>._aligned±  s    ø€ Ø�w‰w×Ñ×<Ñ<¸ZÌÓWÐWrV   c               ó    € V ^8„  d   QhRRRR/# )rO   r¶  rå  rP   rŠ   rR   )rS   s   "rT   rU   rê  ´  s   € ÷ 9ñ 9¨-ð 9¸Dñ 9rV   c                óž   <€ V f   R# V P                   pV P                  pV P                  pS! V P                  4      '       g   R# S! WV4      # )NTF)rf  rÃ  r»   r0  )r¶  rb  Ústridesr»   rî  Ú_is_tma_compatibles   &   €€rT   Ú_is_tma_compatible_layoutÚ.can_use_tma.<locals>._is_tma_compatible_layout´  sG   ø€ ØŠ>ÙØ—‘ˆØ—-‘-ˆØ—‘ˆñ ˜Ÿ™×&Ò&Ùá! %°%Ó8Ð8rV   c               ó    € V ^8„  d   QhRRRR/# )rO   r’  rB   rP   rŠ   rR   )rS   s   "rT   rU   rê  Á  s   € ÷ 9ñ 9 Vð 9°ñ 9rV   c                ó(  <€ V P                  4       pV P                  4       pV P                  4       pV P                  4       SP                  P
                  9   d   R # V P                  4       ;pe   VP                  R8X  d
   S! WV4      # S! WV4      # )FrK   )Úget_sizeÚ
get_strideÚ	get_dtyper3  rÀ  Úunaligned_buffersÚ
get_devicerE  )r’  rb  rò  r»   Úm_devicer½  ró  Ú_is_tma_compatible_xpus   &    €€€rT   Ú_is_tma_compatible_matrixÚ.can_use_tma.<locals>._is_tma_compatible_matrixÁ  sw   ø€ Ø—
‘
“ˆØ—,‘,“.ˆØ—‘“ˆð �:‰:‹<˜1Ÿ7™7×4Ñ4Ô4ÙàŸ™›Ð&ˆHÒ3¸¿¹ÈÔ8NÙ)¨%¸%Ó@Ð@á! %°%Ó8Ð8rV   c               ó(   € V ^8„  d   QhRRRRRRRR/# ©	rO   rb  rV  rò  zSequence[_IntLike]r»   rõ  rP   rŠ   rR   )rS   s   "rT   rU   rê  Ï  s0   € ÷ 3ñ 3Ø#ð3à#ð3ð ð3ð 
ñ	3rV   c                óà  <€ \        V 4      pVP                  pV^8  g   V^8”  d   R# V\        9  d   R# S'       dL   SP                  P                  P                  V 4      pSP                  P                  P                  V4      pMhV  Uu. uF(  pSP                  P                  P                  V4      NK*  	  ppV Uu. uF(  pSP                  P                  P                  V4      NK*  	  pp\        V4       U	Uu. uF5  w  r˜SP                  P                  P                  V^4      '       g   K3  V	NK7  	  p
p	p\        V
4      ^8w  d   R# V
^ ,          p\        V4       F$  w  r˜W›8X  d   K  S! W„,          4      '       d   K#   R# 	  W[,          pS! WÄ,          4      '       g   R# V^8X  d/   SP                  P                  P                  V^ 4      '       g   R# R# u upi u upi u upp	i )rX   FT)
r^   Úitemsizer€   rÀ  rÁ  Úguard_int_seqÚ!replace_backed_symbols_with_hintsrk  Ústatically_known_equalsÚstatically_known_geq)rb  rò  r»   Úrankr  Úsizes_iÚ	strides_ir÷   Úströ   rj  Ú	inner_idxÚ	inner_dimr½  rî  râ  s   &&&          €€€rT   ró  Ú'can_use_tma.<locals>._is_tma_compatibleÏ  sµ  ø€ ô
 �5‹zˆØ—>‘>ˆà�!Œ8�t˜a”xÙàÔ-Ô-ÙçØ—g‘g×&Ñ&×4Ñ4°UÓ;ˆGØŸ™×(Ñ(×6Ñ6°wÓ?‰Iñ PUóÙOTÈ!�—‘× Ñ ×BÑBÀ1ÖEÉuð ð ñ RYóÙQXÈ2�—‘× Ñ ×BÑBÀ2ÖFÑQXð ð ô # 9Ô-ô
á-‘�Ø�w‰w×Ñ×7Ñ7¸¸A×>÷ ˆAÙ-ð 	ñ 
ô
 ˆu‹:˜Œ?ÙØ˜!•Hˆ	ô ˜yÖ)‰EˆAØŒ~ÙÙ˜B�M×*Ô*Úñ	 *ð Õ&ˆ	Ù˜	Õ,×-Ò-Ùð �qŒ= §¡×!1Ñ!1×!FÑ!FÀyÐRT×!UÒ!UÙáùòCùòùó

s   Â.G Ã.G%Ã?0G*Ä4G*c               ó(   € V ^8„  d   QhRRRRRRRR/# r  rR   )rS   s   "rT   rU   rê    s0   € ÷ ñ Ø#ðà#ðð ðð 
ñ	rV   c                ó|  <€ VR,          pSP                   P                  P                  V4      pSP                   P                  P                  V^4      '       g   R# RpV  FV  pSP                   P                  P                  V4      pSP                   P                  P	                  Wu4      '       g   KU   R# 	  R# )rX   FTrÑ   l   ÿÿ )rÀ  rÁ  r  r  Ústatically_known_gt)	rb  rò  r»   Úlast_strideÚlast_stride_hintÚ
MAX_UINT32rf  Ú	size_hintr½  s	   &&&     €rT   rþ  Ú+can_use_tma.<locals>._is_tma_compatible_xpu  s�   ø€ ð ˜b•kˆØŸ7™7×+Ñ+×MÑMØó
Ðð �w‰w×Ñ×7Ñ7Ð8HÈ!×LÒLÙð ˆ
ÛˆDØŸ™×(Ñ(×JÑJÈ4ÓPˆIØ�w‰w×Ñ×3Ñ3°I×JÔJÚñ ñ
 rV   c              3  ó4   <"  € T F  pS! V4      x € K  	  R # 5irš   rR   )rÉ   r’  rÿ  s   & €rT   rË   Úcan_use_tma.<locals>.<genexpr>  s   øé € Ð?±h°Ñ)¨!×,Ð,³hùó   ƒFT)Útorch.utils._tritonrç  r¿  r½  r�   )
rá  râ  rä  rç  rô  r½  rî  ró  rÿ  rþ  s
   $d*  @@@@@rT   Úcan_use_tmar  œ  s‡   ý€ õ" :å÷Xð X÷9ñ 9÷9ò 9÷3ò 3÷jð ñ. 	Ó÷ 	5ð 	5ß‹CÔ?±hÓ?�C�CŠCÔ?±hÓ?Ó?÷	5ð 	5á% mÓ4ðrV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   rb  rV  râ  rŠ   rP   rR   )rS   s   "rT   rU   rU   !  s$   € ÷ ñ ØðØ-1ðà	ñrV   c                ó>  € \         P                  ! \         P                  4      P                  p. pV  FY  p\	        V\
        \        P                  34      '       d   WB8”  d    R # K4  VP                  \        P                  ! WB4      4       K[  	  V'       g   R# ^RI
Hp \        V4      ^8X  d
   V^ ,          M\        P                  ! V!  pV'       d&   VP                  P                  P!                  V4      # VP                  P                  P#                  V4      # )FTr¼  )r\   Úiinfor›  rÚ   rŒ   rƒ   r�   r›   ró  ÚLer¿  r½  r^   ÚAndrÀ  rÁ  Úguard_or_falseÚstatically_known_true)rb  râ  Ú	int32_maxÚ
conditionsrf  r½  Ú	conditions   &&     rT   Ú_descriptor_shape_fits_in_int32r'  !  sÕ   € ô —’œEŸK™KÓ(×,Ñ,€IØ€JÛˆÜ�dœS¤%§-¡-Ð0×1Ò1ØÔÚñ  ð ×ÑœeŸhšh tÓ7Ö8ñ ÷ Ùåä!$ Z£°AÔ!5�
˜1–¼5¿9º9ÀjÑ;Q€I÷ ð 	
�‰×Ñ×'Ñ'¨	Ó2ðð �W‰W×Ñ×3Ñ3°IÓ>ðrV   c               ó(   € V ^8„  d   QhRRRRRRRR/# ©rO   rä  rB   rá  rC   râ  rŠ   rP   rR   )rS   s   "rT   rU   rU   :  s2   € ÷ Oñ OØðOØ&,ðOØ:>ðOà	ñOrV   c                ó`  a€ \         P                  P                  '       g   R # \        ;QJ d    R V 4       F  '       d   K   R M	  RM! R V 4       4      '       g   R # \        ;QJ d    V3R lV 4       F  '       d   K   R M	  RM! V3R lV 4       4      '       g   R # \         P                  P                  '       d    \        V P                  SR7      '       g   R # \        P                  P                  e   R# \         P                  P                  '       d   T MRp\        VRVRS/ # )Fc              3  óZ   "  € T F!  p\        VP                  4       4      ^8H  x € K#  	  R# 5i)rO   N)r^   rø  )rÉ   r’  s   & rT   rË   Ú*use_triton_tma_template.<locals>.<genexpr>?  s    é € Ð8©x¨!Œs�1—:‘:“<Ó  AÖ%«xùs   ‚)+Tc              3  óZ   <"  € T F   p\        VP                  4       SR 7      x € K"  	  R# 5i)©râ  N)r'  rø  )rÉ   r’  râ  s   & €rT   rË   r,  A  s(   øé € ð áˆAô 	(¨¯
©
«À×LÑLÛùs   ƒ(+r.  Nrá  râ  )rv   r±  Úenable_persistent_tma_matmulr�   Úenable_template_tma_storer'  rf  r\   rÝ  r~   r  )rá  râ  rä  r¶  s   $d* rT   Úuse_triton_tma_templater1  :  sÝ   ø€ ô �=‰=×5×5Ð5Ùß‹3Ñ8©xÓ8�3�3Š3Ñ8©xÓ8×8Ò8Ùß‹3ô áó�3�3Š3ô áó÷ ò ñ Ü‡}�}×.×.Ð.Ô7VØ×Ñ z÷8ó 8ñ ô ‡}�}×ÑÒ$ÙÜ$Ÿm™m×E×EÐE‰]È4€FÜ˜ÐN°ÐNÀ:ÑNÐNrV   c               ó(   € V ^8„  d   QhRRRRRRRR/# r)  rR   )rS   s   "rT   rU   rU   R  s2   € ÷ Vñ VØðVØ&,ðVØ:>ðVà	ñVrV   c                óp   € \        VR V RV/ '       g   R# ^ RIHp ^RIHp V! 4       ;'       d    V! 4       # )rá  râ  F)Ú%has_triton_tensor_descriptor_host_tma©Úis_datacenter_blackwell_arch)r1  r  r4  Úcodegen.cuda.cuda_envr6  )rá  râ  rä  r4  r6  s   $$*  rT   Ú!use_triton_blackwell_tma_templater8  R  sF   € ô #Ø	ðØ!.ðØ;E÷ð ñ åIåCñ 1Ó2×UÐUÑ7SÓ7UÐUrV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )rO   Úscale_option_ar9   Úscale_option_bÚscaling_typeszlist[ScalingType]rP   rŠ   rR   )rS   s   "rT   rU   rU   b  s6   € ÷ Oñ OØðOàðOð %ðOð 
ñ	OrV   c                ó"   € W9   ;'       d    W9   # rš   rR   )r:  r;  r<  s   &&&rT   Úuse_triton_scaling_templater>  b  s   € ð
 Ñ*×NÐN¨~Ñ/NÐNrV   )Úmaxsizec               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU   k  s   € ÷ 	ñ 	˜tñ 	rV   c                 ój   €  \         P                  P                  R4      RJ#   \         d     R# i ; i)z¬Check 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=  rR   rV   rT   Úensure_cute_availablerG  j  s3   € ðÜ�~‰~×'Ñ'¨	Ó2¸$Ð>Ð>øÜô Úðúó   ‚ # £2±2c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU   x  s   € ÷ ñ ¨Dñ rV   c                 óž   €  \         P                  P                  R4      RJp T '       d   \        4        T #   \        \        3 d     R# i ; i)zåCheck if NVIDIA Universal GEMM (cutlass.operators) is importable; cache the result for reuse.

Call ensure_nv_universal_gemm_available.cache_clear() after installing
cutlass.operators in the same interpreter to retry the import.
zcutlass.operatorsNF)rD  rE  rF  r=  Ú
ValueErrorÚ_ensure_fp4_dtype_registered)Ú	availables    rT   Ú"ensure_nv_universal_gemm_availablerN  w  sK   € ðÜ—N‘N×,Ñ,Ð-@ÓAÈÐMˆ	÷ Ü$Ô&ØÐøô	 œÐ$ô Úðús   ‚!7 ·AÁAc                 óÌ   aa€ ^ RI Hu Hu Hp   V P	                  \
        P                  4       R#   \        \        3 d"    ^ RI	oT P                  oTT3R lpYn         R# i ; i)a   Ensure cutlass.operators maps torch.float4_e2m1fn_x2 -> cutlass.Float4E2M1FN.

cutlass.operators natively supports this dtype, so this is normally a no-op.
We keep the patch as a safety net (and for backward compat with generated
code that calls this) in case a future version regresses.
Nc                óV   <€ V \         P                  8X  d   SP                  # S! V 4      # rš   )r\   r�  ÚFloat4E2M1FN)r»   Ú_origrB  s   &€€rT   Ú_patchedÚ._ensure_fp4_dtype_registered.<locals>._patched—  s'   ø€ Øœ×.Ñ.Ô.Ø×+Ñ+Ð+Ù˜“<ÐrV   )
Úcutlass.operators.utils.dtypeÚ	operatorsr}  r»   Úcutlass_type_from_torch_typer\   r�  ÚKeyErrorrœ  rB  )Ú_dtype_utilsrS  rR  rB  s     @@rT   rL  rL  ‡  sS   ù€ ÷ 9Ó8ð=Ø×1Ñ1´%×2HÑ2HÖIøÜ”nÐ%ô 
=Ûà×9Ñ9ˆö	 ð
 5=×1ð
=ús   �1 ±.A#Á"A#c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU      s   € ÷ ñ ¨dñ rV   c                 ój   €  \         P                  P                  R4      RJ#   \         d     R# i ; i)a3  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.
ÚnvMatmulHeuristicsNFrC  rR   rV   rT   Ú$ensure_nvmatmul_heuristics_availabler]  Ÿ  s4   € ðÜ�~‰~×'Ñ'Ð(<Ó=ÀTÐIÐIøÜô ÚðúrH  c               ó<   € V ^8„  d   QhRRRRRRRRRRR	R
RR
RR
RR/	# )rO   Úmat_ar   Úmat_br¶  rC   Úa_is_2drŠ   Úb_is_2dÚoffsz
Any | NoneÚbiasÚscale_resultrP   rR   )rS   s   "rT   rU   rU   ¯  sb   € ÷ >ñ >Øð>àð>ð ð>ð ð	>ð
 ð>ð ð>ð ð>ð ð>ð 
ñ>rV   c                óV  € \        4       '       g   R# \        R4      '       g   R# ^RIHp \	        VP
                  P                  4      '       g   R# V! 4       '       g   R# \        P                  .p	\        W)4      '       g   R# \        P                  '       g   \        P                  '       g   R# \        WVR7      '       g   R# \        ;QJ d    R W3 4       F  '       g   K   RM	  RM! R W3 4       4      '       d   R# V'       d	   V'       d   R# Vf   R# Vf   Ve   R# R# )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ÚCUTEDSLr5  )rá  c              3  ó8   "  € T F  p\        V4      x € K  	  R # 5irš   )Ú
is_dynamic©rÉ   ra   s   & rT   rË   Ú3use_blackwell_cutedsl_grouped_mm.<locals>.<genexpr>á  s   é € Ð
1¡.˜QŒ:�a�=ˆ=£.ùrÍ   T)rG  rÄ  r7  r6  r»  r¼   rE  r\   r‘  r¼  rv   rÜ  rÝ  r  rÐ  )
r_  r`  r¶  ra  rb  rc  rd  re  r6  rß  s
   &&&&&&&&  rT   Ú use_blackwell_cutedsl_grouped_mmrl  ¯  sÔ   € ô2 !×"Ò"Ùä  ×+Ò+ÙåCä�&—-‘-×$Ñ$×%Ò%Ùá'×)Ò)Ùä—^‘^Ð$€MÜ  ×7Ò7Ùä××Ð¤6×#;×#;Ð#;Ùô �u°6×:Ó:Ùç
ƒsÑ
1 5¡.Ó
1‡s‡s‚sÑ
1 5¡.Ó
1×1Ò1Ùç—gÙà‚|ÙàÒ˜<Ò3ÙárV   c          
     ó,   € V ^8„  d   QhRRRRRRRRRR/# ©	rO   r¶  rC   r’  rƒ   r   ÚkrP   rŠ   rR   )rS   s   "rT   rU   rU   ð  s/   € ÷ ñ  ð ¨Cð °Cð ¸Cð ÀDñ rV   c                óž  € ^RI Hp VP                  P                  P	                  W,          V,          RR7      pV^ 8:  g    V\
        P                  P                  8  d   R# ^RIH	p \        P                  P                  '       d   R# \        P                  \        P                  \        P                  .p\!        W4      ;'       d<    \
        P"                  ;'       g    \
        P$                  ;'       d    \'        R4      pV'       d>   V! 4       '       g1   \(        P+                  R\
        P                  P,                  4       R# V# )rX   r¼  ©rÚ  F)Útry_import_cutlassÚCUTLASSzŽFailed 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À  rÁ  rä  rv   rB  Úcutlass_backend_min_gemm_sizeÚcodegen.cutlass.utilsrr  r\   rÝ  r~   rÔ   r‘  r›  r¼  rÜ  rÝ  rÄ  rä   rµ  Úcutlass_dir)	r¶  r’  r   ro  r½  Ú	gemm_sizerr  rß  rú   s	   &&&&     rT   Úuse_cutlass_templaterx  ð  sï   € Ýà—‘× Ñ ×2Ñ2°1µ5¸1µ9ÀrÐ2ÓJ€IØ�A„~˜¤V§^¡^×%QÑ%QÔQÙÝ9ô ‡}�}××ÐÙô —]‘]¤E§N¡N´E·K±KÐ@€Mä˜fÓ4÷ 	-ð 	-Ü× Ñ ×<Ð<¤F×$<Ñ$<÷	-ð 	-ä! )Ó,ð ÷ Ù!×#Ò#Ü�K‰Kð4ô —‘×*Ñ*ô	ñ Ø€JrV   c               ó<   € V ^8„  d   QhRRRRRRRRRRR	RR
RRRRR/	# )rO   r¶  rC   r’  Ú_IntLiker   ro  r_  rB   r`  rc  zIRNode | Nonerõ  z_IntLike | NonerP   rŠ   rR   )rS   s   "rT   rU   rU   	  sm   € ÷ Eñ EØðEàðEð ðEð ð	Eð
 ðEð ðEð ðEð ðEð 
ñErV   c                óø  a
a€ ^ RI Ho \        4       '       g   R# \        4       '       g   R# \	        R4      '       g   R# ^RIHo
 S
P                  '       d   R# V P                  P                  R8w  g!   \        P                  P                  '       d   R# \        P                  '       g   \        P                  '       g   R# WV.pVe   VP!                  V4       \"        ;QJ d    V3R lV 4       F  '       g   K   RM	  RM! V3R lV 4       4      '       d   R# WE.p	Ve   V	P!                  V4       \"        ;QJ d    V
3R lV	 4       F  '       g   K   RM	  RM! V
3R lV	 4       4      '       d   R# R# )	aý  
Return True if we can use the NVIDIA Universal GEMM Template.

Required conditions:
    1. NVGEMM backend is enabled
    2. cutlass.operators 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.operators.get_operators() 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).
)Úhas_free_unbacked_symbolsFÚNVGEMMr¼  rJ   c              3  ó4   <"  € T F  pS! V4      x € K  	  R # 5irš   rR   )rÉ   Údimr|  s   & €rT   rË   Ú1use_nv_universal_gemm_template.<locals>.<genexpr>K	  s   øé € Ð
C±]¨cÑ$ S×)Ð)³]ùr  Tc              3  óp   <"  € T F+  qP                  4       SP                  P                  9   x € K-  	  R # 5irš   )r3  rÀ  rû  )rÉ   Útr½  s   & €rT   rË   r€  S	  s&   øé € Ð
OÑ>N¸�:‰:‹<˜1Ÿ7™7×4Ñ4Ö4Ó>Nùs   ƒ36)r`  r|  rG  rN  rÄ  r¿  r½  Úaot_compilationr¼   rE  r\   rÝ  r~   rv   rÜ  rÝ  ró  rÐ  )r¶  r’  r   ro  r_  r`  rc  rõ  Údims_to_checkÚtensors_to_checkr½  r|  s   &&&&&&&&  @@rT   Úuse_nv_universal_gemm_templater†  	  s  ù€ õ< Pä ×"Ò"Ùä-×/Ò/Ùä  ×*Ò*Ùåà××ÐÙà‡}�}×Ñ˜VÔ#¤u§}¡}×'8×'8Ð'8Ùä××Ð¤6×#;×#;Ð#;Ùð
 ˜1�I€MØ‚}Ø×Ñ˜QÔß
ƒsÔ
C±]Ó
C‡s‡s‚sÔ
C±]Ó
C×CÒCÙð �~ÐØÒØ×Ñ Ô%ß
ƒsÔ
OÑ>NÓ
O‡s‡s‚sÔ
OÑ>NÓ
O×OÒOÙárV   c               ó    € V ^8„  d   QhRRRR/# ©rO   Úop_namerQ   rP   rŠ   rR   )rS   s   "rT   rU   rU   Y	  s   € ÷ Jñ J ð J¨ñ JrV   c                óì   € \         P                  P                  P                  4       pVR8X  d   R# V P                  4       VP	                  R4       Uu. uF  q"P                  4       NK  	  up9   # u upi )z8Check if CUTLASS should be used for the given operation.ÚALLTrÒ  )rv   rB  Úcutlass_enabled_opsr(  r§  r  )r‰  Úenabled_opsra   s   &  rT   Ú_use_cutlass_for_oprŽ  Y	  sY   € ä—.‘.×4Ñ4×:Ñ:Ó<€KØ�eÔÙØ�=‰=‹?°+×2CÑ2CÀCÔ2HÓIÑ2H¨QŸw™wžyÑ2HÑIÑIÐIùÒIs   ÁA1r   rz  c          
     ó,   € V ^8„  d   QhRRRRRRRRRR/# )	rO   r’  rz  r   ro  Úthreshold_multiplerƒ   rP   rŠ   rR   )rS   s   "rT   rU   rU   e	  s4   € ÷ ñ ØðØðØ!)ðØ?Bðà	ñrV   c           
     ó  € ^ RI Hp \        P                  P                  V,          pVP
                  P                  P                  \        P                  ! \        P                  ! W%V ,          4      \        P                  ! W%V,          4      4      4      ;'       dd    VP
                  P                  '       * ;'       dA    VP
                  P                  '       * ;'       d    \        P                  P                  ^ 8„  # )r   r¼  )Útorch._inductor.virtualizedr½  rv   r±  Údecompose_k_thresholdrÀ  rÁ  r#  r�   r!  ÚGera  Úcpp_wrapperÚnum_decompose_k_splits)r’  r   ro  r�  r½  r“  s   &&&&  rT   Úuse_decompose_k_choicer—  d	  sÀ   € õ .ä"ŸM™M×?Ñ?ÐBTÕTÐð 	
�‰×Ñ×.Ñ.Ü�IŠIÜ—’˜°AÕ5Ó6Ü—’˜°AÕ5Ó6óó	
÷ 	5ð 	5ð —‘× Ñ Ô ÷	5ð 	5ð —‘×#Ñ#Ô#÷	5ð 	5ô �M‰M×0Ñ0°1Ñ4ð
rV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )rO   r’  rz  r   ro  rP   rŠ   rR   )rS   s   "rT   rU   rU   z	  s(   € ÷ ñ �hð  8ð °ð ¸Tñ rV   c           
     ó  € \         P                  P                  p^ RIHp \        \        P                  P                  4      ;'       d·    VP                  P                  P                  \        P                  ! \        P                  ! W#V ,          4      \        P                  ! W#V,          4      4      4      ;'       d?    VP                  P                  '       * ;'       d    VP                  P                   '       * # )z€
Check if we should use the contiguous subgraph transform.
This transform makes the second matrix contiguous before the matmul.
r¼  )rv   ÚrocmÚcontiguous_thresholdr’  r½  rŠ   r\   rÝ  r~   rÀ  rÁ  r#  r�   r!  r”  ra  r•  )r’  r   ro  r›  r½  s   &&&  rT   Úuse_contiguousrœ  y	  sº   € ô "Ÿ;™;×;Ñ;Ðõ .ô 	ŒU�]‰]×ÑÓ÷ 	$ð 	$Ø�G‰G×Ñ×2Ñ2Ü�IŠIÜ—’˜°1Õ4Ó5Ü—’˜°1Õ4Ó5óó
÷	$ð 	$ð —‘× Ñ Ô ÷	$ð 	$ð —‘×#Ñ#Ô#ð
rV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )rO   r’  rz  r   ro  rP   rÌ  rR   )rS   s   "rT   rU   rU   ’	  s(   € ÷ 5(ñ 5(�Hð 5( ð 5(¨hð 5(¸9ñ 5(rV   c                ó   € \         P                  P                  p. ROp\        V\        P
                  4      '       d   VP                  '       g   V# V^ 8X  d   . # \        V \        P
                  4      '       d   V P                  '       d3   \        V\        P
                  4      '       d   VP                  '       g   RpM\        W ,          W!,          4      p^p\        P                  ! V4      pV Uu. uF  pW…8:  g   K  W†8¼  g   K  VNK  	  pp. . . rºp	V Ft  pW,,          pV^€8  d   K  WÝ^,
          ,          ^ 8X  d   V^€8¼  d   V	P                  V4       KB  V^ ,          ^ 8X  d   V
P                  V4       Kc  VP                  V4       Kv  	  \         P                  R8X  d   Wš,           V,           # Wš,           V,           pVRV # u upi )é   é   Ú
EXHAUSTIVEN)rŸ  é    é@   é€   r   )rv   r±  r–  rŒ   r�   rw  Ú	is_numberÚminÚdivisorsró  Úmax_autotune_gemm_search_space)r’  r   ro  Ú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   &&&            rT   Úget_k_splitsr´  ‘	  sŒ  € ô —]‘]×9Ñ9€Nò .Ðä�!”U—Z‘Z× Ò ¨¯¯¨ØÐØ	˜1Ô	Øˆ	ä�1”e—j‘j×!Ò!¨!¯+¯+¨+Ü�1”e—j‘j×!Ò!¨!¯+¯+¨+à‰ä˜!�& !¥&Ó)ˆà€Kä�~Š~˜aÓ €Hñ  óáˆGØÑ!ô 	à&-Ñ&<÷ 	ˆÙð ð ð =?ÀÀB¨>ÐãˆØ•ˆð �3Œ;Ùð ˜A•IÕ !Ô#¨°¬Ø×$Ñ$ QÖ'à�R�Z˜1Œ_Ø×%Ñ% aÖ(ð ×!Ñ! !Ö$ñ ô" ×,Ñ,°Ô<Ø Õ5¸ÕFÐFà#Õ8¸>ÕI€Kà�˜Ð'Ð'ùò=s   Ã7GÄGÄGc               ó    € V ^8„  d   QhRRRR/# )rO   r¼   rQ   rP   rR   )rS   s   "rT   rU   rU   Ë	  s   € ÷ @ñ @¨#ð @°#ñ @rV   c                óT   € \         P                  P                  V 4      P                  # rš   )r\   rJ   r“  r¶  ©r¼   s   &rT   Ú_rocm_native_device_arch_namer¸  Ê	  s   € ä�:‰:×+Ñ+¨FÓ3×?Ñ?Ð?rV   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU   Ð	  s   € ÷ ñ ˜$ñ rV   c                 ó‚   € \         P                  P                  4       ;'       d    \        R4      P	                  R4      # )zFReturns true if the device is based on RDNA3, otherwise returns false.rJ   Úgfx11)r\   rJ   r]   r¸  r‚  rR   rV   rT   Úusing_rocm_rdna3r¼  Ï	  s8   € ô �:‰:×"Ñ"Ó$÷ ð Ô)FØó*ç�j�ÓðrV   c               ó   € V ^8„  d   QhRR/# )rO   rP   zNtuple[str | None, Callable[[], list[Any]], Callable[[], list[Any]], type[Any]]rR   )rS   s   "rT   rU   rU   Ø	  s   € ÷ Rñ Rð ñ RrV   c                 óÔ   €  ^ RI p ^ RIHpHp ^ RIHp \        P                  P                  V P                  4      pWAW#3#   \         d    R R lpR R lp ! R R	4      pRp L)i ; i)
r   N)Úgen_ops_libraryÚgen_ops_preselected)ÚCKGemmOperationc               ó   € V ^8„  d   QhRR/# ©rO   rP   rê  rR   )rS   s   "rT   rU   Ú'try_import_ck_lib.<locals>.__annotate__è	  s   € ÷ 	ñ 	 ñ 	rV   c                 ó   € . # rš   rR   rR   rV   rT   r¿  Ú*try_import_ck_lib.<locals>.gen_ops_libraryè	  ó   € ØˆIrV   c               ó   € V ^8„  d   QhRR/# rÃ  rR   )rS   s   "rT   rU   rÄ  ë	  s   € ÷ 	ñ 	 Yñ 	rV   c                 ó   € . # rš   rR   rR   rV   rT   rÀ  Ú.try_import_ck_lib.<locals>.gen_ops_preselectedë	  rÇ  rV   c                  ó   € ] tR tRtRtR# )Ú*try_import_ck_lib.<locals>.CKGemmOperationiî	  rR   N)r    r¡   r¢   r£   r¨   rR   rV   rT   rÁ  rÌ  î	  s   † ÛrV   rÁ  )Úck4inductorÚ(ck4inductor.universal_gemm.gen_instancesr¿  rÀ  Úck4inductor.universal_gemm.oprÁ  r¥  r¶  ÚdirnameÚ__file__r=  )rÍ  r¿  rÀ  rÁ  Úpackage_dirnames        rT   Útry_import_ck_librÓ  ×	  sh   € ðÛ÷	
õ	
ô Ÿ'™'Ÿ/™/¨+×*>Ñ*>Ó?ˆð Ð-@ÐQÐQøô ô õ	õ	÷	ñ 	ð Šðús   ‚;A Á#A'Á&A'c               ó    € V ^8„  d   QhRRRR/# ©rO   r¶  rC   rP   rŠ   rR   )rS   s   "rT   rU   rU   õ	  s   € ÷ "ñ "˜Fð " tñ "rV   c                ó„  € \         P                  '       g   \         P                  '       g   R # \        P                  P
                  '       g   R # V P                  P                  R8w  d   R # \        V P                  4      p\         P                  P                   Uu/ uF  q"P                  R4      ^ ,          VbK  	  up;'       g    VP                  R4      ^ ,          V/pVP                  4       \         P                  P                  ,           Uu. uF  pW2,          NK  	  ppV'       g   R # V P                  \        P                  \        P                   \        P"                  39  d   R # \%        4       w  p  pV'       g   \&        P)                  R4       R # V\         P                  n        R# u upi u upi )FrJ   Ú:z,Please pip install Composable Kernel packageT)rv   rÜ  rÝ  r\   rÝ  r~   r¼   rE  r¸  rš  Úarchr§  rÙ  Úck_supported_archr»   rÔ   r‘  r“  rÓ  rä   rµ  Úck_dir)r¶  Únative_archro  Úrequested_archsÚrequested_supported_archsÚck_package_dirnamerñ   s   &      rT   Úuse_ck_templaterß  õ	  sQ  € ä××Ð¤6×#;×#;Ð#;Ùä�=‰=××ÐÙà‡}�}×Ñ˜VÔ#Ùô 0°·±Ó>€KÜ39·;±;×3CÒ3CÓDÑ3C¨a—w‘w˜s“| A•¨Ò)Ñ3CÑD÷ ð Ø×Ñ˜#Ó˜qÕ! ;ðI€Oð
 !×%Ñ%Ó'¬&¯+©+×*GÑ*GÖGó!áGˆAð 	×ÐÙGð ð !÷ %Ùà‡|�|œEŸM™M¬5¯>©>¼5¿=¹=ÐIÔIÙä"3Ó"5ÑÐ˜˜1˜açÜ�‰ÐBÔCÙà+„F‡K�KÔáùò+ Eùò!s   Â!F8ÄF=c          
     ó,   € V ^8„  d   QhRRRRRRRRRR/# rn  rR   )rS   s   "rT   rU   rU   
  s/   € ÷ ñ  ð ¨Cð °Cð ¸Cð ÀDñ rV   c                óÈ   € ^RI Hp \        R4      ;'       dK    \        V 4      ;'       d8    VP                  P
                  P                  W,          V,          RR7      ^ 8„  # )rX   r¼  ÚCKrq  rÑ   ©r¿  r½  rÄ  rß  rÀ  rÁ  rä  ©r¶  r’  r   ro  r½  s   &&&& rT   Úuse_ck_gemm_templaterå  
  s\   € Ýô 	˜dÓ#÷ 	Kð 	KÜ˜FÓ#÷	Kð 	Kà�G‰G×Ñ×.Ñ.¨q­u°q­yÀ2Ð.ÓFÈÑJðrV   c          
     ó,   € V ^8„  d   QhRRRRRRRRRR/# rn  rR   )rS   s   "rT   rU   rU   $
  s/   € ÷ ñ  fð °ð ¸ð Àð Èñ rV   c                óÈ   € ^RI Hp \        R4      ;'       dK    \        V 4      ;'       d8    VP                  P
                  P                  W,          V,          RR7      ^ 8„  # )rX   r¼  ÚCKTILErq  rÑ   rã  rä  s   &&&& rT   Úuse_ck_tile_gemm_templateré  $
  s\   € Ýô 	˜hÓ'÷ 	Kð 	KÜ˜FÓ#÷	Kð 	Kà�G‰G×Ñ×.Ñ.¨q­u°q­yÀ2Ð.ÓFÈÑJðrV   c               ó    € V ^8„  d   QhRRRR/# rÕ  rR   )rS   s   "rT   rU   rU   .
  s   € ÷ Hñ H ð H¨Dñ HrV   c                ó>   € \        R 4      ;'       d    \        V 4      # )râ  )rÊ  rß  ©r¶  s   &rT   Úuse_ck_conv_templaterí  .
  s   € Ü% dÓ+×GÐG´ÀÓ0GÐGrV   c               ó    € V ^8„  d   QhRRRR/# rÕ  rR   )rS   s   "rT   rU   rU   2
  s   € ÷ &ñ & &ð &¨Tñ &rV   c                ó”   € \         P                  ;'       g    \         P                  ;'       d    V P                  P                  R 8H  # )rC  )rv   rÜ  rÝ  r¼   rE  rì  s   &rT   Ú_use_template_for_cpurð  2
  s:   € ä×Ñ×7Ð7œv×7Ñ7÷&ð &à
�-‰-×
Ñ
 Ñ
%ð&rV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )	rO   r¶  rC   Úmat1zReinterpretView | BufferÚmat2rB   rP   rŠ   rR   )rS   s   "rT   rU   rU   8
  s,   € ÷ ñ ØðØ2ðØ:@ðà	ñrV   c                ór  € ^RI Hp \        VP                  V4      '       g"   \	        R\        VP                  4       24      hVP                  P                  pVP                  P                  p\        V 4      ;'       d}    VP                  4       \        P                  8H  ;'       dT    \        V4      ^8H  ;'       d>    \        V4      ^8H  ;'       d(    V^,          V^,          8H  ;'       d    V^,          ^8H  p\        WVRR7      ;'       d%    VP                  P                  4       ;'       g    T# )rX   )rC   zExpected Layout, got F)Úrequire_constant_mat2)rv  rC   rŒ   r¶  r_   rE  rf  rÃ  rð  rú  r\   r“  r^   Úuse_cpp_gemm_templateÚis_contiguous)r¶  rò  ró  rC   Ú	mat1_sizeÚmat1_strideÚmat1_each_batch_is_contiguouss   &&&    rT   Úuse_cpp_bmm_templaterû  8
  s  € õ ä�d—k‘k 6×*Ò*ÜÐ4´T¸$¿+¹+Ó5FÐ4GÐHÓIÐIð
 —‘× Ñ €IØ—+‘+×$Ñ$€Kä˜fÓ%÷ 	"ð 	"Ø�N‰NÓ¤§¡Ñ-÷	"ð 	"ä�‹^˜qÑ ÷	"ð 	"ô �Ó Ñ"÷	"ð 	"ð ˜�^˜y¨�|Ñ+÷		"ð 	"ð
 ˜�^˜qÑ ð "ô ! ¨tÈ5ÔQ÷ ð Ø�‰×!Ñ!Ó#×DÐDÐ'DðrV   c               ó8   € V ^8„  d   QhRRRRRRRRRRR	RR
RRR/# )rO   r¶  rC   rò  rB   ró  Úmat2_transposedrŠ   rõ  Úis_woq_int4Úq_group_sizer§  rP   rR   )rS   s   "rT   rU   rU   R
  sX   € ÷ =ñ =Øð=à
ð=ð ð=ð ð	=ð
  ð=ð ð=ð ð=ð 
ñ=rV   c                ó  € ^RI Hp ^RIHp ^RIHp	 ^RIHp
 \        V 4      '       d   \        R4      '       g   R# \        P                  P                  '       g   R# VP                  4       \        P                  \        P                   39   p\        P"                  \        P$                  \        P&                  \        P                  \        P                   .pT
! TTV'       d   V P(                  MRVVR7      w  rÞrðr\+        Wï34      '       d   R# \-        W'P.                  4      '       d   VP1                  4       pV	! VP                  4       4      w  ppV! R	VVVVP                  4       VP                  4       V\3        4       V'       * VR
7
      pR R lpV P(                  V9   ;'       dY    VRJ;'       dM    V! V4      ;'       d=    \-        W'P4                  4      ;'       d     VP7                  4       ;'       g    V'       * # )rX   ru  )Úcreate_micro_gemm)Ú*get_gemm_template_output_and_compute_dtype)Úmm_argsÚCPPFN)Ú	out_dtyperý  Úuse_4x2_dimÚ
micro_gemm)Úinput_dtypeÚinput2_dtypeÚoutput_dtypeÚnum_threadsÚuse_refrÿ  c               ó    € V ^8„  d   QhRRRR/# )rO   ra   rB   rP   rŠ   rR   )rS   s   "rT   rU   Ú+use_cpp_gemm_template.<locals>.__annotate__…
  s   € ÷ 'ñ '˜vð '¨$ñ 'rV   c                óV   € V P                  4        V P                  4       R,          ^8H  # )rX   rÑ   )Úfreeze_layoutrù  ©ra   s   &rT   Úis_last_dim_stride1Ú2use_cpp_gemm_template.<locals>.is_last_dim_stride1…
  s"   € Ø	�‰ÔØ�|‰|‹~˜bÕ! QÑ&Ð&rV   )rw  rv  Úcodegen.cpp_micro_gemmr  Úcodegen.cpp_utilsr  Úkernel.mm_commonr  rð  rÄ  rv   ÚcppÚweight_prepackrú  r\   rŸ  r—  r“  r‘  Úhalfr»   Úhas_free_symbolsrŒ   ÚBaseViewÚunwrap_viewÚparallel_num_threadsr³  Úis_module_buffer)r¶  rò  ró  rý  rõ  rþ  rÿ  rv  r  r  r  Ú	int8_gemmrß  r’  r   ro  r
  rñ   r  r  s   &&&&&&&             rT   rö  rö  R
  s¥  € õ Ý9ÝMÝ)ä  ×(Ò(Ô0EÀe×0LÒ0LÙä�:‰:×$×$Ð$Ùà—‘Ó ¤U§[¡[´%·*±*Ð$=Ñ=€IÜ—]‘]¤E§N¡N´E·J±JÄÇÁÌUÏZÉZÐX€MÙ")ØØß"+�&—,’,°Ø'Øô#Ñ€Aˆ!�Tô ˜˜×ÒÙä�$Ÿ™×$Ò$Ø×ÑÓ!ˆá@ÀÇÁÓAQÓR�O€L�!Ù"ØØ	Ø	Ø	Ø—N‘NÓ$Ø—^‘^Ó%Ø!Ü(Ó*Ø”Ø!ô€Jõ'ð
 	�‰˜Ñ%÷ 	Cð 	CØ˜dÐ"÷	Cð 	Cá Ó%÷	Cð 	Cô �tŸ]™]Ó+÷	Cð 	Cð ×"Ñ"Ó$×AÐAÐ,AÔ(AðrV   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU   ’
  s   € ÷ 'ñ '˜tñ 'rV   c                 ó‚   € \         P                  ;'       g    \         P                  '       * ;'       g    \        R 4      # )ÚATEN)rv   rÜ  rÝ  rÄ  rR   rV   rT   Úuse_aten_gemm_kernelsr#  ’
  s5   € ä×Ñ×7Ð7œv×7Ñ7ô÷ 'ð 'ä	˜vÓ	&ð'rV   c                  ój   € ] tR tRt$ ]P
                  ! ^ 4      tR]R&   R R ltR R lt	R R	 lt
R
tR# )ÚDebugDirManageri˜
  rQ   Úprev_debug_namec               ó   € V ^8„  d   QhRR/# ry  rR   )rS   s   "rT   rU   ÚDebugDirManager.__annotate__œ
  s   € ÷ 0ñ 0˜$ñ 0rV   c                	óB   € \        \        P                  4      V n        R # rš   )re  r%  Úcounterr¬   r  s   &rT   r  ÚDebugDirManager.__init__œ
  s   € Ü”×.Ñ.Ó/ˆŽrV   c               ó   € V ^8„  d   QhRR/# ry  rR   )rS   s   "rT   rU   r(  Ÿ
  s   € ÷ <ñ <˜4ñ <rV   c                	óê   € \         P                  P                  P                  V n        V P                   R V P
                   2V n        V P                  \         P                  P                  n        R# )Ú_tmp_N)r\   Ú_dynamorv   Údebug_dir_rootr&  r¬   Únew_namer  s   &rT   Ú	__enter__ÚDebugDirManager.__enter__Ÿ
  sM   € Ü$Ÿ}™}×3Ñ3×BÑBˆÔØ×/Ñ/Ð0°°d·g±g°YÐ?ˆŒØ.2¯m©mŒ�‰×ÑÖ+rV   c               ó    € V ^8„  d   QhRRRR/# )rO   r’   rk  rP   r  rR   )rS   s   "rT   rU   r(  ¤
  s   € ÷ Cñ C˜fð C¨ñ CrV   c                	ó˜   € \         P                  ! V P                  4       V P                  \        P
                  P                  n        R # rš   )rÂ  rÃ  r1  r&  r\   r/  rv   r0  )rH  r’   s   &*rT   Ú__exit__ÚDebugDirManager.__exit__¤
  s*   € Ü�Š�d—m‘mÔ$Ø.2×.BÑ.BŒ�‰×ÑÖ+rV   )r¬   r1  r&  N)r    r¡   r¢   r£   r  rª  r*  r²   r  r2  r6  r¨   rR   rV   rT   r%  r%  ˜
  s,   ‡ Ø�oŠo˜aÓ €GØÓõ0õ<÷
Cñ CrV   r%  c               ó(   € V ^8„  d   QhRRRRRRRR/# ©	rO   r´   úCallable[P, _T]r’   rD  rë  rE  rP   ztuple[_T, list[str]]rR   )rS   s   "rT   rU   rU   ©
  s0   € ÷ &ñ &Øð&àð&ð ð&ð ñ	&rV   c                ó&  a€ ^RI Hp \        4       oR V3R llp\        P                  P                  VRV4      ;_uu_ 4        \        P                  P                  4        V ! V/ VB pRRR4       X\        S4      3#   + '       g   i     L; i)rX   r>   c               ó    € V ^8„  d   QhRRRR/# ©rO   ÚcoderQ   rP   r  rR   )rS   s   "rT   rU   Ú&run_and_get_code.<locals>.__annotate__²
  s   € ÷ ñ ˜sð  tñ rV   c                ó*   <€ SP                  V 4       R # rš   )r  ©r>  Úsource_codess   &€rT   Úsave_output_codeÚ*run_and_get_code.<locals>.save_output_code²
  s   ø€ Ø×Ñ˜ÖrV   rC  N)
rÀ  r?   r"   r   Úpatchrk  r\   r/  Úresetr©  )r´   r’   rë  r?   rC  r1  rB  s   &*,   @rT   Úrun_and_get_coderG  ©
  sw   ø€ õ
 %ä$.£L€L÷ð ô 
�‰×	Ñ	˜=Ð*<Ð>N×	OÕ	OÜ�‰×ÑÔÙ�TÐ$˜VÑ$ˆ÷ 
Pð ”4˜Ó%Ð%Ð%÷ 
P×	Oús   Á'B Â B	c               ó(   € V ^8„  d   QhRRRRRRRR/# r9  rR   )rS   s   "rT   rU   rU   »
  s,   € ÷ ñ ØðØ &ðØ2:ðàñrV   c                óæ  € VP                  R R4      p\        V .VO5/ VB w  rE. pV F¿  p\        P                  '       dU   \        P                  P
                  RJd7   VP                  \        P                  ! RV\        P                  4      4       M5VP                  \        P                  ! RV\        P                  4      4       V'       g   K¬  V Uu. uF  qˆ^R NK	  	  ppKÁ  	  WF3# u upi )Úremove_quoteFTzR"TRITON\((.*?)\)TRITON"z	'''.*?'''éýÿÿÿ)
r`   rG  rv   r•  r±  Úautotune_at_compile_timerÞ  rê   ÚfindallÚDOTALL)	r´   r’   rë  rJ  r1  rB  Úkernelsr>  rŠ  s	   &*,      rT   Úrun_and_get_kernelsrP  »
  sÃ   € ð —:‘:˜n¨eÓ4€Lä+¨BÐ@°Ò@¸Ñ@Ñ€FØ€GÛˆÜ××Ð¤&§-¡-×"HÑ"HÐPTÓ"Tð �N‰Nœ2Ÿ:š:Ð&AÀ4ÌÏÉÓSÕTà�N‰Nœ2Ÿ:š: l°D¼"¿)¹)ÓDÔEß‰<Ù29Ó:±'¨˜a “|±'ˆGÐ:‰Gñ ð ˆ?Ðùò ;s   ÃC.c               ó    € V ^8„  d   QhRRRR/# )rO   r´   r  rP   ztuple[Any, list[str]]rR   )rS   s   "rT   rU   rU   Î
  s   € ÷ /ñ /Ð1ð /Ð6Kñ /rV   c                ó,   a € R  V 3R llp\        V4      # )c               ó   € V ^8„  d   QhRR/# )rO   rP   r   rR   )rS   s   "rT   rU   Ú,run_fw_bw_and_get_code.<locals>.__annotate__Ï
  s   € ÷ ñ ˜sñ rV   c                 óR   <€ S! 4       p V P                  4       P                  4        V # rš   )r^  rˆ  )r1  r´   s    €rT   Úrun_with_backwardÚ1run_fw_bw_and_get_code.<locals>.run_with_backwardÏ
  s!   ø€ Ù“ˆØ�
‰
‹×ÑÔØˆrV   )rG  )r´   rV  s   f rT   Úrun_fw_bw_and_get_coderX  Î
  s   ø€ ÷ð ô
 Ð-Ó.Ð.rV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )	rO   r´   r:  r’   rD  rë  rE  rP   r¯   rR   )rS   s   "rT   rU   rU   ×
  s(   € ÷ (ñ (�ð (¨ð (¸8ð (È	ñ (rV   c           	     ó   aa€ ^RI Hp . oR V3R lloR V3R llp\        P                  P	                  VRV4      ;_uu_ 4        \        P                  P	                  VRS4      ;_uu_ 4        \
        P                  P                  4        V ! V/ VB pRRR4       RRR4       S#   + '       g   i     L; i  + '       g   i     S# ; i)	zLGet the inductor-generated code, but skip any actual compilation or running.r>   c               ó    € V ^8„  d   QhRRRR/# r=  rR   )rS   s   "rT   rU   Úget_code.<locals>.__annotate__Ý
  s   € ÷ "ñ "˜sð " tñ "rV   c                ó*   <€ SP                  V 4       R # rš   ©ró  rA  s   &€rT   rC  Ú"get_code.<locals>.save_output_codeÝ
  s   ø€ Ø×Ñ˜DÖ!rV   c               ó    € V ^8„  d   QhRRRR/# )rO   rH  r?   rP   r   rR   )rS   s   "rT   rU   r\  à
  s   € ÷ ñ ¨ð ¸#ñ rV   c                óæ   <€  ! R  R4      pV P                   '       d   V P                  4       MV P                  4       w  r#S! VP                  4       V'       d   S! VP                  4       V! 4       # )c                  ó2   € ] tR tRtRtR R ltR R ltRtR# )	Ú@get_code.<locals>.patched_compile_to_module.<locals>.DummyModuleiá
  z4This is empty to replace the generated triton modulec               ó   € V ^8„  d   QhRR/# ry  rR   )rS   s   "rT   rU   ÚMget_code.<locals>.patched_compile_to_module.<locals>.DummyModule.__annotate__ä
  s   € ÷ ñ  $ñ rV   c                	ó   € R # rš   rR   r  s   &rT   r  ÚIget_code.<locals>.patched_compile_to_module.<locals>.DummyModule.__init__ä
  s   € ÙrV   c               ó$   € V ^8„  d   QhRRRRRR/# ©rO   r’   r   rë  rP   r  rR   )rS   s   "rT   rU   re  ç
  s!   € ÷ ñ  #ð °ð ¸ñ rV   c                	ó   € R # rš   rR   rG  s   &*,rT   ÚcallÚEget_code.<locals>.patched_compile_to_module.<locals>.DummyModule.callç
  s   € árV   rR   N)r    r¡   r¢   r£   r¤   r  rk  r¨   rR   rV   rT   ÚDummyModulerc  á
  s   † ÙFõ÷ñ rV   rm  )r•  Úcodegen_with_cpp_wrapperÚcodegenr—   )rH  rm  Úwrapper_codeÚkernel_coderC  s   &   €rT   Úpatched_compile_to_moduleÚ+get_code.<locals>.patched_compile_to_moduleà
  s^   ø€ ÷	ñ 	ð 04×/?×/?Ð/?ˆD×)Ñ)Ô+ÀTÇ\Á\Ã^ñ 	"ˆñ 	˜×+Ñ+Ô,ßÙ˜[×.Ñ.Ô/á‹}ÐrV   Úcompile_to_modulerC  N)rÀ  r?   r   rE  rk  r\   r/  rF  )r´   r’   rë  r?   rr  rñ   rC  rB  s   &*,   @@rT   Úget_coderu  ×
  s­   ù€ å$à €L÷"ð "÷ð ô, 	�
‰
×ÑØÐ.Ð0I÷	
õ 	
ô 	�
‰
×Ñ˜-Ð);Ð=M×NÕNä�‰×ÑÔá�Ð˜Ñˆ÷	 	O÷	
ð Ð÷ 	O×Nú÷	
ö 	
ð Ðús$   Á*B<Á0'B)	ÂB<Â)B9Â4B<Â<C	c               ó(   € V ^8„  d   QhRRRRRRRR/# ©	rO   r´   r:  r’   rD  rë  rE  rP   rQ   rR   )rS   s   "rT   rU   rU     s(   € ÷ ñ ˜ð °ð À(ð Èsñ rV   c                ó”   € \        V .VO5/ VB p^\        V4      u;8:  d   ^8:  g   M \        R\        V4       24      hV^ ,          # ©rX   z%expected one or two code outputs got )ru  r^   r_   )r´   r’   rë  rB  s   &*, rT   Úget_triton_coderz    sQ   € ä˜BÐ0 Ò0¨Ñ0€Là”�\Ó"Ö' aÖ'ÜØ3´C¸Ó4EÐ3FÐGó
ð 	
ð ˜�?ÐrV   c               ó(   € V ^8„  d   QhRRRRRRRR/# rw  rR   )rS   s   "rT   rU   rU     s,   € ÷ 
ñ 
Øð
Ø &ð
Ø2:ð
àñ
rV   c                ó˜   € \        V .VO5/ VB w  r4^\        V4      u;8:  d   ^8:  g   M \        R\        V4       24      hV^ ,          # ry  )rG  r^   r_   )r´   r’   rë  rñ   rB  s   &*,  rT   Úrun_and_get_triton_coder}    sU   € ô ' rÐ;¨DÒ;°FÑ;�O€Aà”�\Ó"Ö' aÖ'ÜØ3´C¸Ó4EÐ3FÐGó
ð 	
ð ˜�?ÐrV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )	rO   r´   r:  r’   rD  rë  rE  rP   ztuple[Any, list[GraphLowering]]rR   )rS   s   "rT   rU   rU     s,   € ÷ #ñ #Øð#Ø &ð#Ø2:ð#à$ñ#rV   c                óú   aaa€ ^ RI Ho ^ RIHp VP                  o. oR VVV3R llp\
        P                  P                  VRV4      ;_uu_ 4        V ! V/ VB pRRR4       VS3#   + '       g   i     XS3# ; i)r   r>   rF   c               ó$   € V ^8„  d   QhRRRRRR/# ri  rR   )rS   s   "rT   rU   Ú0run_and_get_graph_lowering.<locals>.__annotate__#  s!   € ÷ &ñ &˜ð &¨ð &°ñ &rV   c                 óž   <€ S! V / VB  V ^,          p\        VS4      '       g   \        R\        V4       24      hSP                  V4       R# )rO   zExpected GraphLowering, got N)rŒ   r_   rE  ró  )r’   rë  rÀ  r?   Úgraph_loweringsÚ	real_inits   *, €€€rT   Ú	fake_initÚ-run_and_get_graph_lowering.<locals>.fake_init#  sK   ø€ Ù�4Ð"˜6Ò"Ø�Q•ˆÜ˜% ×/Ò/Ü Ð#?ÄÀUÃ¸}Ð!MÓNÐNØ×Ñ˜uÖ%rV   r  N)Útorch._inductor.graphr?   Útorch._inductor.output_coderG   r  r   rE  rk  )	r´   r’   rë  rG   r…  r1  r?   rƒ  r„  s	   &*,   @@@rT   Úrun_and_get_graph_loweringr‰    sy   ú€ õ 4Ý;à×(Ñ(€IØ€O÷&ò &ô 
�‰×	Ñ	˜?¨J¸	×	BÕ	BÙ�TÐ$˜VÑ$ˆ÷ 
Cð �?Ð"Ð"÷ 
CÖ	Bð �?Ð"Ð"ús   Á	A'Á'A:	c               ó$   € V ^8„  d   QhRRRRRR/# )rO   Úaten_opr  Úoverride_fnrP   r£  rR   )rS   s   "rT   rU   rU   1  s$   € ÷ .ñ .Øð.Ø.@ð.àñ.rV   c              #  óÒ   "  € ^ RI Hp VP                  V ,          p \        P                  ! W4      VP                  V &   Rx € W2P                  V &   R#   Y2P                  T &   i ; i5i)zs
Override the lowering of aten_op with override_fn.
The first argument of override_fn is the original lowering fn.
)ÚloweringN)Útorch._inductorrŽ  Ú	loweringsrM  Úpartial)r‹  rŒ  rŽ  Úorig_fns   &&  rT   Úoverride_loweringr“  0  sY   é € õ )à× Ñ  Õ)€Gð.Ü&/×&7Ò&7¸Ó&Mˆ×Ñ˜7Ñ#Ûà&-×Ñ˜7Ó#ø g×Ñ˜7Ò#üs   ‚A'�'A ÁA'ÁA$Á$A'c               ó$   € V ^8„  d   QhRRRRRR/# )rO   Úpre_fnr  Úpost_fnzCallable[..., Any] | NonerP   r   rR   )rS   s   "rT   rU   rU   B  s)   € ÷ Fñ FØðFØ)BðFàñFrV   c                ó˜   a aa€ ^ RI Hp VP                  oR VVV 3R llp\        P                  P
                  P                  VRV4      # )zf
Add hook functions to be called at the beginning and end of Scheduler.__init__.
Used for unit tests.
)Ú	Schedulerc               ó$   € V ^8„  d   QhRRRRRR/# )rO   r2  r   rÖ  rP   rR   )rS   s   "rT   rU   Ú-add_scheduler_init_hook.<locals>.__annotate__M  s!   € ÷ ñ ˜3ð  sð ¨sñ rV   c                óH   <€ S! W4       S! W4      pS'       d	   S! W4       V# rš   rR   )r2  rÖ  Úoutr’  r–  r•  s   && €€€rT   r_  Ú(add_scheduler_init_hook.<locals>.wrapperM  s%   ø€ ÙˆyÔ Ù�iÓ'ˆßÙ�IÔ%Øˆ
rV   r  )Útorch._inductor.schedulerr˜  r  Úunittestr   rE  rk  )r•  r–  r˜  r_  r’  s   ff  @rT   Úadd_scheduler_init_hookr   B  s>   ú€ õ 4à× Ñ €G÷ò ô �=‰=×Ñ×%Ñ% i°¸WÓEÐErV   c               ó    € V ^8„  d   QhRRRR/# )rO   ÚmsgrQ   rP   r  rR   )rS   s   "rT   rU   rU   W  s   € ÷ 	ñ 	˜3ð 	 4ñ 	rV   c                óŠ   € \         P                  '       d   \        P                  V 4       R# \        P	                  V 4       R# )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)rv   Údeveloper_warningsrä   rµ  Úinfo)r¢  s   &rT   Údeveloper_warningr¦  W  s'   € ô × × Ð Ü�‰�CÖä�‰�ŽrV   c               ó   € V ^8„  d   QhRR/# )rO   rP   r¬  rR   )rS   s   "rT   rU   rU   c  s   € ÷ ñ ˜Jñ rV   c                 ó  €  \         P                  P                  R4      p V ^,           \        \         P                  4      8  du   \        \         P                  V ^,           ,          4      ^ 8”  dI   \         P                  V ^,           ,          ^ ,          R8w  d   \         P                  V ^,           ,          # \         P                   F*  pTP                  R4      '       g   K  T\        R4      R u # 	  R#   \         d     LMi ; i)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--onlyr  z--only=N)r€  ÚargvrB  r^   rK  r‚  )rä  rï  s     rT   Úget_benchmark_namerª  c  s¾   € ð	Ü�h‰h�n‰n˜XÓ&ˆà�!�G”cœ#Ÿ(™(“mÔ#Ü”C—H‘H˜S 1�WÕ%Ó&¨Ô*Ü—‘˜˜q�Õ! !Õ$¨Ô+ä—8‘8˜C !�GÕ$Ð$ô �xŒxˆØ�>‰>˜)×$Ô$Ø”s˜9“~Ð'Ð(Ò(ñ ñ øô ô Ùðús   ‚B7C: Ã:DÄDc               ó    € V ^8„  d   QhRRRR/# ©rO   rÚ  r  rP   rŠ   rR   )rS   s   "rT   rU   rU   „  s   € ÷ &ñ &�=ð & Tñ &rV   c                ój   € \         ;QJ d    R  V  4       F  '       d   K   R# 	  R# ! R  V  4       4      # )c              3  ó*   "  € T F	  q^8H  x € K  	  R# 5i©rX   NrR   rj  s   & rT   rË   Úis_ones.<locals>.<genexpr>…  ó   é € Ð%™u˜!�AŽv›uùó   ‚FT©r�   ©rÚ  s   &rT   Úis_onesrµ  „  ó+   € ß‹3Ñ%™uÓ%�3Œ3Ð%Š3Ð%ˆ3Ñ%™uÓ%Ó%Ð%rV   c               ó    € V ^8„  d   QhRRRR/# r¬  rR   )rS   s   "rT   rU   rU   ˆ  s   € ÷ &ñ &�Mð & dñ &rV   c                ój   € \         ;QJ d    R  V  4       F  '       d   K   R# 	  R# ! R  V  4       4      # )c              3  ó*   "  € T F	  q^ 8H  x € K  	  R# 5i)r   NrR   rj  s   & rT   rË   Úis_zeros.<locals>.<genexpr>‰  r±  r²  FTr³  r´  s   &rT   Úis_zerosr»  ˆ  r¶  rV   c               ó    € V ^8„  d   QhRRRR/# )rO   ÚinputszSequence[torch.Tensor]rP   rŠ   rR   )rS   s   "rT   rU   rU   Œ  s   € ÷ ñ Ð0ð °Tñ rV   c                ój   € \         ;QJ d    R  V  4       F  '       d   K   R# 	  R# ! R  V  4       4      # )c              3  ó¬   "  € T FJ  p\        V\        P                  4      '       g   K%  VP                  \        P                  ! R 4      8H  x € KL  	  R# 5ir  )rŒ   r\   rù  r¼   )rÉ   rã   s   & rT   rË   Ú is_cpu_device.<locals>.<genexpr>�  s9   é € ð áˆDÜ�dœEŸL™L×)ô 	+ˆ�‰”u—|’| EÓ*Ö*Ûùs
   ‚"A©+AFTr³  )r½  s   &rT   Úis_cpu_devicerÁ  Œ  s=   € ß‹3ñ áó�3Œ3ð Š3ð ˆ3ñ áóó ð rV   c               ó    € V ^8„  d   QhRRRR/# )rO   r³  r‰   rP   rõ  rR   )rS   s   "rT   rU   rU   ”  s   € ÷ ñ ˜jð ¨[ñ rV   c                ó¼   € \        V \        P                  4      '       g   \        R 4      hV P                  '       d   \
        P                  # \
        P                  # )z8only support sympy.Expr as input to get_sympy_Expr_dtype)rŒ   r�   rw  r_   r¦   r\   r�  r•  r¸  s   &rT   Úget_sympy_Expr_dtyperÄ  ”  s<   € Ü�cœ5Ÿ:™:×&Ò&ÜÐWÓXÐXØ
‡~‡~€~Ü�{‰{Ðä�}‰}ÐrV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )rO   Úshould_profilerŠ   r’   r   rë  rP   zIterator[Any]rR   )rS   s   "rT   rU   rU   ž  s(   € ÷ ñ  $ð ¨sð ¸cð Àmñ rV   c              /  ó¼   "  € V '       d8   \         P                  P                  ! V/ VB ;_uu_ 4       pVx € R R R 4       R # R x € R #   + '       g   i     R # ; i5irš   )r\   rÛ   rÜ   )rÆ  r’   rë  rõ   s   &*, rT   Úmaybe_profilerÈ  �  sA   é € çÜ�^‰^×#Ò# TÐ4¨V×4Ò4¸ØŠG÷ 5Ñ4õ 	÷ 5×4Ð4üs   ‚1A³A¸AÁA	Á	Ac               ó   € V ^8„  d   QhRR/# r�  rR   )rS   s   "rT   rU   rU   ¦  s   € ÷ ñ ˜cñ rV   c                 ór   € \         P                  P                  p V ^8  d   \        P                  ! 4       p V # rŸ   )rv   r  Úthreadsr\   Úget_num_threads)rË  s    rT   r  r  ¦  s+   € Ü�j‰j× Ñ €GØ�„{Ü×'Ò'Ó)ˆØ€NrV   c               ó   € V ^8„  d   QhRR/# r�  rR   )rS   s   "rT   rU   rU   ®  s   € ÷ Dñ D ñ DrV   c                 óŒ   € ^RI Hp  V ! 4       pTP                  R\        P                  P
                  '       d   ^4      # ^4      # )rX   )Úget_backend_optionsÚ
num_stages)Úruntime.triton_helpersrÏ  rØ  r\   rÝ  r~   )rÏ  Úoptionss     rT   Úget_backend_num_stagesrÓ  ­  s5   € å;á!Ó#€GØ�;‰;�|¬%¯-©-×*;×*;Ð*; QÓCÐCÀÓCÐCrV   c               ó    € V ^8„  d   QhRRRR/# )rO   r»   rõ  rP   r¸   rR   )rS   s   "rT   rU   rU   ¶  s   € ÷ 26ñ 26˜[ð 26¨Uñ 26rV   c                óˆ  € \         P                  P                  P                  P                  R8H  p\         P
                  P                  4       '       d%   \         P                  P                  P                  p\        WR7      pVe   V# \         P                  P                  4       '       g   \        P                  R4       R# ^ RIHpHp \         P                  P                  4       ;'       d"    \         P                  P                  4       R	8¬  pV \         P                   \         P"                  \         P$                  39  d   \'        RV  24      h\(        P*                  ! V4      P,                  P/                  R4      '       d¤   ^ RIHp V! 4       pV \         P                   \         P"                  39   d   V'       d	   V! W4      # \         P                  P                  P                  P                  R8X  d   V! \         P$                  V4      # V! \         P$                  V4      # V \         P                   \         P"                  39   d   V'       d	   V! V 4      # \         P                  P                  P                  P                  R8X  d   V! \         P$                  4      # V! \         P$                  4      # )
úœ
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_tf32z‡get_device_tflops: no Triton fallback available for non-CUDA devices. Returning 0.0; roofline estimates will use memory bandwidth only.r  )Úget_max_simd_tflopsÚget_max_tensorcore_tflopsz8Expected dtype to be float16, bfloat16, or float32, got Ú
clock_rate)Úmax_clock_rate)rŒ  r   )r\   ÚbackendsrJ   ÚmatmulÚfp32_precisionrK   r]   ÚmkldnnÚ
allow_tf32r   rä   rµ  Útriton.testingrÙ  rÚ  r�  rÔ   r‘  r“  r_   ÚinspectÚ	signatureÚ
parametersrØ  Útorch._utils_internalrÜ  )r»   rØ  Úds_topsrÙ  rÚ  ÚSM80OrLaterrÜ  Úsm_clocks   &       rT   Úget_device_tflopsrê  µ  sè  € ô �n‰n×!Ñ!×(Ñ(×7Ñ7¸6ÑA€GÜ‡y�y×Ñ×ÒÜ—.‘.×'Ñ'×2Ñ2ˆÜ˜UÔ4€GØÒØˆä�:‰:×"Ñ"×$Ò$Ü�‰ðPô	
ñ çMä—*‘*×)Ñ)Ó+÷ ð ´·
±
×0PÑ0PÓ0Rð Wñ 1€Kð
 ”U—]‘]¤E§N¡N´E·M±MÐBÔBÜØFÀuÀgÐNó
ð 	
ô ×ÒÐ,Ó-×8Ñ8×<Ñ<¸\×JÒJå8á!Ó#ˆØ”U—]‘]¤E§N¡NÐ3Ô3¿Ù,¨UÓ=Ð=ä�>‰>×Ñ×%Ñ%×4Ñ4¸Ô>Ù,¬U¯]©]¸HÓEÐEá&¤u§}¡}°hÓ?Ð?à”U—]‘]¤E§N¡NÐ3Ô3¿Ù,¨UÓ3Ð3ä�>‰>×Ñ×%Ñ%×4Ñ4¸Ô>Ù,¬U¯]©]Ó;Ð;á&¤u§}¡}Ó5Ð5rV   c               ó   € V ^8„  d   QhRR/# )rO   rP   r¸   rR   )rS   s   "rT   rU   rU   ì  s   € ÷ ñ ˜5ñ rV   c                 óB   € ^RI Hp  V ! 4       pVe   V# ^ RIHp V! 4       # )rÖ  )Údatasheet_dram_bw_gbs)Úget_dram_gbps)Úanalysis.device_inforí  râ  rî  )rí  Úds_bwrî  s      rT   Úget_gpu_dram_gbpsrñ  ë  s%   € õ <á!Ó#€EØÒØˆå,á‹?ÐrV   c               ó   € V ^8„  d   QhRR/# r�  rR   )rS   s   "rT   rU   rU   ü  s   € ÷ Qñ Q˜sñ QrV   c                 óx   € ^ RI Hp  V P                  P                  P	                  ^ 4      P                  R^ 4      # )r   ©r�  Úmax_shared_mem)Útriton.runtimer�  rŽ  r}  r“  rØ  rô  s    rT   Úget_gpu_shared_memoryr÷  ü  s.   € Ý%à�=‰=×Ñ×4Ñ4°QÓ7×;Ñ;Ð<LÈaÓPÐPrV   c               ó   € V ^8„  d   QhRR/# r�  rR   )rS   s   "rT   rU   rU     s   € ÷ ñ ˜#ñ rV   c                 óh  € \         P                  P                  4       '       dy   \         P                  ! R \         P                  P	                  4       4      p \
        P                  ! V 4      pVP                  pVP                  pVf   \        R4      hW2,          # \        P                  R4       ^ # )rJ   z(expected max_threads_per_block to be setz5CUDA is not available; defaulting max num warps to 32)r\   rJ   r]   r¼   rG  r   rŠ  Úwarp_size_or_defaultÚmax_threads_per_blockr_   rä   rå   )r¼   ÚpropsÚ	warp_sizerû  s       rT   Úget_max_numwarpsrþ    s„   € Ü‡z�z×Ñ× Ò Ü—’˜f¤e§j¡j×&?Ñ&?Ó&AÓBˆÜ ×'Ò'¨Ó/ˆØ×.Ñ.ˆ	Ø %× ;Ñ ;ÐØ Ò(Ü Ð!KÓLÐLØ$Õ1Ð1ä‡I�IÐEÔFÙrV   c               ó    € V ^8„  d   QhRRRR/# )rO   Úreduction_typerQ   rP   rŠ   rR   )rS   s   "rT   rU   rU     s   € ÷ 0ñ 0¨ð 0°ñ 0rV   c                ó$   € V P                  R 4      # )Úwelford)r‚  ©r   s   &rT   Úis_welford_reductionr    s   € Ø×$Ñ$ YÓ/Ð/rV   c               ó    € V ^8„  d   QhRRRR/# )rO   r   rQ   rP   rƒ   rR   )rS   s   "rT   rU   rU     s   € ÷ 
ñ 
¨#ð 
°#ñ 
rV   c                ó>   € \        V 4      '       d   ^# V R9   d   ^# ^# )é   )Úargmax_with_valueÚargmin_with_valueÚonline_softmax_reduce)r  r  s   &rT   Úreduction_num_outputsr    s%   € Ü˜N×+Ò+ÙØ	ð ô 
ñ
 árV   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU   !  s   € ÷ (ñ (�$ñ (rV   c                 ó2   € \         P                  ! 4       R 8H  # )ÚLinux)ÚplatformÚsystemrR   rV   rT   Úis_linuxr  !  s   € Ü�?Š?Ó Ñ'Ð'rV   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU   %  s   € ÷ #ñ #�Dñ #rV   c                 ó(   € \         P                  R 8H  # )rx   )r€  r  rR   rV   rT   r�  r�  %  s   € Ü�<‰<˜7Ñ"Ð"rV   c               ó    € V ^8„  d   QhRRRR/# )rO   Úitrr	  rP   rŠ   rR   )rS   s   "rT   rU   rU   )  s   € ÷ Kñ K˜-ð K¨Dñ KrV   c                ój   € \         ;QJ d    R  V  4       F  '       g   K   R# 	  R# ! R  V  4       4      # )c              3  óˆ   "  € T F8  p\        V\        P                  4      ;'       d    VP                  '       * x € K:  	  R # 5irš   )rŒ   r�   rw  r¥  rj  s   & rT   rË   Ú#has_free_symbols.<locals>.<genexpr>*  s,   é € ÐJÁcÀŒz˜!œUŸZ™ZÓ(×<Ð<°·±¬_Ô<Ãcùs
   ‚%A¨ATFr[  )r  s   &rT   r  r  )  s+   € ß‹3ÑJÁcÓJ�3Œ3ÐJŠ3ÐJˆ3ÑJÁcÓJÓJÐJrV   c               ó    € V ^8„  d   QhRRRR/# rW  rR   )rS   s   "rT   rU   rU   -  s   € ÷ ñ �cð ˜dñ rV   c            	     óÆ  € ^RI Hp V  FÔ  p\        W!P                  VP                  VP
                  VP                  VP                  34      '       dX   \        VP                  4       ;'       g    R4      '       g*   \        VP                  4       ;'       g    R4      '       d    R# K¢  \        W!P                  4      '       g   K¿  \        R\        V4       24      h	  R# )rX   ru  Tzunexpected type for is_dynamic FrR   )rw  rv  rŒ   r±  r³  r  rÿ  r@   r  Úmaybe_get_sizeÚmaybe_get_striderB   Ú	TypeErrorrE  )r’   rv  r‚  s   *  rT   ri  ri  -  s¹   € ÝãˆÜØ—‘˜bŸm™m¨R¯[©[¸"×:KÑ:KÈRÏYÉYÐW÷
ò 
ô   × 0Ñ 0Ó 2× 8Ð 8°b×9Ò9Ô=MØ×"Ñ"Ó$×*Ð*¨÷>ò >ò ñ>ô ˜AŸy™y×)Ò)ÙäÐ=¼dÀ1»g¸YÐGÓHÐHñ ñ rV   c                  ó   € ] tR tRtRtRtRtR# )ÚPlaceholderiA  ÚKERNEL_NAMEÚDESCRIPTIVE_NAMErR   N)r    r¡   r¢   r£   r   r!  r¨   rR   rV   rT   r  r  A  s   † ð  €Kð *ÔrV   r  c          
     ó,   € V ^8„  d   QhRRRRRRRRR	R
/# )rO   r$  r  rþ  r7   Úinpr  r¢  rQ   rP   r  rR   )rS   s   "rT   rU   rU   K  s4   € ÷  
ñ  
Ø
ð 
Ø"-ð 
Ø4Að 
ØHKð 
à	ñ 
rV   c           	     óÐ  € ^RI Hp \        P                  ! RRR7      ;_uu_ 4       p\        P
                  ! 4       p\        P
                  ! 4       p\        V\        V4      R7      P                  ! V!   \        RVP                   2VR7       \        VP                  VR7       \        P                  ! 4       p\        W4      ;_uu_ 4        V ! VP                  4       RRR4       \        P                  ! 4       V,
          p	V! VP                  4       VP                  P                  4        VP                  4        \        R	VP                   2VR7       \        VP                  VR7       VP!                  4       VP!                  4       8H  p
\"        P%                  R
VVP&                  V
V	4       RRR4       R#   + '       g   i     Lç; i  + '       g   i     R# ; i)rX   )Ústable_topological_sortrî  r‹  )ÚmodeÚencoding)rþ  Ú	fake_modezBefore:
)ÚfileNzAfter:
zZ%s, save before/after graph to %s, graph before/after are the same = %s, time elapsed = %s)Úpattern_matcherr%  r¼  ÚNamedTemporaryFileÚioÚStringIOrj   rf   Ú	propagater  rÀ  r   Únowri   ÚlintÚ	recompileÚgetvaluerä   r¥  rì   )r$  rþ  r#  r¢  r%  rÇ  Ú	before_ioÚafter_ioÚ
start_timeÚtime_elapsedr‚  s   &&&&       rT   Úpass_execution_and_saver7  K  sR  € õ 9ä	×	$Ò	$ØØ÷
õ 
ð 
Ü—K’K“Mˆ	Ü—;’;“=ˆÜ�RÔ#3°CÓ#8Ô9×CÒCÀSÒIÜ�	˜"Ÿ(™(˜Ð$¨1Õ-Üˆb�h‰h˜YÕ'Ü—\’\“^ˆ
Ü# B×,Õ,Ù�—‘ŒN÷ -ä—|’|“~¨
Õ2ˆá §¡Ô)Ø
�‰�‰ŒØ
�‰Œä�˜Ÿ™˜
Ð#¨!Õ,Üˆb�h‰h˜XÕ&Ø×ÑÓ  H×$5Ñ$5Ó$7Ñ7ˆÜ�‰ØhØØ�F‰FØØô	
÷+
ñ 
÷ -×,ú÷
÷ 
ð 
ús%   §B(GÃG	Ã"CGÇGÇGÇG%	c               ó    € V ^8„  d   QhRRRR/# ©rO   Ú	input_bufzBuffer | Operation | NonerP   rŠ   rR   )rS   s   "rT   rU   rU   n  s   € ÷ 	ñ 	Ð)Bð 	Àtñ 	rV   c                óh   € ^RI Hp \        WP                  4      ;'       d    V P	                  4       # )z:
Check if input buffer is a multi-outputs template buffer
ru  )rw  rv  rŒ   r  Úis_multi_outputs_template©r:  rv  s   & rT   r<  r<  n  s2   € õ ô 	�9×/Ñ/Ó0÷ 	2ð 	2Ø×/Ñ/Ó1ðrV   c               ó    € V ^8„  d   QhRRRR/# r9  rR   )rS   s   "rT   rU   rU   z  s   € ÷ ñ Ø(ðà	ñrV   c                óÀ   € ^RI Hp \        WP                  4      ;'       d=    \	        V P
                  4      ^8H  ;'       d    \        V P
                  ^ ,          4      # )zD
Check if input buffer is a output of multi-outputs template buffer
ru  )rw  rv  rŒ   ÚMultiOutputr^   r½  r<  r=  s   & rT   Ú#is_output_of_multi_outputs_templaterA  z  sT   € õ ô 	�9Ÿn™nÓ-÷ 	;ð 	;Ü�	× Ñ Ó! QÑ&÷	;ð 	;ä% i×&6Ñ&6°qÕ&9Ó:ðrV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r¶  zNode | Operation | NonerÉ  ztorch._ops.OperatorBase | NonerP   rŠ   rR   )rS   s   "rT   rU   rU   ‰  s&   € ÷ #ñ #Ø
!ð#à&ð#ð 
ñ#rV   c                óª  € V f   R# ^RI Hp \        WP                  4      ;'       dG    \        WP                  4      '       * ;'       d%    VR J ;'       Eg}    V P
                  VJ ;'       Egf    \        V 4      VP                  J ;'       EdF    \        \        P                  P                  R4      ;'       dD    V P
                  \        P                  P                  P                  P                  8H  ;'       g×    \        \        P                  P                  R4      ;'       dD    V P
                  \        P                  P                  P                  P                  8H  ;'       gh    \        \        P                  P                  R4      ;'       d<    V P
                  \        P                  P                  P                  P                  8H  # )NFru  Úall_to_all_singleÚall_gather_into_tensorÚreduce_scatter_tensor)rw  rv  rŒ   Ú_CollectiveKernelÚ_WaitKernelÚop_overloadrE  ÚFallbackKernelr;  r\   r<  ÚtorchrecrD  ÚdefaultrE  rF  ©r¶  rÉ  rv  s   && rT   Úis_collectiverN  ‰  s  € ð ‚|Ùåô 	�4×-Ñ-Ó.÷ 	3ð 	3Ü˜4§¡Ó0Ô0÷	3ð 	3à�4ˆZ÷ð ð 2˜4×+Ñ+¨rÐ1÷ñ ô 	ˆT‹
�b×'Ñ'Ð'÷ 	
ñ 	
ô
 œŸ	™	×*Ñ*Ð,?Ó@÷ Uð UØ×$Ñ$¬¯	©	×(:Ñ(:×(LÑ(L×(TÑ(TÑT÷ð ô
 œŸ	™	×*Ñ*Ð,DÓE÷ Eð EØ×$Ñ$Ü—9‘9×%Ñ%×<Ñ<×DÑDñE÷ð ô œŸ	™	×*Ñ*Ð,CÓD÷ Yð YØ×$Ñ$¬¯	©	×(:Ñ(:×(PÑ(P×(XÑ(XÑXð/rV   c               ó    € V ^8„  d   QhRRRR/# )rO   r¶  zIRNode | Operation | NonerP   rŠ   rR   )rS   s   "rT   rU   rU   ¯  s   € ÷ (ñ (Ð+ð (°ñ (rV   c                ó<   € ^RI Hp \        V 4      VP                  J # ©rX   ru  )rw  rv  rE  rH  )r¶  rv  s   & rT   Úis_waitrR  ¯  s   € Ýä�‹:˜Ÿ™Ð'Ð'rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   ÚsnoderH   Ú	filter_fnz*Callable[[BaseSchedulerNode], bool] | NonerP   rŠ   rR   )rS   s   "rT   rU   rU   µ  s+   € ÷ 	Qñ 	QØð	Qà9ð	Qð 
ñ	QrV   c                ó"  € ^ RI Hp \        W4      '       dI   \        ;QJ d&    R V P                   4       F  '       g   K   R# 	  R# ! R V P                   4       4      # \        V P                  4      ;'       d    VRJ ;'       g	    V! V 4      # )r   ©ÚGroupedSchedulerNodec              3  ó8   "  € T F  p\        V4      x € K  	  R # 5irš   )Úcontains_collectiverj  s   & rT   rË   Ú&contains_collective.<locals>.<genexpr>¼  s   é € Ð@±<¨aÔ& q×)Ð)³<ùrÍ   TFN)rž  rX  rŒ   rÐ  ÚsnodesrN  r¶  )rT  rU  rX  s   && rT   rZ  rZ  µ  sl   € õ ?ä�%×.Ò.ß‹sÑ@°5·<²<Ó@�sŒsÐ@ŠsÐ@ˆsÑ@°5·<²<Ó@Ó@Ð@ä˜Ÿ™Ó$×PÐP¨)°tÐ*;×*OÐ*O¹yÈÓ?OÐPrV   c               ó    € V ^8„  d   QhRRRR/# )rO   rT  rH   rP   rŠ   rR   )rS   s   "rT   rU   rU   Á  s   € ÷ #ñ #Ð*ð #¨tñ #rV   c                óê   € ^ RI Hp \        W4      '       dI   \        ;QJ d&    R V P                   4       F  '       g   K   R# 	  R# ! R V P                   4       4      # \        V P                  4      # )r   rW  c              3  ó8   "  € T F  p\        V4      x € K  	  R # 5irš   )Úcontains_waitrj  s   & rT   rË   Ú contains_wait.<locals>.<genexpr>Å  s   é € Ð:©\¨”= ×#Ð#«\ùrÍ   TF)rž  rX  rŒ   rÐ  r\  rR  r¶  )rT  rX  s   & rT   r`  r`  Á  sP   € Ý>ä�%×.Ò.ß‹sÑ:¨U¯\ª\Ó:�sŒsÐ:ŠsÐ:ˆsÑ:¨U¯\ª\Ó:Ó:Ð:ä�u—z‘zÓ"Ð"rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r¶  zOperation | NonerÉ  z9torch._ops.OpOverload | Collection[torch._ops.OpOverload]rP   rŠ   rR   )rS   s   "rT   rU   rU   Ê  s+   € ÷ Jñ JØ
ðJàAðJð 
ñJrV   c                óÀ   € ^RI Hp \        V\        P                  P
                  4      '       d   V.p\        WP                  4      ;'       d    V P                  V9   # rQ  )rw  rv  rŒ   r\   râ  rã  rJ  rI  rM  s   && rT   Úis_fallback_oprd  Ê  sI   € õ ä�"”e—j‘j×+Ñ+×,Ò,ØˆTˆÜ�d×-Ñ-Ó.×IÐI°4×3CÑ3CÀrÑ3IÐIrV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )rO   Úbuf_namerQ   Úname_to_bufrì  Úname_to_fused_noderP   r   rR   )rS   s   "rT   rU   rU   Õ  s2   € ÷ Lñ LØðLØ .ðLØDRðLàñLrV   c                óP   € W!V ,          P                   P                  4       ,          # rš   )Údefining_opr3  )rf  rg  rh  s   &&&rT   Úbuf_name_to_fused_snoderk  Õ  s!   € ð ¨(Õ3×?Ñ?×HÑHÓJÕKÐKrV   c                ó   € R # rÖ  rR   ©rT  s   &rT   r×  r×  à  ó   € ±urV   c               ó0   € V ^8„  d   QhRRRRRRRRR	R
RR/# ©rO   rT  rH   Úcollected_node_setzMutableSet[BaseSchedulerNode]rg  zdict[str, SchedulerBuffer]rh  zdict[str, BaseSchedulerNode]Úcriteria_cbzCallable[[Any], bool]rP   r  rR   )rS   s   "rT   rU   rU   Û  sD   € ÷ 
ñ 
Øð
à5ð
ð ,ð
ð 5ð	
ð
 'ð
ð 
ñ
rV   c           	     óÊ   € V! V 4      '       d   R # VP                  V 4       V P                   F1  p\        VP                  W#4      pWa9   d   K!  \	        VVVVVR7       K3  	  R # )N©rr  )r  Úunmet_dependenciesrk  rì   Úfind_recursive_deps_of_node)rT  rq  rg  rh  rr  ÚdepÚdefining_op_for_deps   &&&&&  rT   rv  rv  Û  sf   € ñ �5×ÒÙØ×Ñ˜5Ô!Ø×'Ô'ˆÜ5Ø�H‰H�kó
Ðð Ô4ÙÜ#ØØØØØ#÷	
ó (rV   c                ó   € R # rÖ  rR   rm  s   &rT   r×  r×  ù  rn  rV   c               ó0   € V ^8„  d   QhRRRRRRRRR	R
RR/# rp  rR   )rS   s   "rT   rU   rU   ô  sD   € ÷ ñ Øðà5ðð ,ðð 5ð	ð
 'ðð 
ñrV   c           
     óÂ  € V! V 4      '       d   R # VP                  V 4       V P                  4        F©  pVP                   F–  pVP                  f   \	        R4      hVP                  P                  4       R8X  d   K=  VP                  P                  4       V9  d   K^  W6P                  P                  4       ,          pWq9   d   K†  \        VVVVVR7       K˜  	  K«  	  R # )Nz,user.node is None in collect_defined_kernelsÚOUTPUTrt  )r  Úget_outputsrå  r¶  r_   r3  Úfind_recursive_users_of_node)rT  rq  rg  rh  rr  Úorø  Úuser_ops   &&&&&   rT   r~  r~  ô  s½   € ñ �5×ÒÙØ×Ñ˜5Ô!Ø×ÑÖ ˆØ—G”GˆDØ�y‰yÒ Ü$Ð%SÓTÐTØ�y‰y×!Ñ!Ó# xÔ/ÙØ�y‰y×!Ñ!Ó#Ð+=Ô=ÙØ(¯©×);Ñ);Ó)=Õ>ˆGØÔ,ÙÜ(ØØ"ØØ"Ø'÷ó ó !rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   Údynamo_gm_num_inputsrƒ   Úaot_fw_gm_num_inputsrP   rR   )rS   s   "rT   rU   rU     s'   € ÷ Tñ T°ð TÈCð TÐTWñ TrV   c                ó|   € \         P                  P                  P                  '       d   ^M^ pW,
          V,
          # )zaComputes the number of inputs to the aot fw graph which have fixed addresses (params and buffers))r\   Ú
_functorchrv   Úfunctionalize_rng_ops)r‚  rƒ  Únum_rng_seed_offset_inputss   && rT   Únum_fw_fixed_argumentsrˆ    s6   € ô ×Ñ×$Ñ$×:×:Ð:‰Àð ð  Õ6Ð9SÕSÐSrV   c               ó    € V ^8„  d   QhRRRR/# )rO   Úfx_gr^  rP   rƒ   rR   )rS   s   "rT   rU   rU     s   € ÷  ñ  Ð-ð  °#ñ  rV   c                ó<  € R R lp^ p. pV P                   P                   F>  pVP                  R8X  g   K  V! V4      '       d   VP                  V4       V^,          pK@  	  V\	        \        \        V4      4      4      8w  d   \        RV 24      h\        V4      # )z6
Infers which inputs are static for a backwards graph
c               ó    € V ^8„  d   QhRRRR/# )rO   ra   r8   rP   rŠ   rR   )rS   s   "rT   rU   Ú$count_tangents.<locals>.__annotate__"  s   € ÷ 
ñ 
˜4ð 
 Dñ 
rV   c                óª   € R V P                   9  ;'       d>    RV P                   9  ;'       d'    RV P                   9  ;'       d    RV P                   9  # )ÚtangentsÚbwd_seedÚbwd_base_offsetÚbwd_rng_staterL  r  s   &rT   Úis_saved_tensorÚ'count_tangents.<locals>.is_saved_tensor"  sW   € à˜aŸf™fÑ$÷ .ð .Ø !§&¡&Ñ(÷.ð .à!¨¯©Ñ/÷.ð .ð   q§v¡vÑ-ð		
rV   rô  z6Expected static_arg_idxs to be contiguous from 0, got )rÀ  rÖ  rÉ  ró  r©  r×   r^   r_   )rŠ  r“  Ú	arg_countÚstatic_arg_idxsr   s   &    rT   Úcount_tangentsr—    s“   € õ

ð €IØ€OØ�Z‰Z×ÔˆØ�4‰4�=Ö Ù˜q×!Ò!Ø×&Ñ& yÔ1Ø˜�NŠIñ	 ð œ$œu¤S¨Ó%9Ó:Ó;Ô;ÜØDÀ_ÐDUÐVó
ð 	
ô ˆÓÐrV   c               ó    € V ^8„  d   QhRRRR/# )rO   rŠ  r^  rP   rÌ  rR   )rS   s   "rT   rU   rU   9  s   € ÷ ñ Ð#7ð ¸Iñ rV   c                óê   € . p\        V P                  P                  4       FM  w  r#VP                  R8w  d    V# VP                  P                  R4      '       g   K<  VP                  V4       KO  	  V# )zû
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À  rÖ  rÉ  rì   r‚  ró  )rŠ  Ústatic_idxsrä  r   s   &   rT   Úget_static_bw_input_idxsrœ  9  sh   € ð €KÜ˜DŸJ™J×,Ñ,Ö-‰ˆØ�4‰4�=Ô Øð Ðð �6‰6×Ñ˜Z×(Ô(Ø×Ñ˜sÖ#ñ	 .ð
 ÐrV   c                  óD   € ] tR tRt$ R]R&   R R lt]R R l4       tRtR	# )
Ú	BoxedBooliI  rŠ   r—   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   ÚBoxedBool.__annotate__M  s   € ÷ ñ ˜$ñ rV   c                	ó   € V P                   # rš   )r—   r  s   &rT   Ú__bool__ÚBoxedBool.__bool__M  s   € Ø�z‰zÐrV   c               ó    € V ^8„  d   QhRRRR/# )rO   r  r   rP   zBoxedBool | boolrR   )rS   s   "rT   rU   r   Q  s   € ÷ ñ �Sð Ð-ñ rV   c                	óD   € \        V \        4      '       d
   R V n        V # R # rÖ  )rŒ   rž  r—   r�  s   &rT   ÚdisableÚBoxedBool.disableP  s   € ä�cœ9×%Ò%ØˆCŒIØˆJÙrV   rR   N)	r    r¡   r¢   r£   r²   r¢  rK  r¦  r¨   rR   rV   rT   rž  rž  I  s    ‡ àƒKõð ôó ôrV   rž  c               ó    € V ^8„  d   QhRRRR/# )rO   Úkernel_listr¯   rP   r£  rR   )rS   s   "rT   rU   rU   Y  s   € ÷ ñ ¨ð °~ñ rV   c              #  óâ   a a"  € ^RI Hp VP                  oRR V V3R lllp\        P                  P                  VRV4      ;_uu_ 4        Rx € RRR4       R#   + '       g   i     R# ; i5i)rX   r;   Nc               ó4   € V ^8„  d   QhRRRRRRRRRR	R
RRR/# )rO   rH  r<   Úkernel_namerQ   rq  rç  r¬  ÚgpurŠ   Úcpp_definitionrP   r   rR   )rS   s   "rT   rU   Ú-collect_defined_kernels.<locals>.__annotate__^  sN   € ÷ 
ñ 
Ø"ð
àð
ð ð
ð ð	
ð
 ð
ð #ð
ð 
ñ
rV   c                ó:   <€ SP                  V4       S! WW#WE4      # rš   r^  )rH  r¬  rq  rç  r­  r®  r©  Úorig_define_kernels   &&&&&&€€rT   Údefine_kernelÚ.collect_defined_kernels.<locals>.define_kernel^  s'   ø€ ð 	×Ñ˜;Ô'Ù!Ø˜{°có
ð 	
rV   r²  )NTN)Úcodegen.wrapperr<   r²  r   rE  rk  )r©  r<   r²  r±  s   f  @rT   Úcollect_defined_kernelsrµ  X  sM   ùé € å5à-×;Ñ;Ð÷
ó 
ô 
�‰×	Ñ	Ð/°À-×	PÕ	PÛ÷ 
Q×	P×	PÒ	Püs   „AA/ÁAÁ
A/ÁA,	Á&	A/c               ó    € V ^8„  d   QhRRRR/# ©rO   rì   rQ   rP   rR   )rS   s   "rT   rU   rU   o  s   € ÷ !ñ !¨3ð !°3ñ !rV   c                ó   € V R ,           # )Ú__original__rR   rL  s   &rT   Ú get_cloned_parameter_buffer_namerº  o  s   € Ø�.Õ Ð rV   c               ó    € V ^8„  d   QhRRRR/# )rO   r¼   r¬  rP   rŠ   rR   )rS   s   "rT   rU   rU   s  s   € ÷ ñ �:ð  $ñ rV   c                ó   € V \         9   # rš   )rZ   r·  s   &rT   r»  r»  s  s   € Ø”YÑÐrV   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU   w  s   € ÷ )ñ )�ñ )rV   c                 ó:   € \         P                  P                  RJ# )z,Check if we're running on ROCm/HIP platform.N)r\   rÝ  r~   rR   rV   rT   Úis_rocmr¿  w  s   € ä�=‰=×Ñ DÐ(Ð(rV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )	rO   Úbackend_namerQ   rØ  z	int | strrý  rƒ   rP   ztuple[str, ...] | NonerR   )rS   s   "rT   rU   rU   }  s,   € ÷ 'ñ 'Øð'Ø&ð'Ø36ð'àñ'rV   c                ó   € \        4       '       g   R #  ^ RIHp ^ RIHp V! WV4      pV! V4      P                  / 4      p\        TRR 4      pTf   R # \        T4      #   \        \        \        \        3 d     R # i ; i)N)Ú	GPUTarget)Úmake_backendÚsupported_fp8_dtypes)r%   Útriton.backends.compilerrÃ  Útriton.compiler.compilerrÄ  Úparse_optionsrœ  r=  rX  r  r[   rq  )rÁ  rØ  rý  rÃ  rÄ  rá  rÒ  rÅ  s   &&&     rT   Ú_triton_supported_fp8_dtypesrÉ  |  s‚   € ô ×ÒÙðÝ6Ý9ñ ˜<¨yÓ9ˆÙ˜vÓ&×4Ñ4°RÓ8ˆô # 7Ð,BÀDÓIÐØÒ#ÙÜÐ%Ó&Ð&øô œK¬´9Ð=ô Úðús   ”,A ÁA=Á<A=Útriton_backendÚtriton_archrý  c               ó0   € V ^8„  d   QhRRRRRRRRR	R
RR/# )rO   r»   rõ  r¼   ztorch.device | str | NonerÊ  r¬  rË  zint | str | Nonerý  r§  rP   rŠ   rR   )rS   s   "rT   rU   rU   “  sD   € ÷ 10ñ 10Øð10à%ð10ð ð	10ð
 "ð10ð ð10ð 
ñ10rV   c               óF  € V \         P                  \         P                  \         P                  \         P                  39  d   R # \        V 4      P                  R4      pVEf   Vf   R # \         P                  ! V4      pVP                  R8w  d   R # \         P                  P                  4       '       g   R # VP                  pVf   \         P                  P                  4       p\         P                  P                  V4      p\        4       '       d&   RpVP                  P!                  R4      ^ ,          pM&RpVP"                  ^
,          VP$                  ,           p\'        VRR4      pVf   \)        R4      hVf   VR8X  d   ^@M^ p\+        W#V4      pVf   R # WX9   # )Tr¦  NrJ   r~   r×  rý  z0triton_arch must be provided with triton_backend)r\   r‡  rˆ  r‰  r‹  r¬  Úremoveprefixr¼   rE  rJ   r]   rB  rG  r“  r¿  r¶  r§  r‹  Úminorr[   rK  rÉ  )	r»   r¼   rÊ  rË  rý  Útriton_dtypeÚdevice_indexÚ
propertiesrÅ  s	   &&$$$    rT   Úis_triton_fp8_dtype_supportedrÓ  “  sc  € ð Ü×ÑÜ×ÑÜ×ÑÜ×Ñð	ô ñ ä˜E“?×/Ñ/°Ó4€LàÓØŠ>ÙÜ—’˜fÓ%ˆØ�;‰;˜&Ô ÙÜ�z‰z×&Ñ&×(Ò(Ùà—|‘|ˆØÒÜ Ÿ:™:×4Ñ4Ó6ˆLÜ—Z‘Z×5Ñ5°lÓCˆ
Ü�9Š9Ø"ˆNØ$×0Ñ0×6Ñ6°sÓ;¸AÕ>‰Kà#ˆNØ$×*Ñ*¨RÕ/°*×2BÑ2BÕBˆKÜ˜J¨°TÓ:ˆ	àÒÜÐKÓLÐLØÒØ(¨EÔ1‘B°rˆ	ä7Ø YóÐð Ò#ÙØÑ/Ð/rV   c               ó    € V ^8„  d   QhRRRR/# )rO   r¼   rQ   rP   rŠ   rR   )rS   s   "rT   rU   rU   Ç  s   € ÷ .ñ .˜cð . dñ .rV   c                ó2   € V R 8g  ;'       d    \        V 4      # )Úmps)r»  r·  s   &rT   Údevice_need_guardr×  Ç  s   € Ø�U‰?×-Ð-œv f›~Ð-rV   c               ó    € V ^8„  d   QhRRRR/# )rO   r»   rõ  rP   rŠ   rR   )rS   s   "rT   rU   rU   Ë  s   € ÷ 2ñ 2¸ð 2Èñ 2rV   c                ór  € V \         P                  8X  dF   \         P                  P                  4       '       d"   \         P                  P	                  4       R8  # V \         P                  8X  d'   \         P
                  P                  4       '       d   R# V \         P                  \         P                  39   # )é	   T)rÚ  r   )r\   r‘  rJ   r]   r�  rK   r�  rŠ   r÷  s   &rT   Ú,needs_fallback_due_to_atomic_add_limitationsrÛ  Ë  sq   € Ø”—‘Ô¤5§:¡:×#:Ñ#:×#<Ò#<Ü�z‰z×/Ñ/Ó1°FÑ:Ð:Ø	”%—.‘.Ô	 ¤U§Y¡Y×%;Ñ%;×%=Ò%=ÙàœŸ™¤e§j¡jÐ1Ñ1Ð1rV   c               ó4   € V ^8„  d   QhRRRRRRRRRR	R
RRR/# )rO   rI  rÊ  r   r¬  Ú
self_dtyperõ  Ú	src_dtypeÚsrc_device_typerQ   Úsrc_is_tensorrŠ   rP   rR   )rS   s   "rT   rU   rU   Ô  sN   € ÷ $ñ $Ø&ð$àð$ð ð$ð ð	$ð
 ð$ð ð$ð 
ñ$rV   c                ó°  € V P                   \        P                  P                  P                  \        P                  P                  P
                  39   d   Vf   R# V P                   \        P                  P                  P                  8X  d   RMRpVR V39  ;'       Eg4    T;'       d(    \        V4      ;'       d    \        V4      ;'       Eg    V P                   \        P                  P                  P                  8H  ;'       d~    VR8H  ;'       dq    T;'       dg    VR8H  ;'       dZ    \        P                  P                  ;'       d8    \        P                  P                  ;'       gb    \        4       ^8g  ;'       gM    W8H  ;'       d,    V\        P                  \        P                  39   ;'       g    \        P                   ! 4       # )NFr  r^  rC  )Úoverloadpacketr\   r<  ÚatenÚscatter_reduce_Úscatter_reduceÚscatter_r»  rÛ  rv   r  Úfallback_scatter_reduce_sumÚdynamic_threadsr  rŠ   r�  Ú$are_deterministic_algorithms_enabled)rI  r   rÝ  rÞ  rß  rà  Ú	reduce_tys   &&&&&& rT   Úuse_scatter_fallbackrë  Ô  s¡  € ð 	×"Ñ"Ü�I‰I�N‰N×*Ñ*¬E¯I©I¯N©N×,IÑ,IÐJô	KàÒ"áð ×+Ñ+¬u¯y©y¯~©~×/FÑ/FÔF‰ÈEð ð
 	˜t YÐ/Ñ/÷ 	8ñ 	8à÷ Hð HÜ�Ó'÷Hð Hä<¸YÓG÷		8ñ 	8ð ×&Ñ&¬%¯)©)¯.©.×*HÑ*HÑH÷ Lð LØ %Ñ'÷Lð Là÷Lð Lð   5Ñ(÷Lð Lô —
‘
×6Ñ6÷	Lð Lô
 —‘×+Ñ+÷	8ð KÔ/CÓ/EÈÑ/J÷	8ð 	8ð Ñ'×SÐS¨J¼5¿:¹:ÄuÇ{Á{Ð:SÑ,S÷	8ð 	8ô ×5Ò5Ó7ð!rV   c               ó    € V ^8„  d   QhRRRR/# )rO   rr  r|  rP   r  rR   )rS   s   "rT   rU   rU   û  s   € ÷ !Hñ !HÐ&Að !HÀdñ !HrV   c                óì  € ^ RI HpHp ^ RIHp \        R\        V 4       R24       \        V 4       EF=  w  rE\        RVR R24       WRJ d   \        R4       K(  WQJ d   \        R	4       K:  \        WS4      '       dß   VP                  4       p\        V'       d   R
MR R24       V'       dF   VP                  f   \        R4      h\        RVP                  P                  P                   24       \        R4       VP                  P                   F  p\        V4       K  	  \        R4       VP                  P                    F  p\        V4       K  	  EK)  \#        R\%        V4       24      h	  R# )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.
)ÚDisableReductionÚEnableReduction)ÚSchedulerNodezNode schedule with z nodesr  Ú3r×  zenable reductionzdisable reductionÚredÚpwz scheduler nodeNz-SchedulerNode.node is None for reduction nodezoriginal reduction hint zReadDep:z	WriteDep:zUnrecognized node type: )Útorch._inductor.codegen.simdrî  rï  rž  rð  r  r^   rk  rŒ   Úis_reductionr¶  r_   r²  Úreduction_hintrÊ  rË  rÌ  r  rE  )rr  rî  rï  rð  rä  r¶  Úis_redrw  s   &       rT   Údump_node_schedulerø  û  s2  € ÷
 OÝ7ä	Ð¤ MÓ 2Ð3°6Ð
:Ô;Ü˜}×-‰	ˆÜ��#�a�˜ˆlÔàÓ"ÜÐ$Ö%àÓ%ÜÐ%Ö&Ü˜×,Ò,Ø×&Ñ&Ó(ˆFÜŸf‘U¨$Ð/¨Ð?Ô@ßØ—9‘9Ò$Ü(ØGóð ô Ð0°·±·±×1NÑ1NÐ0OÐPÔQÜ�*ÔØ×'Ñ'×-Ô-�Ü�c–
ñ .ä�+ÔØ×'Ñ'×.Ô.�Ü�c–
ô /ô Ð!9¼$¸t»*¸ÐFÓGÐGó1 .rV   c               ó    € V ^8„  d   QhRRRR/# )rO   rà   rð  rP   rŠ   rR   )rS   s   "rT   rU   rU     s   € ÷ ñ ˜lð ¨tñ rV   c                óŠ   € ^ RI Hp V! V P                  4       \        V P                  4      ,          \
        ,          ^ 8H  4      # )r   )r#  )r`  r#  Ústorage_offsetrù  r»   ÚGPU_ALIGN_BYTES)rà   r#  s   & rT   Útensor_is_alignedrý    s:   € õ Lá Ø	×	Ñ	Ó	 ¤>°&·,±,Ó#?Õ	?Ä?ÕRÐVWÑWóð rV   c               ó    € V ^8„  d   QhRRRR/# )rO   Úexample_inputrð  rP   rŠ   rR   )rS   s   "rT   rU   rU   -  s   € ÷ Lñ L¨|ð LÀñ LrV   c                ó–   € \        V P                  P                  4      '       g   R # \        P                  ;'       g    \        V 4      # rÖ  )r»  r¼   rE  rv   Úassume_aligned_inputsrý  )rÿ  s   &rT   Úshould_assume_input_alignedr  -  s8   € ô �-×&Ñ&×+Ñ+×,Ò,ÙÜ×'Ñ'×KÐKÔ+<¸]Ó+KÐKrV   c               ó   € V ^8„  d   QhRR/# )rO   rP   z'contextlib.AbstractContextManager[None]rR   )rS   s   "rT   rU   rU   6  s   € ÷ 'ñ 'Ð-Tñ 'rV   c                 ó^  € \         P                  P                  P                  4       p V '       g   \        P
                  ! 4       # V P                  '       d   V P                  P                  '       g   \        P
                  ! 4       # V P                  P                  pVP                  4       # rš   )	r\   Ú_guardsÚTracingContextÚtry_getr6  Únullcontextr(  rÂ  Úsuppress_guards)Útracing_contextrÂ  s     rT   Ú#maybe_get_suppress_shape_guards_ctxr  6  s|   € ô
 —m‘m×2Ñ2×:Ñ:Ó<€OßÜ×%Ò%Ó'Ð'ð ×$×$Ð$¨O×,EÑ,E×,O×,OÐ,OÜ×%Ò%Ó'Ð'Ø×)Ñ)×3Ñ3€IØ×$Ñ$Ó&Ð&rV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )	rO   r´   r:  r’   rD  rë  rE  rP   ztuple[_T, str]rR   )rS   s   "rT   rU   rU   F  s,   € ÷ ñ ØðØ &ðØ2:ðàñrV   c                ó6  € \         P                  P                  P                  \        R R4      ;_uu_ 4        \
        P                  P                  4        ^ RIp^ RI	pVP                  ! 4       pVP                  ! V4      p^ RIHp VP                  V4       VP                  pVP!                  VP"                  4       V ! V/ VB p	VP%                  4       p
VP!                  V4       VP'                  V4       RRR4       Wš3#   + '       g   i     X	X
3# ; i)rå   TN)Úoutput_code_log)rŸ  r   rE  rk  rv   r\   r/  rF  r,  r+  r-  ÚStreamHandlerÚtorch._inductor.codecacher  Ú
addHandlerÚlevelÚsetLevelr,  r2  ÚremoveHandler)r´   r’   rë  r,  r+  Úlog_capture_stringÚchr  Ú
prev_levelr1  r÷   s   &*,        rT   Úrun_and_get_cpp_coder  F  sá   € ô 
�‰×	Ñ	×	#Ñ	#¤F¨G°T×	:Õ	:Ü�‰×ÑÔÛÛàŸ[š[›]ÐØ×"Ò"Ð#5Ó6ˆÝ=à×"Ñ" 2Ô&Ø$×*Ñ*ˆ
Ø× Ñ  §¡Ô/Ù�TÐ$˜VÑ$ˆØ×'Ñ'Ó)ˆØ× Ñ  Ô,Ø×%Ñ% bÔ)÷ 
;ð  ˆ9Ð÷! 
;Ö	:ð  �1ˆ9Ðús   ¸CDÄD	c               ó    € V ^8„  d   QhRRRR/# )rO   r½  úSequence[InputType]rP   zShapeEnv | NonerR   )rS   s   "rT   rU   rU   _  s   € ÷ ñ Ð"5ð ¸/ñ rV   c                óJ  € \        V 4      pVe   VP                  # V  EF   p\        V\        P                  4      '       d   VP
                  P                  u # \        V\        P                  4      '       g   K^  VP                  4        F=  p\        V\        P                  4      '       g   K%  VP
                  P                  u u # 	  VP                  4        F=  p\        V\        P                  4      '       g   K%  VP
                  P                  u u # 	  EK  	  R # rš   )	rf   rÂ  rŒ   r\   r5   r¶  rù  rf  rÃ  )r½  r(  Úinputrf  rÃ  s   &    rT   Úshape_env_from_inputsr  _  sÉ   € Ü  Ó(€Ið ÒØ×"Ñ"Ð"ô ˆÜ�eœUŸ\™\×*Ò*Ø—:‘:×'Ñ'Ò'ô �eœUŸ\™\×*Ô*ØŸ
™
ž�Ü˜d¤E§L¡L×1Ô1ØŸ9™9×.Ñ.Ô.ñ %ð  Ÿ,™,ž.�Ü˜f¤e§l¡l×3Ô3Ø!Ÿ;™;×0Ñ0Ô0ô )ñ ñ rV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )rO   r  zCallable[[list[InputType]], _T]Úinputs_to_checkr)  Úmutated_input_idxszOrderedSet[int]rP   rR   )rS   s   "rT   rU   rU   |  s0   € ÷ ñ Ø*ðà"ðð (ðð %ñ	rV   c                óF   a aa€ \        S4      ^ 8X  d   S # R VV V3R llpV# )r   c               ó    € V ^8„  d   QhRRRR/# )rO   Ú
new_inputsúlist[InputType]rP   r   rR   )rS   s   "rT   rU   Ú2align_inputs_from_check_idxs.<locals>.__annotate__„  s   € ÷ ñ ˜ð ¨Cñ rV   c                ó„   <€ \        V SS4      w  rS! V 4      p\        V4      '       d   \        P                  ! W4       V# rš   )Úcopy_misaligned_inputsr^   r\   Ú_foreach_copy_)r#  Úold_tensorsÚnew_tensorsrœ  r  r  r   s   &   €€€rT   r‹  Ú)align_inputs_from_check_idxs.<locals>.run„  sD   ø€ Ü#9Ø˜Ð);ó$
Ñ ˆñ �JÓˆô ˆ{×ÒÜ× Ò  Ô:àˆ
rV   )r^   )r  r  r   r‹  s   fff rT   Úalign_inputs_from_check_idxsr,  |  s(   ú€ ô
 ˆ?Ó˜qÔ Øˆ÷ò ð €JrV   c               ó    € V ^8„  d   QhRRRR/# )rO   ra   rð  rP   rR   )rS   s   "rT   rU   rU   ”  s   € ÷ 	:ñ 	:˜lð 	:¨|ñ 	:rV   c                ób  € ^ V P                  4       9   d   ^ pM?\        R \        V P                  4       V P                  4       4       4       4      ^,           p\        P
                  ! W3R4      P                  4       p\        P
                  ! W P                  4       V P                  4       4      # )r   c              3  óF   "  € T F  w  rV^,
          V,          x € K  	  R# 5ir¯  rR   )rÉ   rº  rÃ  s   &  rT   rË   Ú)clone_preserve_strides.<locals>.<genexpr>š  s    é € ÐTÑ:S©¨�˜•˜f×$Ò$Ó:Sùs   ‚!rŸ   )rf  r^  rá   rÃ  r\   Ú
as_stridedÚclone)ra   Úneeded_sizer­  s   &  rT   Úclone_preserve_stridesr4  ”  s€   € ØˆA�F‰F‹H„}à‰ô ÑT¼#¸a¿f¹f»hÈÏÉË
Ô:SÓTÓTÐWXÕXð 	ô ×Ò˜a °Ó6×<Ñ<Ó>€FÜ×Ò˜F§F¡F£H¨a¯h©h«jÓ9Ð9rV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )	rO   r#  r$  Úcheck_inputs_idxsr)  Úreturn_pair_idxszOrderedSet[int] | NonerP   z-tuple[list[torch.Tensor], list[torch.Tensor]]rR   )rS   s   "rT   rU   rU      s0   € ÷ $ñ $Øð$à$ð$ð -ð$ð 3ñ	$rV   c                ó|  € . p. pVRJpV F¬  pW,          p\        V\        P                  4      '       g   \        R\	        V4       24      hVP                  4       \        ,          '       g   Ke  \        V4      W&   V'       g   K|  Wb9   g   K„  VP                  V4       VP                  W,          4       K®  	  W43# )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Œ   r\   rù  r_   rE  Údata_ptrÚ	ALIGNMENTr4  ró  )r#  r6  r7  r)  r*  Úret_pair_definedrö   Ú_inps   &&&     rT   r'  r'     s¥   € ð ')€KØ&(€Kð (¨tÐ3ÐÛˆØ�}ˆÜ˜$¤§¡×-Ò-Ü Ð#CÄDÈÃJÀ<Ð!PÓQÐQØ�=‰=‹?œY×&Ö&Ü2°4Ó8ˆJ‰MçÑ AÖ$9Ø×"Ñ" 4Ô(Ø×"Ñ" :¥=Ö1ñ ð Ð#Ð#rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r½  r  Ústatic_input_idxsr)  rP   rR   )rS   s   "rT   rU   rU   ½  s&   € ÷ ñ Øðà$ðð ñrV   c                ó  € . pV F`  pW,          p\        V\        P                  4      '       g   K-  VP                  4       \        ,          ^ 8X  g   KO  VP                  V4       Kb  	  \        V4      \        V4      8w  d   V# V# )zO
We require all inputs to be aligned, so introduce a copy for any
that aren't.
)rŒ   r\   rù  r9  r:  ró  r^   )r½  r>  Úaligned_static_input_idxsrä  r  s   &&   rT   Úremove_unaligned_input_idxsrA  ½  sp   € ð !#ÐÛ ˆØ•ˆÜ�eœUŸ\™\×*Ô*°·±Ó0@Ä9Õ0LÐQRÖ/RØ%×,Ñ,¨SÖ1ñ !ô Ð$Ó%¬Ð->Ó)?Ô?Ø(Ð(ØÐrV   c               ó    € V ^8„  d   QhRRRR/# )rO   rø   r‰   rP   rŠ   rR   )rS   s   "rT   rU   rU   Ï  s   € ÷ *Iñ *I˜jð *I¨Tñ *IrV   c                ó�  € ^RI Hp \        P                  ! \        P                  4      P
                  pVP                  P                  P                  pVP                  P                  P                  P                  p\        P                  '       d(   VP                  P                  P                  W4       R# VP                  P                  P                  W8*  4      '       d   R# VP                  '       d1   VP                  P                  P                  V R8  4      '       d   R# V! V 4      ;'       d    V! V 4      V8*  # )zíCheck 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.
r¼  Tg@Œµx¯DF)r¿  r½  r\   r  r›  rÚ   rÀ  rÁ  Úguarding_hint_or_throwrÂ  Úhas_guarding_hintrv   Úassume_32bit_indexingÚ	check_leqr#  rƒ  )rø   r½  Úint_maxrD  rE  s   &    rT   Úexpr_fits_within_32bitrI  Ï  sî   € õ ä�kŠkœ%Ÿ+™+Ó&×*Ñ*€GØŸW™W×-Ñ-×DÑDÐØŸ™×(Ñ(×2Ñ2×DÑDÐä×#×#Ð#Ø	�‰×Ñ×"Ñ" 1Ô.Ùð 	‡w�w×Ñ×-Ñ-¨a©l×;Ò;Ùð 	××Ðð �7‰7×Ñ×1Ñ1°!°d±(×;Ò;ñ ñ ˜QÓ×HÐHÑ$:¸1Ó$=ÀÑ$HÐHrV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r  r  Úcompiled_graphrG   rP   r  rR   )rS   s   "rT   rU   rU   ü  s$   € ÷ ñ Ø!ðØ3Bðà	ñrV   c                óØ  aaa€ \         P                  P                  P                  4       pVEe:   VP                  Ee)   \        VP                  4      ^ 8w  d"   \        R\        VP                  4       24      h\        V 4      oVP                  f   \        R4      hVP                   F·  pVf   VP                  P                  R 4       K$  Ro\         P                  P                  P                  4       ;p'       d   VP                  oR VV3R lloVP                  P                  \        ;QJ d    . V3R lV 4       F  NK  	  5M! V3R lV 4       4      4       K¹  	  R # R # R # )Nz*Expected empty output_strides, got length z%compiled_graph.output_strides is NoneFc               ó    € V ^8„  d   QhRRRR/# )rO   rø   r   rP   z)float | int | SymInt | SymFloat | SymBoolrR   )rS   s   "rT   rU   Ú8set_tracing_context_output_strides.<locals>.__annotate__  s   € ÷ 9ñ 9 ð 9Ð(Qñ 9rV   c                óv   <€ Sf   \        V 4      # S'       d   SP                  V 4      # SP                  V 4      # rš   )rƒ   Údeserialize_symexprÚevaluate_symexpr)rø   Úfakify_first_callrÂ  s   &€€rT   Úmap_exprÚ4set_tracing_context_output_strides.<locals>.map_expr  s7   ø€ Ø Ò(Ü" 1›v˜ß(Ø(×<Ñ<¸QÓ?Ð?Ø$×5Ñ5°aÓ8Ð8rV   c              3  ó4   <"  € T F  pS! V4      x € K  	  R # 5irš   rR   )rÉ   rø   rS  s   & €rT   rË   Ú5set_tracing_context_output_strides.<locals>.<genexpr>  s   øé € Ð5©u¨!™( 1Ÿ+˜+«uùr  )r\   r  r  r  Úoutput_stridesr^   r_   r  ró  rR  rq  )r  rK  Úcontextrò  r^  rR  rS  rÂ  s   &&   @@@rT   Ú"set_tracing_context_output_stridesrY  ü  s&  ú€ ô �m‰m×*Ñ*×2Ñ2Ó4€GØÓ˜w×5Ñ5ÓAÜˆw×%Ñ%Ó&¨!Ô+Ü Ø<¼SÀ×AWÑAWÓ=XÐ<YÐZóð ô *¨.Ó9ˆ	Ø×(Ñ(Ò0Ü Ð!HÓIÐIØ#×2Ô2ˆEØŠ}Ø×&Ñ&×-Ñ-¨dÖ3à$)Ð!ÜŸ-™-×6Ñ6×>Ñ>Ó@Ð@�3Ö@Ø(+×(=Ñ(=Ð%÷9ñ 9ð ×&Ñ&×-Ñ-ß”EÔ5©uÓ5—E‘EÔ5©uÓ5Ó5öó 3ñ  BÑrV   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU     s   € ÷ ñ ¨$ñ rV   c                 óB  € \         P                  e   \         P                  # \         P                  ! 4       '       g   R# \        P                  P                  4       '       d   R#  ^ RIHp  T \        P                  P                  R4      8¬  #   \         d     R# i ; i)NF©ÚREMOTE_CACHE_VERSIONz.pytorch/remote_cache:fx_graph_memcache_version)
rv   Úfx_graph_remote_cacheÚ	is_fbcoder\   Ú_utils_internalÚis_fb_unit_testÚtorch._inductor.fb.remote_cacher]  ÚModuleNotFoundErrorÚjustknobs_getval_intr\  s    rT   Ú should_use_remote_fx_graph_cachere    s…   € Ü×#Ñ#Ò/Ü×+Ñ+Ð+Ü×Ò×ÒÙä×Ñ×,Ñ,×.Ò.ÙðÝHð  ¤5×#8Ñ#8×#MÑ#MØ8ó$ñ ð øô ô Úðús   Á'B ÂBÂBc               ó    € V ^8„  d   QhRRRR/# r·  rR   )rS   s   "rT   rU   rU   0  s   € ÷ /ñ /˜ð / ñ /rV   c                ó2   € \         P                  ! R RV 4      # )z[^a-zA-Z0-9_]rñ   )rê   ÚsubrL  s   &rT   Únormalize_nameri  0  s   € Ü�6Š6Ð" C¨Ó.Ð.rV   ztl.boolztl.int1ztl.float8_e4m3fnztl.float8e4nvztl.float8_e5m2ztl.float8e5ztl.float8_e4m3fnuzztl.float8e4b8ztl.float8_e5m2fnuzztl.float8e5b16ztl.float8_e8m0fnuztl.uint8ztl.float4_e2m1fn_x2z^.*[.]c               ó    € V ^8„  d   QhRRRR/# )rO   r»   rõ  rP   rQ   rR   )rS   s   "rT   rU   rU   F  s   € ÷ Hñ H�{ð H sñ HrV   c                ój   € \         P                  R\        V 4      4      p\        P	                  W4      # )z"Convert torch.dtype to triton typeútl.)Ú_triton_type_rerh  rQ   Ú_triton_type_mappingrØ  )r»   Útriton_type_names   & rT   Útriton_typerp  F  s+   € ä&×*Ñ*¨5´#°e³*Ó=ÐÜ×#Ñ#Ð$4ÓGÐGrV   c               ó    € V ^8„  d   QhRRRR/# )rO   r»   rQ   rP   rõ  rR   )rS   s   "rT   rU   rU   L  s   € ÷ ñ  ð ¨ñ rV   c                óâ   € \         P                  W 4      pVP                  R R4      p\        \        V4      p\        V\        P                  4      '       g   \        R\        V4       24      hV# )rl  rw  zExpected torch.dtype, got )	Ú_torch_triton_mappingrØ  r~  r[   r\   rŒ   r»   r_   rE  )r»   Úadjusted_typeÚ	type_namer  s   &   rT   Útriton_type_to_torchrv  L  s_   € Ü)×-Ñ-¨eÓ;€MØ×%Ñ% e¨RÓ0€IÜœ˜yÓ)€IÜ�i¤§¡×-Ò-ÜÐ9¼$¸y»/Ð9JÐKÓLÐLØÐrV   c               ó$   € V ^8„  d   QhRRRRRR/# ©rO   r²  rð  r—   rP   rŠ   rR   )rS   s   "rT   rU   rU   U  s!   € ÷ 	ñ 	˜ð 	¨lð 	¸tñ 	rV   c                ó(  € V P                   '       * ;'       dû    V P                  4       VP                  4       8H  ;'       dÒ    V P                  4       VP                  4       8H  ;'       d©    V P                  VP                  8H  ;'       dˆ    V P                  VP                  8H  ;'       dg    V P                  4       P                  4       VP                  4       P                  4       8H  ;'       d"    V P                  4       VP                  4       8H  # rš   )Ú	is_mkldnnrf  rÃ  r»   r¼   Úuntyped_storager9  rû  ©r²  r—   s   &&rT   Úis_same_tensorr}  U  sã   € à�N‰NÔ÷ 	<ð 	<Ø�I‰I‹K˜5Ÿ:™:›<Ñ'÷	<ð 	<à�K‰K‹M˜UŸ\™\›^Ñ+÷	<ð 	<ð �J‰J˜%Ÿ+™+Ñ%÷	<ð 	<ð �K‰K˜5Ÿ<™<Ñ'÷		<ð 	<ð
 × Ñ Ó"×+Ñ+Ó-°×1FÑ1FÓ1H×1QÑ1QÓ1SÑS÷	<ð 	<ð ×ÑÓ! U×%9Ñ%9Ó%;Ñ;ðrV   c               ó$   € V ^8„  d   QhRRRRRR/# rx  rR   )rS   s   "rT   rU   rU   a  s!   € ÷ ñ  ð °\ð Àdñ rV   c                ó¦  € V P                   ;'       d¿    V P                  4       VP                  4       8H  ;'       d–    V P                  VP                  8H  ;'       du    V P                  VP                  8H  ;'       dT    \        P
                  P                  P                  V 4      \        P
                  P                  P                  V4      8H  # rš   )rz  rf  r»   r¼   r\   r<  rà  r9  r|  s   &&rT   Úis_same_mkldnn_tensorr€  a  s¨   € à�‰÷ 	Pð 	PØ�I‰I‹K˜5Ÿ:™:›<Ñ'÷	Pð 	Pà�J‰J˜%Ÿ+™+Ñ%÷	Pð 	Pð �K‰K˜5Ÿ<™<Ñ'÷	Pð 	Pô �I‰I×Ñ×%Ñ% dÓ+¬u¯y©y×/?Ñ/?×/HÑ/HÈÓ/OÑOðrV   c               ó   € V ^8„  d   QhRR/# )rO   rP   ztuple[str, ...]rR   )rS   s   "rT   rU   rU   l  s   € ÷ ñ �_ñ rV   c                 ó   € R# )Úisinf)rƒ  ÚisnanÚlogical_notÚlogical_andÚsignbitÚand_ÚleÚltÚgeÚgtÚeqÚnerx  ÚxorrR   rR   rV   rT   Úboolean_opsr�  k  s   € ðð rV   c                  ó,   € ] tR tRt$ R]R&   R]R&   RtR# )ÚOpDtypeRulei  r6   Útype_promotion_kindr|  Úoverride_return_dtyperR   N)r    r¡   r¢   r£   r²   r¨   rR   rV   rT   r’  r’    s   ‡ à8Ó8Ø-×-rV   r’  zdict[str, OpDtypeRule]Úop_dtype_propagation_rulesc               ó(   € V ^8„  d   QhRRRRRRRR/# )	rO   rì   rQ   r“  r6   r”  r|  rP   r  rR   )rS   s   "rT   rU   rU   ˆ  s0   € ÷ ñ Ø
ðà8ðð .ðð 
ñ	rV   c                ó*   € \        W4      \        V &   R # rš   )r’  r•  )rì   r“  r”  s   &&&rT   Ú#register_op_dtype_propagation_rulesr˜  ˆ  s   € ô
 (3Øó(Ô˜tÓ$rV   zOrderedSet[str]Úop_requires_libdevice_fp64c               ó    € V ^8„  d   QhRRRR/# )rO   rì   rQ   rP   r  rR   )rS   s   "rT   rU   rU   •  s   € ÷ )ñ )¨cð )°dñ )rV   c                ó0   € \         P                  V 4       R # rš   )r™  r  rL  s   &rT   Ú#register_op_requires_libdevice_fp64rœ  •  s   € Ü×"Ñ" 4Ö(rV   c               ó    € V ^8„  d   QhRRRR/# )rO   ré   r¬  rP   rQ   rR   )rS   s   "rT   rU   rU   ™  s   € ÷ #ñ # Zð #¸3ñ #rV   c                ó"  € ^ RI Hp V '       g%   VP                  P                  4       P                  p V R8X  d   \
        P                  # V R8X  d   R# V R8X  d   \
        P                  # V R8X  d   \
        P                  # \
        P                  # )r   r¼  rC  rÖ  rK   Útpu)
r’  r½  rÀ  Úget_current_device_or_throwrE  rv   Úcpu_backendÚxpu_backendÚtpu_backendÚcuda_backend)ré   r½  s   & rT   Úget_current_backendr¥  ™  ss   € Ý-çØ—g‘g×9Ñ9Ó;×@Ñ@ˆØ�eÔÜ×!Ñ!Ð!Ø	˜Ô	ÙØ	˜Ô	Ü×!Ñ!Ð!Ø	˜Ô	Ü×!Ñ!Ð!ä×"Ñ"Ð"rV   c               ó    € V ^8„  d   QhRRRR/# )rO   r¼   ztorch.device | NonerP   rŠ   rR   )rS   s   "rT   rU   rU   ª  s   € ÷ ñ Ð!4ð ¸ñ rV   c                ó‚   € V e;   V P                   R8X  d*   \        P                  P                  V 4      P                  # R# )z+Check if the given device supports float64.rK   T)rE  r\   rK   r“  Úhas_fp64r·  s   &rT   Údevice_supports_fp64r©  ª  s1   € àÒ˜fŸk™k¨UÔ2Ü�y‰y×.Ñ.¨vÓ6×?Ñ?Ð?ÙrV   c               ó    € V ^8„  d   QhRRRR/# )rO   r»   rõ  rP   rR   )rS   s   "rT   rU   rU   ±  s   € ÷ ñ ˜{ð ¨{ñ rV   c                óÎ   € V \         P                  \         P                  39   d@   \        P                  P
                  '       d    \        4       R8X  d   \         P                  # V # )z"Maybe upcast [b]float16 to float32r±  )r\   rÔ   r‘  rv   r±  Úcodegen_upcast_to_fp32r¥  r“  r÷  s   &rT   Úupcast_compute_typer­  ±  sC   € ð 	”%—-‘-¤§¡Ð0Ô0Ü�M‰M×0×0Ð0ÜÓ! XÔ-ä�}‰}ÐØ€LrV   ÚKeyTypeÚValTypec                  óŠ   € ] tR tRtRtR R ltR R ltR R ltR	 R
 ltRR R llt	R R lt
R R ltR R ltR R ltRtR# )Ú
ScopedDictiÀ  zÔ
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               ó   € V ^8„  d   QhRR/# )rO   Úoriginal_dictúMapping[KeyType, ValType]rR   )rS   s   "rT   rU   ÚScopedDict.__annotate__È  s   € ÷ 4ñ 4Ð&?ñ 4rV   c                	ó    € Wn         / V n        R # rš   ©r³  Ú	new_items)rH  r³  s   &&rT   r  ÚScopedDict.__init__È  s   € Ø*ÔØ13ˆŽrV   c               ó    € V ^8„  d   QhRRRR/# )rO   r{  r®  rP   r¯  rR   )rS   s   "rT   rU   rµ  Ì  s   € ÷ 'ñ '˜wð '¨7ñ 'rV   c                	ón   € WP                   9   d   V P                   V,          # V P                  V,          # rš   ©r¸  r³  rX  s   &&rT   rÏ  ÚScopedDict.__getitem__Ì  s,   € Ø—.‘.Ô Ø—>‘> #Õ&Ð&Ø×!Ñ! #Õ&Ð&rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   r{  r®  r—   r¯  rP   r  rR   )rS   s   "rT   rU   rµ  Ñ  s!   € ÷ $ñ $˜wð $¨wð $¸4ñ $rV   c                	ó"   € W P                   V&   R # rš   )r¸  )rH  r{  r—   s   &&&rT   Ú__setitem__ÚScopedDict.__setitem__Ñ  s   € Ø#�‰�sÓrV   c               ó    € V ^8„  d   QhRRRR/# )rO   r{  rk  rP   rŠ   rR   )rS   s   "rT   rU   rµ  Ô  s   € ÷ Bñ B ð B¨4ñ BrV   c                	óJ   € WP                   9   ;'       g    WP                  9   # rš   r¼  rX  s   &&rT   Ú__contains__ÚScopedDict.__contains__Ô  s    € Ø—n‘nÑ$×AÐA¨×/AÑ/AÑ(AÐArV   Nc               ó$   € V ^8„  d   QhRRRRRR/# )rO   r{  r®  rL  úValType | NonerP   rR   )rS   s   "rT   rU   rµ  ×  s!   € ÷ 4ñ 4�wð 4¨ð 4À>ñ 4rV   c                	ó~   € WP                   9   d   V P                   V,          # V P                  P                  W4      # rš   )r¸  r³  rØ  )rH  r{  rL  s   &&&rT   rØ  ÚScopedDict.get×  s2   € Ø—.‘.Ô Ø—>‘> #Õ&Ð&Ø×!Ñ!×%Ñ% cÓ3Ð3rV   c               ó   € V ^8„  d   QhRR/# r�  rR   )rS   s   "rT   rU   rµ  Ü  s   € ÷ ñ ˜ñ rV   c                	óŒ   € \        V P                  4      pV P                   F  pW P                  9  g   K  V^,          pK   	  V# rŸ   )r^   r³  r¸  )rH  r   ro  s   &  rT   Ú__len__ÚScopedDict.__len__Ü  s<   € Ü�×"Ñ"Ó#ˆØ—”ˆAØ×*Ñ*Ö*Ø�Q•’ñ  ð ˆrV   c               ó   € V ^8„  d   QhRR/# )rO   rP   zIterator[KeyType]rR   )rS   s   "rT   rU   rµ  ã  s   € ÷ ñ Ð+ñ rV   c              #  	óŒ   "  € V P                    R j  x€L
  V P                   F  pWP                   9  g   K  Vx € K  	  R #  L/5irš   r·  )rH  ro  s   & rT   Ú__iter__ÚScopedDict.__iter__ã  s8   é € Ø×%Ñ%×%Ð%Ø—”ˆAØ×*Ñ*Ö*Ø”ó  ñ 	&ùs   ‚A’A“ A¸Ac               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rµ  é  s   € ÷ :ñ :˜$ñ :rV   c                	óT   € \        V P                  ;'       g    V P                  4      # rš   )rŠ   r³  r¸  r  s   &rT   r¢  ÚScopedDict.__bool__é  s   € Ü�D×&Ñ&×8Ð8¨$¯.©.Ó9Ð9rV   c               ó    € V ^8„  d   QhRRRR/# )rO   r{  r®  rP   r  rR   )rS   s   "rT   rU   rµ  ì  s   € ÷ "ñ "˜wð "¨4ñ "rV   c                	ó   € \         hrš   )rŽ  rX  s   &&rT   Ú__delitem__ÚScopedDict.__delitem__ì  s   € Ü!Ð!rV   r¼  rš   )r    r¡   r¢   r£   r¤   r  rÏ  rÀ  rÄ  rØ  rÌ  rÐ  r¢  r×  r¨   rR   rV   rT   r±  r±  À  s:   † ñõ4õ'õ
$õB÷4õ
õõ:÷"ñ "rV   r±  )Úfrozen_defaultr©   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   rœ   ztype[Any] | Noner©   rŠ   rP   r   rR   )rS   s   "rT   rU   rU   ñ  s"   € ÷ ñ Ð&ð ¸Tð ÈSñ rV   c              ó2   a€ R  V3R llpV f   V# V! V 4      # )c               ó    € V ^8„  d   QhRRRR/# )rO   rœ   rz   rP   rR   )rS   s   "rT   rU   Ú"ir_dataclass.<locals>.__annotate__ò  s   € ÷ Gñ G�"ð G˜ñ GrV   c                ó6   <€ \         P                  ! V R SR7      # )T)Úkw_onlyr©   )ÚdataclassesÚ	dataclass)rœ   r©   s   &€rT   ÚwrapÚir_dataclass.<locals>.wrapò  s   ø€ Ü×$Ò$ S°$¸vÔFÐFrV   rR   )rœ   r©   râ  s   "d rT   Úir_dataclassrä  ð  s%   ø€ ÷Gð Gð ‚{ØˆÙ�‹9ÐrV   c               ó   € V ^8„  d   QhRR/# )rO   rP   zlist[int] | NonerR   )rS   s   "rT   rU   rU   ú  s   € ÷ ñ Ð*ñ rV   c                 ó®   € \         P                  P                  P                  4       p V e)   V P                  '       d   V P                  P
                  # R # rš   )r\   r  r  r  Úfw_metadataÚbw_donated_idxs)r
  s    rT   Úget_donated_idxsré  ú  s@   € Ü—m‘m×2Ñ2×:Ñ:Ó<€OØÒ" ×'B×'BÐ'BØ×*Ñ*×:Ñ:Ð:ÙrV   c                  ó*   € ] tR tRt^ t^t^t^t^tRt	R# )ÚTritonAttrsDescriptorVersioni  rR   N)
r    r¡   r¢   r£   ÚV0_NO_TRITONÚV1_COMPILERÚV2_BACKENDSÚV3_BACKENDS_TUPLEÚV4_DICTr¨   rR   rV   rT   rë  rë    s    † Ø€LØ€KØ€Kà	ð ð „GrV   rë  c               ó   € V ^8„  d   QhRR/# )rO   rP   rë  rR   )rS   s   "rT   rU   rU     s   € ÷ 4ñ 4Ð-Iñ 4rV   c                 ól  € \         P                  P                  R 4      f   \        P                  # ^ RIp ^ RIp \        V P                  P                  R4      '       d   \        P                  # \        V P                  P                  R4      '       d   \        P                  # \        P                  # )r±  NÚAttrsDescriptor)rD  rE  rF  rë  rì  rÆ  rÇ  r;  rÝ  Úcompilerrî  rí  rð  )r±  s    rT   Ú#get_triton_attrs_descriptor_versionrõ    s�   € ä‡~�~×Ñ Ó)Ò1Ü+×8Ñ8Ð8ã#Û#äˆv�‰×'Ñ'Ð):×;Ò;ô ,×7Ñ7Ð7Ü	�—‘×)Ñ)Ð+<×	=Ò	=ä+×7Ñ7Ð7ô ,×3Ñ3Ð3rV   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU   %  s   € ÷ Yñ Y¨ñ YrV   c                 ó8   € \        4       \        P                  8H  # rš   )rõ  rë  rð  rR   rV   rT   Útriton_version_uses_attrs_dictrø  %  s   € Ü.Ó0Ô4P×4XÑ4XÑXÐXrV   c               ó    € V ^8„  d   QhRRRR/# )rO   rÉ  ztorch._ops.OperatorBaserP   r�  rR   )rS   s   "rT   rU   rU   )  s   € ÷ 5ñ 5Ð,ð 5°ñ 5rV   c                ó¢   € V P                  4       p\        V \        P                  P                  4      '       d   V R V P
                   2MTpW3# )r   )rì   rŒ   r\   râ  rã  Ú_overloadname)rÉ  Úop_overload_packet_nameÚop_overload_names   &  rT   Úget_op_namesrþ  )  sR   € Ø#%§7¡7£9Ðô �bœ%Ÿ*™*×/Ñ/×0Ò0ð #Ð
# 1 R×%5Ñ%5Ð$6Ñ7à$ð ð
 #Ð4Ð4rV   c               ó    € V ^8„  d   QhRRRR/# ©rO   r.  útorch.fx.NoderP   rŠ   rR   )rS   s   "rT   rU   rU   3  s   € ÷ !ñ !¸-ð !ÈDñ !rV   c                ó¢  € ^ RI Hp V P                  p\        V\        P
                  P                  4      '       g   R# V\        P                  P                  P                  P                  \        P                  P                  P                  P                  \        P                  P                  P                  P                  39   dƒ   V! W P                  V P                  RR7      pVe`   Vw  rEVR,          pV FL  pVf   K	  VP                  R,          P                   \        P"                  \        P$                  39   g   KK   R# 	  R# )zî
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).
)Únormalize_functionFT)Únormalize_to_only_use_kwargsra  r³  )Útorch.fx.operator_schemasr	  rá  rŒ   r\   râ  rã  r<  rã  Ú	index_putrL  Ú
index_put_Ú_unsafe_index_putr’   rë  rD  r»   rŠ   rŸ  )r.  r	  rá  Ú
normalizedrñ   rë  ra  rä  s   &       rT   Ú,_fx_node_is_input_dependent_cudagraph_unsafer
	  3  sô   € õ =à�^‰^€FÜ�fœeŸj™j×3Ñ3×4Ò4Ùð Ü�	‰	�‰× Ñ ×(Ñ(Ü�	‰	�‰×!Ñ!×)Ñ)Ü�	‰	�‰×(Ñ(×0Ñ0ðô ñ
 (Ø—L‘L '§.¡.Ètô
ˆ
ð Ò!Ø"‰IˆAØ˜YÕ'ˆGÛ�Ø”? s§x¡x°¥×'<Ñ'<Ü—J‘JÜ—K‘KðAö (ò  ñ ñ rV   c               ó    € V ^8„  d   QhRRRR/# r 	  rR   )rS   s   "rT   rU   rU   W  s   € ÷ "ñ "¨ð "¸4ñ "rV   c                ó.  € V P                   p\        V4      \        9   d   R# \        V\        P
                  P                  4      '       d6   \        P                  P                  P                  VP                  9   d   R# \        V 4      '       d   R# V P                  P                  R4      ;peb   \        V\        \        34      '       g   V.MTpV F:  p\        V\        P                   4      '       g   K%  VP"                  '       g   K9   R# 	  R# )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á  rQ   ÚFORBIDDEN_CUDAGRAPH_OPSrŒ   r\   râ  rã  r9  ræ  Úcudagraph_unsaferè  r
	  rD  rØ  r©  rq  rù  Ú	is_sparse)r.  rá  r³  Úvalsrˆ   s   &    rT   ra  ra  W  sÅ   € ð �^‰^€Fô ˆ6ƒ{Ô-Ô-Ùô 	�6œ5Ÿ:™:×0Ñ0×1Ò1Ü�H‰H�L‰L×)Ñ)¨V¯[©[Ô8áô 4°G×<Ò<Ùð �|‰|×Ñ Ó&Ð&ˆÒ3Ü& s¬T´5¨M×:Ò:�‰uÀˆÛˆAÜ˜!œUŸ\™\×*Ô*¨q¯{¯{ª{Úñ ñ rV   c               ó    € V ^8„  d   QhRRRR/# )rO   r¶  rD   rP   rŠ   rR   )rS   s   "rT   rU   rU   |  s   € ÷ ñ  ð ¨tñ rV   c                óþ   € ^RI Hp \        WP                  VP                  34      '       d   R# \        WP
                  VP                  34      '       g   R# \        V RR4      pVe   \        V4      '       d   R# R# )aC  
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 (Switch, WhileLoop)
- Ops with sparse tensor outputs
ru  TFr.  N)	rw  rv  rŒ   ÚSwitchÚ	WhileLooprJ  rA   r[   ra  )r¶  rv  r.  s   &  rT   Úis_cudagraph_unsafe_opr	  |  sf   € õ ô �$Ÿ™ B§L¡LÐ1×2Ò2Ùä�d×.Ñ.°·±Ð@×AÒAÙä�d˜I tÓ,€GØÒÔ:¸7×CÒCÙárV   c               ó   € V ^8„  d   QhRR/# rN   rR   )rS   s   "rT   rU   rU   –  s   € ÷ 
ñ 
˜Sñ 
rV   c                 ó<  € \         P                  P                  R R4      p \        P                  ! 4       '       da   ^ RIHp V! 4       pV'       dL   \         P                  P                  VRR4      pV '       d!   \         P                  P                  W0.4      MTp V # )ÚLD_LIBRARY_PATHrw  )Úget_runtime_pathr…  Úlib)
r¥  r¦  rØ  rv   r_  Úlibfb.py.parutilr	  r¶  r�  Úpathsep)r¶  r	  Úruntime_pathÚlib_paths       rT   Úget_ld_library_pathr	  –  sh   € Ü�:‰:�>‰>Ð+¨RÓ0€DÜ×Ò×ÒÝ5á'Ó)ˆßÜ—w‘w—|‘| L°)¸UÓCˆHß8<”2—:‘:—?‘? HÐ#3Ô4À(ˆDà€KrV   c               ó    € V ^8„  d   QhRRRR/# )rO   r_  r<   rP   rŠ   rR   )rS   s   "rT   rU   rU   £  s   € ÷ ñ Ð1Eð È$ñ rV   c                óP   € ^ RI Hp \        W4      ;'       d    V P                  RJ# )r   )ÚSubgraphPythonWrapperCodegenN)Útorch._inductor.codegen.wrapperr"	  rŒ   Úpartition_signatures)r_  r"	  s   & rT   Ú#is_codegen_graph_partition_subgraphr%	  £  s,   € ÝLô 	�7Ó9÷ 	5ð 	5Ø×(Ñ(°Ð4ðrV   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU   ¬  s   € ÷ 1ñ 1 dñ 1rV   c                 óê   € \         P                  P                  P                  P                  ;'       g    \
        P                  R J;'       d%    \         P                  P                  P                  # rš   )r\   r„  rv   r±  Ú
cudagraphsÚ&_unstable_customized_partition_wrapperr_  Úgraph_partitionrR   rV   rT   Úis_using_cudagraph_partitionr+	  ¬  sY   € ä�‰×Ñ×%Ñ%×0Ñ0÷ 	Fð 	FÜ1×9Ñ9ÀÐE÷1ð 1ô �/‰/×
 Ñ
 ×
0Ñ
0ð1rV   c               ó    € V ^8„  d   QhRRRR/# )rO   rf  rƒ   rP   rõ  rR   )rS   s   "rT   rU   rU   ³  s   € ÷ ñ ˜#ð  +ñ rV   c                óþ   € ^RI Hp VP                  P                  P	                  V R4      '       d=   VP                  P                  P                  V R4      '       d   \        P                  # \        P                  # )rX   r¼  l        i   €)	r¿  r½  rÀ  rÁ  Ústatically_known_ltr  r\   r›  r�  )rf  r½  s   & rT   Údtype_from_sizer/	  ³  sX   € Ýà‡w�w×Ñ×+Ñ+Øˆe÷ò à
�'‰'×
Ñ
×
/Ñ
/°°h×
?Ò
?Ü�{‰{Ðä�{‰{ÐrV   c               ó    € V ^8„  d   QhRRRR/# ©rO   ré   rQ   rP   rŠ   rR   )rS   s   "rT   rU   rU   Á  ó   € ÷ 	ñ 	¨#ð 	°$ñ 	rV   c                óv   € V R8X  d)   \         P                  P                  P                  4       # RV 9   d   R# R# )z3
Returns True if the device supports MKL-DNN BF16.
rC  rK   TF)r\   r<  rà  Ú_is_mkldnn_bf16_supported©ré   s   &rT   Úis_mkldnn_bf16_supportedr6	  Á  ó3   € ð �eÔÜ�y‰y×Ñ×9Ñ9Ó;Ð;Ø	�+Ô	áÙrV   c               ó    € V ^8„  d   QhRRRR/# r1	  rR   )rS   s   "rT   rU   rU   Í  r2	  rV   c                óv   € V R8X  d)   \         P                  P                  P                  4       # RV 9   d   R# R# )z3
Returns True if the device supports MKL-DNN FP16.
rC  rK   TF)r\   r<  rà  Ú_is_mkldnn_fp16_supportedr5	  s   &rT   Úis_mkldnn_fp16_supportedr;	  Í  r7	  rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   ÚelementszSequence[Sequence[T]]ÚheaderszSequence[T]rP   rQ   rR   )rS   s   "rT   rU   rU   Ù  s"   € ÷ ñ Ð/ð ¸+ð È#ñ rV   c           
     óê  € V Uu. uF  p\        \        V4      4      NK  	  ppV  Fy  p\        V4      \        V4      8w  d$   \        R \        V4       R\        V4       24      h\        V4       F+  w  rR\	        W5,          \        \        V4      4      4      W5&   K-  	  K{  	  . pVP                  RP                  R \        W4       4       4      4       \        V4      \        V4      ^,          ,           \        V4      ^,
          ,           pVP                  RV,          4       V  F3  pVP                  RP                  R \        WC4       4       4      4       K5  	  RP                  V4      # u upi )zRow length z != header length Ú|c              3  ó4   "  € T F  w  rR W  R 2x € K  	  R# 5i©r  NrR   )rÉ   rš  rî  s   &  rT   rË   Útabulate_2d.<locals>.<genexpr>ã  s   é € ÐHÑ3G©4¨1˜A˜a ˜W A�,Ó3Gùó   ‚r  c              3  ó4   "  € T F  w  rR W  R 2x € K  	  R# 5irB	  rR   )rÉ   rø   rî  s   &  rT   rË   rC	  è  s   é € ÐHÑ7G©t¨q  ! C ¨�lÓ7GùrD	  rÓ  )	r^   rQ   r_   rk  rÚ   ró  r�  rá   r^  )r=	  r>	  rø   ÚwidthsÚrowrö   r  Útotal_widths   &&      rT   Útabulate_2drI	  Ù  s!  € Ù#*Ó+¡7˜aŒc”#�a“&Žk¡7€FÐ+ÛˆÜˆs‹8”s˜7“|Ô#Ü Øœc #›h˜ZÐ'9¼#¸g»,¸ÐHóð ô ˜c–N‰DˆAÜ˜F�I¤s¬3¨q«6£{Ó3ˆF‹Ió #ñ ð €EØ	‡L�L�—‘ÑH´3°wÔ3GÓHÓHÔIä�f“+¤ V£¨q¥Õ1´S¸³[À1µ_ÕE€KØ	‡L�L��{Õ"Ô#ÛˆØ�‰�S—X‘XÑH´s¸3Ô7GÓHÓHÖIñ à�9‰9�UÓÐùò ,s   …E0c          
     ó,   € V ^8„  d   QhRRRRRRRRRR/# )	rO   Údict1r´  Údict2Ú
d1_defaultrÇ  Ú
d2_defaultrP   zEGenerator[tuple[KeyType, ValType | None, ValType | None], None, None]rR   )rS   s   "rT   rU   rU   ì  s;   € ÷  
ñ  
Ø$ð 
à$ð 
ð ð 
ð ð	 
ð
 Kñ 
rV   c              #  óø   "  € \        V P                  4       4      \        VP                  4       4      ,          pV F8  pV P                  V4      pVP                  V4      pTVe   TMTVe   TMT3x € K:  	  R# 5i)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"   rÙ  rØ  )rK	  rL	  rM	  rN	  Úall_keysr{  Úvalue1Úvalue2s   &&&&    rT   Ú	zip_dictsrS	  ì  sp   é € ô( ˜%Ÿ*™*›,Ó'¬*°U·Z±Z³\Ó*BÕB€Hó ˆà—‘˜3“ˆØ—‘˜3“ˆð ØÒ(‰F¨jØÒ(‰F¨jð
ô 	
ó ùs   ‚A8A:c               ó    € V ^8„  d   QhRRRR/# )rO   Úconfig_patchesrì  rP   rR   )rS   s   "rT   rU   rU     s   € ÷ Rñ R°ð RÀNñ RrV   c                ó`  € R R lpR R lpV P                  R\        P                  P                  4      pV P	                  4       p V'       dr   V! V RR4       V! V RR4       V! V R	\
        P                  P                  '       * 4       V! V R
R4       V! V R\        P                  P                  4       V! V RR4       V P                  R\        P                  P                  4      pV P                  R\        P                  P                  4      pVR8X  d   V'       d   \        R4      hV # )a
  
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               ó(   € V ^8„  d   QhRRRRRRRR/# ©	rO   rU	  rì  Úconfig_namerQ   Úconfig_valuer   rP   r  rR   )rS   s   "rT   rU   Ú2maybe_aoti_standalone_config.<locals>.__annotate__   s,   € ÷ 	ñ 	Ø&ð	Ø58ð	ØHKð	à	ñ	rV   c                óž   € V P                  V\        \        V4      4      pVf   W V&   R # V'       g   W28w  d   \        RV RV R24      hR # R # )NzInvalid config: Ú=z3 when aot_inductor_mode.compile_standalone is True.)rØ  r[   rv   r  ©rU	  rY	  rZ	  r—   s   &&& rT   Úpatch_configÚ2maybe_aoti_standalone_config.<locals>.patch_config   sY   € ð ×"Ñ" ;´¼ÀÓ0LÓMˆØŠ=Ø*6˜;Ó'ß˜5Ô0ÜØ" ; -¨q°°Ð>qÐróð ñ 1‘rV   c               ó(   € V ^8„  d   QhRRRRRRRR/# rX	  rR   )rS   s   "rT   rU   r[	  +  s,   € ÷ 
3ñ 
3Ø&ð
3Ø58ð
3ØHKð
3à	ñ
3rV   c                óˆ   € V P                  V\        \        V4      4      pW28w  d   \        P	                  R VV4       W V&   R# )zDOverriding: %s=%s when aot_inductor_mode.compile_standalone is True.N)rØ  r[   rv   rä   rµ  r^	  s   &&& rT   Úforce_patch_configÚ8maybe_aoti_standalone_config.<locals>.force_patch_config+  sB   € ð ×"Ñ" ;´¼ÀÓ0LÓMˆØÔ Ü�K‰KØVØØôð
 '3�{Ó#rV   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Úwindowszªconfig.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Ø  rv   Úaot_inductor_modeÚcompile_standaloneÚcopyr\   rÝ  r~   Útest_configsÚuse_libtorchÚaot_inductorÚcross_target_platformÚpackage_constants_in_sor  )rU	  r_	  rc	  rh	  rm	  rn	  s   &     rT   Úmaybe_aoti_standalone_configro	    s+  € õ"	õ
3ð (×+Ñ+Ø.Ü× Ñ ×3Ñ3óÐð
 $×(Ñ(Ó*€Nßá�^Ð%DÀdÔKá�^Ð%GÈÔNáØÐAÄuÇ}Á}×GXÑGXÔCXô	
ñ 	ØÐIÈ<ô	
ñ 	ØØ(Ü×Ñ×,Ñ,ô	
ñ
 	˜>Ð+IÈ5ÔQà*×.Ñ.Ø,Ü×Ñ×1Ñ1óÐð
 -×0Ñ0Ø.Ü×Ñ×3Ñ3óÐð
  	Ô)×.EÜð]ó
ð 	
ð
 ÐrV   c               ó    € V ^8„  d   QhRRRR/# )rO   Úconsts_sizerƒ   rP   ztuple[bool, bool]rR   )rS   s   "rT   rU   rU   d  s   € ÷ &2ñ &2¨3ð &2Ð3Dñ &2rV   c                óÔ  € \         P                  P                  '       d+   \         P                  P                  R8X  d   \	        R4      h\         P                  P                  '       d2   \         P                  P
                  R8X  d   \	        R4      hRpRpW!3# \         P                  P                  R8X  d   RpRpW!3# V R8:  d   R# Rp\         P                  ! 4       '       * pW!3# )	zý
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_blobz‚config.aot_inductor.package_constants_on_disk_format = binary_blob and config.aot_inductor.force_mmap_weights cannot both be True.rf	  zKwhen cross_target_platform is windows, use_mmap_weights should not be true.TFi ”5w)FF)rv   rl	  Úforce_mmap_weightsÚ package_constants_on_disk_formatr  rm	  r_  )rq	  Úuse_mmap_weightsÚuse_external_weightss   &  rT   Údetermine_aoti_mmap_flagsrx	  d  sä   € ô 	×Ñ×.×.Ð.Ü×Ñ×@Ñ@ÀMÔQäðJó
ð 	
ô
 ×Ñ×-×-Ð-Ü×Ñ×4Ñ4¸	ÔAÜØ]óð ð  ÐØ$ÐØ#Ð5Ð5ä×Ñ×;Ñ;¸}ÔLØ#ÐØ ÐØ#Ð5Ð5à�mÔ#ØÐà ÐÜ!×+Ò+Ó-Ô-ÐàÐ1Ð1rV   c               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   rU   �  s   € ÷ ñ  $ñ rV   c                 óî   € ^ RI Hp  V P                  P                  pVf   R# \	        V\
        4      '       g   \        R4      hVR8X  d   R# \        P                  ! RV4      '       g   \        R4      hR# )zD
Validates if a model name is suitable for use in code generation.

ru   Tz4Invalid AOTI model name: Model name must be a stringrw  z^[a-zA-Z_][a-zA-Z0-9_]*$zVInvalid AOTI model name: Model name can only contain letters, numbers, and underscores)	r�  rv   rl	  Úmodel_name_for_generated_filesrŒ   rQ   rK  rê   rë   )rv   Ú
model_names     rT   Úis_valid_aoti_model_namer}	  �  sn   € õ
 'à×$Ñ$×CÑC€JàÒÙä�j¤#×&Ò&ÜÐOÓPÐPà�RÔÙô �8Š8Ð/°×<Ò<ÜØdó
ð 	
ñ rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   ra   r-   Úunbacked_onlyrŠ   rP   zOrderedSet[sympy.Symbol]rR   )rS   s   "rT   rU   rU   ¨  s"   € ÷ ñ ˜ð °Tð Ð>Vñ rV   c                ó>   € V'       d   \        V 4      # \        V 4      # rš   )r+   r*   )ra   r	  s   &&rT   Úget_free_symbolsr�	  ¨  s   € ßÜ$ QÓ'Ð'ä˜A‹ÐrV   c               ó   € V ^8„  d   QhRR/# )rO   rP   zdict[str, str]rR   )rS   s   "rT   rU   rU   ¯  s   € ÷ ñ ˜~ñ rV   c            	     ó  € \         P                  P                  \         P                  P                  \         P                  P                  \        P
                  4      4      4      p / \         P                  CR\         P                  P                  R\         P                  P                  V .\        P                  O54      4      /Cp\        P                  ! 4       '       d   \        P                  ! R4      VR&   V# )z9
Get a base environment for running Python subprocesses.
Ú
PYTHONPATHÚTORCH_CUSTOM_PYTHONPATHr²  Ú
PYTHONHOME)r¥  r¶  rÐ  Úabspathr\   rÑ  r¦  rØ  r	  r�  r€  rv   r_  Ú	sysconfigÚget_path)Útorch_package_rootÚenvs     rT   Úpython_subprocess_envrŒ	  ¯  s¯   € ô
 Ÿ™Ÿ™Ü
�‰�‰œŸ™Ÿ™¬¯©Ó7Ó8óÐðä
�*‰*ðð 	”b—j‘j—n‘nØ%Ü�J‰J�O‰OÐ/Ð;´#·(±(Ñ;Ó<ó
ñ	€Cô" ×Ò×ÒÜ%×.Ò.¨vÓ6ˆˆLÑà€JrV   c               ó    € V ^8„  d   QhRRRR/# )rO   Ú	extra_envúMapping[str, str | None] | NonerP   r  rR   )rS   s   "rT   rU   rU   Î  s   € ÷ $ñ $Ð$Cð $Èñ $rV   c                ó¶   € V f   R# V P                  4        F>  w  rVf#   \        P                  P                  VR4       K+  V\        P                  V&   K@  	  R# )z×
Apply environment updates sent from a parent process to a persistent worker.
A None value means the variable is absent in the parent and should be
removed from the worker, rather than leaving a stale value behind.
N)rÚ  r¥  r¦  r`   )rŽ	  r{  r—   s   &  rT   Úapply_subprocess_envr‘	  Î  sD   € ð ÒÙà—o‘oÖ'‰
ˆØŠ=Ü�J‰J�N‰N˜3 Ö%à#ŒB�J‰J�s‹Oó	 (rV   c               ó    € V ^8„  d   QhRRRR/# )rO   rŽ	  r�	  rP   r£  rR   )rS   s   "rT   rU   rU   ß  s   € ÷ #ñ #Ø.ð#àñ#rV   c              #  ór  "  € V f   Rx € R# \        \        P                  4      p\        V 4        Rx € \        P                  P	                  4        \        P                  P                  V4       R#   \        P                  P	                  4        \        P                  P                  T4       i ; i5i)zN
Temporarily apply parent process environment updates in a persistent worker.
N)r¾  r¥  r¦  r‘	  rG  r¨  )rŽ	  Úold_envs   & rT   Úpatch_subprocess_envr•	  Þ  sx   é € ð ÒÛÙä”2—:‘:Ó€GÜ˜Ô#ð#Ûä
�
‰
×ÑÔÜ
�
‰
×Ñ˜'Ö"øô 	�
‰
×ÑÔÜ
�
‰
×Ñ˜'Õ"üs   ‚/B7²A5 ¶?B7Á5?B4Â4B7c                  ó0   € ] tR tRt$ RtR]R&   R]R&   RtR# )ÚCUDAGraphWrapperMetadataiò  zˆ
Metadata for Customized CUDAGraphWrapper.

Currently assumes there is 1 dynamo graph and will extend to
multiple graphs in the future.
rƒ   Únum_partitionsÚpartition_indexrR   Nr±   rR   rV   rT   r—	  r—	  ò  s   ‡ ñð Óð ×rV   r—	  .c                  ó&   € ] tR tRt$ RtR]R&   RtR# )ÚCUDAGraphWrapperi	  NzCUDAGraphWrapperType | Noner_  rR   )r    r¡   r¢   r£   r_  r²   r¨   rR   rV   rT   r›	  r›	  	  s   ‡ Ø+/€GÐ(×/rV   r›	  c               ó    € V ^8„  d   QhRRRR/# )rO   r_  ÚCUDAGraphWrapperTyperP   r  rR   )rS   s   "rT   rU   rU     s   € ÷ =ñ =Ð/Cð =Èñ =rV   c                ó   € V \         n        R # rš   )r)	  r_  )r_  s   &rT   Ú!set_customized_partition_wrappersrŸ	    s   € Ø5<Ô*Ö2rV   c               ó    € V ^8„  d   QhRRRR/# )rO   rT  rH   rP   z tuple[list[Any], dict[str, Any]]rR   )rS   s   "rT   rU   rU     s   € ÷ +ñ +Ð.ð +Ð3Sñ +rV   c                ó  a€ \        R V P                  4      p\        V\        P                  P
                  P                  4      '       d*   VP                  VP                  VP                  4      w  r#M&. VP                  OVP                  OpVP                  pVP                  W#4      p\        V4      p\        VRR4      p\        V\        P                  P                  4      '       dj   VP                   P"                   Uu. uF"  qUP$                  '       d   K  VP&                  NK$  	  ppVR\)        V4        F  pVP+                  VR4       K  	  \,        P.                  ! W#34      w  r‰R R lp
V Uu. uF>  pV
! V4      '       d,   \        P                  P
                  P1                  VRR7      MTNK@  	  ppR R loR	 V3R
 llpV Uu. uF
  q¼! V4      NK  	  pp\,        P2                  ! W‰4      w  r#W#3# u upi u upi u upi )rA   rI  Nc               ó   € V ^8„  d   QhRR/# r3  rR   )rS   s   "rT   rU   Ú'snode_args_kwargs.<locals>.__annotate__2  s   € ÷ 
ñ 
˜Dñ 
rV   c                ó  € \        V \        P                  P                  P                  4      ;'       dX    \        V \        P                  P                  P
                  \        P                  P                  P                  34      '       * # rš   )rŒ   r\   r„  rv  rB   ÚGeneratorStateÚOpaqueObjectStater  s   &rT   Ú_is_tensor_irÚ(snode_args_kwargs.<locals>._is_tensor_ir2  sc   € Ü˜!œUŸ_™_×/Ñ/×6Ñ6Ó7÷ 
ð 
Ä
ØÜ�_‰_×Ñ×.Ñ.´·±×0BÑ0B×0TÑ0TÐUóA
ô =
ð 	
rV   T)Úreplace_symbols_with_hintsc               ó   € V ^8„  d   QhRR/# )rO   rP   rð  rR   )rS   s   "rT   rU   r£	  ?  s   € ÷ =ñ =¨ñ =rV   c                ó2   € \         P                  ! WVR 7      # )rº   )r\   rÓ   )rf  r»   r¼   s   &&&rT   Ú_tensorÚ"snode_args_kwargs.<locals>._tensor?  s   € Ü�{Š{˜4°VÔ<Ð<rV   c               ó    € V ^8„  d   QhRRRR/# )rO   rø   r   rP   rR   )rS   s   "rT   rU   r£	  B  s   € ÷ ñ ˜#ð  #ñ rV   c                ó¤   <€ \        V \        P                  4      '       g   V # S! V P                  4       V P                  V P
                  4      pV# rš   )rŒ   r\   rù  rf  r»   r¼   )rø   rœ  r¬	  s   & €rT   Úto_real_tensorÚ)snode_args_kwargs.<locals>.to_real_tensorB  s:   ø€ Ü˜!œUŸ\™\×*Ò*ØˆHÙ�a—f‘f“h §¡¨¯©Ó2ˆØˆ
rV   )r   r¶  rŒ   r\   r„  rv  rJ  Úunflatten_argsr½  Úconstant_argsrë  Úfill_non_provided_argsr¾  r[   râ  rã  rÑ  rÒ  Ú
kwarg_onlyrì   r^   r`   Úpytreer#   Úir_node_to_tensorÚtree_unflatten)rT  r¶  r’   rë  rI  rï  Úpositional_arg_namesÚarg_nameÚ	flat_argsÚflat_args_pytree_specr§	  rZ  r°	  r¬	  s   &            @rT   Úsnode_args_kwargsr½	    sÊ  ø€ Ü� §
¡
Ó+€DÜ�$œŸ™×*Ñ*×9Ñ9×:Ò:Ø×*Ñ*¨4¯;©;¸×8JÑ8JÓK‰ˆˆfà2�—‘Ð2˜t×1Ñ1Ð2ˆØ—‘ˆà×&Ñ& tÓ4€DÜ�&‹\€FÜ˜$ ¨tÓ4€KÜ�+œuŸz™z×4Ñ4×5Ò5à +× 3Ñ 3× =Ò =ó 
Ù =˜Ç^Å^ŒHˆC�HŒHÑ =ð 	ð  
ð -¨[¬s°4«yÓ9ˆHØ�J‰J�x Ö&ñ :ä'-×':Ò':¸D¸>Ó'JÑ$€Iõ
ñ ó	ñ ˆAñ ˜×Òô 	�‰×Ñ×,Ñ,¨QÈ4Ð,ÔPàò	ñ ð	 ð õ=÷ð ñ -6Ó6©I q� Ö"©I€IÐ6Ü×(Ò(¨ÓJ�L€DØˆ<Ðùò? 
ùòùò  7s   ÄH ÄH Å:HÆ1HÇH
c               ó    € V ^8„  d   QhRRRR/# )rO   rw  r=   rP   rŠ   rR   )rS   s   "rT   rU   rU   M  s   € ÷ 
ñ 
 ð 
¨ñ 
rV   c                óØ   € ^RI Hp V P                  pVP                  P                  '       d-   VP	                  VP                  P                  R,           4      pVP                  R4      # )rX   r¼  rñ   )rš  rï  Úfwd_rng_stater’  r�  )r¿  r½  rì   rÀ  rÎ  r‚  )rw  r½  Údep_names   &  rT   Úis_nonfreeable_buffersrÂ	  M  sP   € Ýà�x‰x€Hð 	‡w�w‡|‡|€|Ø×(Ñ(¨¯©¯©¸Õ);Ó<ˆØ×ÑØIóð rV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   rì   rQ   Útemplate_dirr2   rP   rR   )rS   s   "rT   rU   rU   [  s&   € ÷ 
Lñ 
L˜ð 
L¨4ð 
L°Cñ 
LrV   c                óì   € W R2,          p \        VRR7      ;_uu_ 4       pVP                  4       uuRRR4       #   + '       g   i     R# ; i  \         d   p\        RT RT 24      ThRp?ii ; i)z,Load a template file and return its content.z	.py.jinjar‹  )r'  Nz	Template z is not valid UTF-8: )ÚopenÚreadÚUnicodeDecodeErrorrK  )rì   rÄ	  r¶  rÇ  rø   s   &&   rT   Úload_templaterÉ	  [  sl   € à˜F )Ð,Õ,€DðLÜ�$ ×)Õ)¨QØ—6‘6“8÷ *×)×)Ó)ûäô Lô ˜9 T FÐ*?À¸sÐCÓDÈ!ÐKûð	Lús1   �A ¢=²
A ½A	ÁA ÁA ÁA3ÁA.Á.A3c               ó    € V ^8„  d   QhRRRR/# )rO   r¶  r	  rP   rŠ   rR   )rS   s   "rT   rU   rU   h  s   € ÷ #-ñ #- ]ð #-°tñ #-rV   c                ó²  € V P                   p\        V\        P                  P                  \        P                  P
                  34      '       g   \        R\        V4       24      h\        P                  '       g   R# \        \        P                  P                  P                  P                  \        P                  P                  P                  P                  .4      pW9   d   R# \        \        P                  P                   P"                  .4      p\        V\        P                  P
                  4      '       d   W9   # \%        V 4      '       * # )zLDecide whether fallback for a node. This is only used in inductor lite mode.z6Expected OpOverload or HigherOrderOperator, but found F)rá  rŒ   r\   râ  rã  r‚  r_   rE  rv   Úfallback_by_defaultr"   r<  rã  Ú_assert_scalarrL  Úlift_fresh_copyÚhigher_orderÚ triton_kernel_wrapper_functionalr   )r¶  rá  Ú"skip_fallback_due_to_dynamic_shapeÚfallback_hopss   &   rT   Úshould_fallback_by_defaultrÓ	  h  sú   € à�[‰[€Fä�fœuŸz™z×4Ñ4´e·j±j×6TÑ6TÐU×VÒVÜØDÄTÈ&Ã\ÀNÐSó
ð 	
ô ×%×%Ð%Ùô *4ä�I‰I�N‰N×)Ñ)×1Ñ1Ü�I‰I�N‰N×*Ñ*×2Ñ2ð	
ó*Ð&ð Ô3Ùô Ü	�‰×	Ñ	×	@Ñ	@ÐAó€Mô �&œ%Ÿ*™*×8Ñ8×9Ò9ØÑ&Ð&ä& tÓ,Ô,Ð,rV   c               ó    € V ^8„  d   QhRRRR/# rˆ  rR   )rS   s   "rT   rU   rU   ¡  s   € ÷ %ñ %˜cð % dñ %rV   c                ó   € V \         9   # )z0Check if an operation is a collective operation.)ÚCOLLECTIVE_OPS)r‰  s   &rT   Úis_collective_opr×	  ¡  s   € à”nÑ$Ð$rV   c               ó   € V ^8„  d   QhRR/# )rO   rP   r¯   rR   )rS   s   "rT   rU   rU   §  s   € ÷ ñ ˜yñ rV   c                 ó:   €  ^ RI Hp  V #   \         d    . u # i ; i)r   ©Útlx_only_cuda_options)Ú+triton.language.extra.tlx.inductor.registryrÛ	  r=  rÚ	  s    rT   rÛ	  rÛ	  ¦  s$   € ðåUà$Ð$øäô ØŠ	ðús   ‚
 Š™c               ó$   € V ^8„  d   QhRRRRRR/# )rO   ra   rƒ   r8  rP   rR   )rS   s   "rT   rU   rU   ²  s!   € ÷ "ñ "�ð "˜ð " ñ "rV   c                ó<   € W,           ^,
          V,          V,          # )z(Round x up to the nearest multiple of y.rR   )ra   r8  s   &&rT   Ú	_round_uprß	  ²  s   € à�U�Q�Y˜1Õ Õ!Ð!rV   c               ó4   € V ^8„  d   QhRRRRRRRRR	RR
RRR/# )rO   Úmat_sizeztuple[Any, Any]Ú
scale_sizeztuple[Any, ...]Úscale_numelr   Ú	mat_dtyperõ  Úscale_dtypeÚeq_fnzCallable[[Any, Any], bool]rP   útuple[Any | None, Any | None]rR   )rS   s   "rT   rU   rU   ·  sW   € ÷ Yñ YØðYàðYð ðYð ð	Yð
 ðYð &ðYð #ñYrV   c                óÐ	  € ^ RI HpHp V! V^4      '       d   VP                  VP                  3# \        V4      ^8¼  Ed…   V! V^ ,          V ^ ,          4      '       d   V! V^,          ^4      '       g4   V! V^ ,          ^4      '       d6   V! V^,          V ^,          4      '       d   VP                  VP                  3# V! V^ ,          V ^ ,          4      '       d(   V! V^,          \        V ^,          ^€4      4      '       gE   V! V^,          V ^,          4      '       d@   V! V^ ,          \        V ^ ,          ^€4      4      '       d   VP                  VP                  3# V! V^ ,          \        V ^ ,          ^€4      4      '       d@   V! V^,          \        V ^,          ^€4      4      '       d   VP                  VP                  3# V\        P                  8X  d   ^M^pV\        P                  8X  Edi   V\        P                  8X  EdS   \        P                  P                  4       '       d…   \        V ^ ,          ^4      V,          V ^,          ,          p	\        W€^,          ,          ^4      V ^ ,          ,          p
V! W)4      '       g   V! W*4      '       d   VP                  VP                  3# Mª\!        V ^ ,          ^€4      \!        \        W€^,          ,          ^4      ^4      ,          p	\!        V ^,          ^€4      \!        \        W€^ ,          ,          ^4      ^4      ,          p
V! W)4      '       g   V! W*4      '       d   VP                  VP"                  3# V\        P$                  8X  Edu   \        P&                  P(                  '       gÒ   \        P                  P                  4       '       g®   \!        V ^ ,          ^€4      \!        \        W€^,          ,          ^ 4      ^4      ,          p	\!        V ^,          ^€4      \!        \        W€^ ,          ,          ^ 4      ^4      ,          p
V! W)4      '       g   V! W*4      '       d   VP*                  VP"                  3#  R# \        V ^ ,          ^ 4      V,          V ^,          ,          p	\        W€^,          ,          ^ 4      V ^ ,          ,          p
V! W)4      '       g   V! W*4      '       d   VP*                  VP                  3# R# )z2
Core implementation for scale/swizzle inference.
)r9   ÚSwizzleType©NN)Útorch.nn.functionalr9   ré	  Ú
TensorWiseÚ
NO_SWIZZLEr^   ÚRowWiserw   ÚBlockWise1x128ÚBlockWise128x128r\   r�  r‡  rK   Ú_is_compiledÚBlockWise1x16rß	  ÚSWIZZLE_32_4_4r�  rÝ  r~   ÚBlockWise1x32)rá	  râ	  rã	  rä	  rå	  ræ	  r9   ré	  ÚK_multiplierÚexpected_numel_aÚexpected_numel_bs   &&&&&&     rT   Ú_infer_scale_swizzle_implrø	  ·  s  € ÷ =ñ ˆ[˜!×ÒØ×%Ñ% {×'=Ñ'=Ð=Ð=ô ˆ:ƒ˜!ÕÙ�*˜Q•- ¨!¥×-Ò-±%¸
À1½Àq×2IÒ2IÙ�*˜Q•- ×#Ò#©¨j¸­m¸XÀa½[×(IÒ(Ià×&Ñ&¨×(>Ñ(>Ð>Ð>ñ �*˜Q•- ¨!¥×-Ò-Ù�j •m¤W¨X°a­[¸#Ó%>×?Ò?á�*˜Q•- ¨!¥×-Ò-Ù�j •m¤W¨X°a­[¸#Ó%>×?Ò?à×-Ñ-¨{×/EÑ/EÐEÐEñ �˜A•¤¨°­°SÓ 9×:Ò:¹uØ�q�Mœ7 8¨A¥;°Ó4÷@
ò @
ð ×/Ñ/°×1GÑ1GÐGÐGð "¤U×%;Ñ%;Ô;‘1À€Lð ”E×*Ñ*Õ*¨{¼e×>QÑ>QÕ/QÜ�9‰9×!Ñ!×#Ò#ä& x°¥{°BÓ7¸,ÕFÈÐRSÍÕTÐÜ& |¸qµkÕ'AÀ2ÓFÈÐRSÍÕTÐÙ�[×3Ò3±uØ÷8ò 8ð #×0Ñ0°+×2HÑ2HÐHÐHð8ô  )¨°!­°cÓ:¼YÜ˜°¥{Õ2°BÓ7¸ó>õ  Ðô  )¨°!­°cÓ:¼YÜ˜°¥{Õ2°BÓ7¸ó>õ  Ðñ �[×3Ò3±uØ÷8ò 8ð #×0Ñ0°+×2LÑ2LÐLÐLð ”e×*Ñ*Õ*Ü�}‰}× × Ð ¬¯©×)?Ñ)?×)AÒ)Aä(¨°!­°cÓ:¼YÜ˜°¥{Õ2°BÓ7¸ó>õ  Ðô  )¨°!­°cÓ:¼YÜ˜°¥{Õ2°BÓ7¸ó>õ  Ðñ �[×3Ò3±uØ÷8ò 8ð #×0Ñ0°+×2LÑ2LÐLÐLð8ð Ðô  ' x°¥{°BÓ7¸,ÕFÈÐRSÍÕTÐÜ& |¸qµkÕ'AÀ2ÓFÈÐRSÍÕTÐÙ�[×3Ò3±uØ÷8ò 8ð #×0Ñ0°+×2HÑ2HÐHÐHàÐrV   c               ó$   € V ^8„  d   QhRRRRRR/# )rO   Úmatrð  ÚscalerP   rç	  rR   )rS   s   "rT   rU   rU     s$   € ÷ ñ Ø	ðØ*ðà"ñrV   c           	     óØ   € \        V P                  ^ ,          V P                  ^,          3\        VP                  4      VP                  4       V P                  VP                  R R7      # )aR  
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
c                ó
   € W8H  # rš   rR   r&  s   &&rT   r×  Ú%infer_scale_swizzle.<locals>.<lambda>-  s   € ˜1š6rV   ©rá	  râ	  rã	  rä	  rå	  ræ	  )rø	  rº  rq  Únumelr»   )rú	  rû	  s   &&rT   Úinfer_scale_swizzler
    sO   € ô( %Ø—)‘)˜A•, §	¡	¨!¥Ð-Ü˜Ÿ™Ó%Ø—K‘K“MØ—)‘)Ø—K‘KÙ!ôð rV   c               ó(   € V ^8„  d   QhRRRRRRRR/# )rO   rú	  r@   rû	  Ú	transposerŠ   rP   rç	  rR   )rS   s   "rT   rU   rU   1  s0   € ÷ "ñ "Ø	ð"àð"ð ð"ð #ñ	"rV   c           	     ó¨  a€ ^ RI Ho V P                  4       pVP                  4       pV'       d   V^,          V^ ,          3pV'       d'   \        P                  ! \
        P                  V^4      M^pR V3R llp\        \        V4      ^8¼  d   V^ ,          V^,          3M
V^ ,          ^3\        V4      VV P                  VP                  VR7      # )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¼  c               ó$   € V ^8„  d   QhRRRRRR/# )rO   rZ  r   r[  rP   rŠ   rR   )rS   s   "rT   rU   Ú,infer_scale_swizzle_ir.<locals>.__annotate__H  s!   € ÷ >ñ >�sð >˜sð > tñ >rV   c                óN   <€ SP                   P                  P                  W4      # )z5Compare values using symbolic equality when possible.)rÀ  rÁ  r  )rZ  r[  r½  s   &&€rT   Úsymbolic_eqÚ+infer_scale_swizzle_ir.<locals>.symbolic_eqH  s   ø€ à�w‰w×Ñ×7Ñ7¸Ó=Ð=rV   rÿ	  )r’  r½  rø  rM  rN  rO  rP  rø	  r^   rq  r»   )rú	  rû	  r
  rá	  râ	  rã	  r
  r½  s   &&&    @rT   Úinfer_scale_swizzle_irr

  1  s´   ø€ õ .à�|‰|‹~€HØ—‘Ó!€J÷ Ø˜Q•K ¨!¥Ð-ˆ÷ DN”)×"Ò"¤8§<¡<°¸QÔ?ÐST€K÷>ð >ô %Ü/2°8«}ÀÔ/A�(˜1•+˜x¨�{Ñ+ÈÐQRÍÐUVÐGWÜ˜Ó$ØØ—)‘)Ø—K‘KØôð rV   )rJ   rÖ  rK   Úmtia)é   éd   )r
  r
  Fr’  )rX   rJ   )rR   rœ  rœ  g      ð?rJ   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._linalg_eigh.defaultzaten._assert_scalar)NNTrJ  )rŸ  rÖ  rê	  rŸ   )FTFN)rC  rK   )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.default(¿  Ú__conditional_annotations__Ú
__future__r   rÔ  r6  rà  ÚenumrM  r‘  rD  rã  r,  r  r+  rO  r¥  r  rê   rÂ  rí   r€  rˆ	  r¼  r
  rŸ  Úcollections.abcr   r   r   r   r   r	   r
   r   r   Útypingr   r   r   r   r   r   r   r   r   Útyping_extensionsr   r   r   r   r�   r\   Útorch.utils._pytreer}  Ú_pytreer¶	  Ú$torch._inductor.analysis.device_infor   Útorch._inductor.runtime.hintsr   Ú!torch.fx.passes.regional_inductorr   Útorch.utils._dtype_abbrsr   Útorch.utils._indented_bufferr   r   r    r!   Útorch.utils._ordered_setr"   r#   r$   r  r%   ÚOPTIMUS_EXCLUDE_POST_GRADÚ#torch.fx.experimental._size_hintingr)   r`  r*   r+   r,   r-   r.   r/   r0   r1   Úpathlibr2   r3   r4   r5   Útorch._prims_commonr6   Útorch.fxr7   Útorch.fx.noder8   rë	  r9   r­  r:   r´  r<   Údependenciesr=   rÀ  r?   rv  r@   rA   rB   rC   rD   rE   Úoutput_coderG   r2  rH   rI   rZ   rL   rî   rd   Útorch._dynamo.device_interfacere   Útorch._dynamo.utilsrf   Útorch.autogradrg   Útorch.autograd.profiler_utilrh   Ú(torch.fx.passes.graph_transform_observerri   Útorch.fx.passes.shape_proprj   Útorch.utils._sympy.functionsrk   rl   rm   rn   ro   rp   Útorch.utils._sympy.symbolrq   rr   Útorch.utils._sympy.value_rangesrs   rt   rw  rv   Úruntime.runtime_utilsrw   ry  Ú_IS_WINDOWSÚ	getLoggerr    rä   r  rz   r¾  rw  Ú	VarRangesrù  rƒ   Ú	InputTypeÚgetenvÚXPU_KERNEL_FORMATÚGPU_KERNEL_BIN_EXTSrü  r:  rí  r¯  rŸ  r—  r   r™  r¢  r›  r�  rÔ   r‘  r“  r•  r‡  rˆ  r‰  r‹  r€   r²   r…   r_   r†   r‘   ÚFunctionr“   rá  r«   rü   r  r$  r  r>  rI  rS  r_  rg  rl  rp  r¬  r±  r¹  rÄ  rÇ  rÔ  rÞ  rÿ  rÒ   r  r  r  r#  r'  r+  r8  r9  r:  ÚFN_TYPEr=  r`  rc  rp  rz  r›  rñ  rù  r  r%  r/  r5  r:  rC  rI  rM  rQ  rU  r\  Ú	frozensetrÝ  r~   r	  rb  ri  rw  r�  r–  r—  rž  r¡  r7  r¨  rÈ  Úclear_on_fresh_inductor_cacheÚclear_inductor_cachesÚfresh_inductor_cacherÓ  ró  rù  rû  rO  rb  rd  rt  rv  rŽ  r•  r™  rž  r£  r±  r·  r¼  rÄ  rÊ  rÎ  rÓ  rà  r  r'  r1  r8  r>  rG  rN  rL  r]  rl  rx  r†  rŽ  rz  r—  rœ  r´  r¸  r¼  rÓ  rß  rå  ré  rí  rð  rû  rö  r#  r%  rG  rP  rX  ru  rz  r}  r‰  r“  r   r¦  rª  rµ  r»  rÁ  rÄ  rÈ  r  rÓ  rê  rñ  r÷  rþ  r  r  r  r�  r  ri  ÚEnumr  r7  r<  rA  rN  rR  rZ  r`  rd  rk  rv  r~  rˆ  r—  rœ  rž  rµ  rº  r»  r¿  rÉ  rÓ  r×  rÛ  rë  rø  rý  r  r  r  r  r,  r4  r'  rA  rI  rY  re  ri  rn  rÚ  rs  Úcompilerm  rp  rv  r}  r€  r�  r’  r•  r˜  r™  rœ  r¥  r©  r­  r®  r¯  r±  rä  ré  rë  rõ  rø  rþ  r
	  ra  r	  r	  r%	  r+	  r/	  ÚSUPPORTED_MKLDNN_DEVICESr6	  r;	  rI	  rS	  ro	  rx	  r}	  r�	  rŒ	  r‘	  r•	  r—	  ÚPartitionFnTyper�	  r›	  r)	  rŸ	  r½	  rÂ	  rÉ	  rÓ	  rÖ	  r×	  rÛ	  rß	  rø	  r
  r

  )ro  rˆ   r
  s   00@rT   Ú<module>rA
     sg  øðß "Ð "ã Û Û Û Û Û Û Û Û 	Û Û Û Û 	Û Û 	Û Û Û 
Û Û Û Û ÷÷ ñ õ Ý ÷
÷ 
õ 
÷ FÑ EÝ ã ã ß $Ð $Ý ?Ý :Ý EÝ 0÷
ó õ 0ß ;Ý 2ð (Ø$ðÐ õ
 <÷õ ÷ ß>Ñ>Ýç/Ñ/ÝCÝ$Ý"Ý/å,Ý5Ý!Ý$ßT×TÝ,ß=ó +€	ÙˆCƒL€ð
 ‡�ôó ðõ DÝ 0Ý %Ý 2Ý KÝ 0÷÷ ÷ 8ß Då Ý =ð �l‰l˜gÑ%€à×Ó˜Ó!€ð  CÐ ñ ˆTƒ]€Ø�—’˜UŸZšZÐ'Õ(€	Ø�LŠL˜3Õ §¡Õ-°Õ4€	÷ �E˜bŸi›iÐ(IÈ7ÓSð ð
 ˆHØ	ˆ8Ø	ˆQÐ Ð!Ð"ðÐ ð €Ø€	à€ØÐ ñ
 2<à�ŠØ�
Š
Ø�ŠØ�ŠØ�ŠØ�ŠØ�ŠØ�ŠØ�ŠØ�ŠØ�ŠØ×ÒØ×ÒØ×ÒØ×Òðó2Ð Ð.ó ð( €Ø	˜ a�Õ	(¨QÔ.°;À!Ô3CÙ
Ð-Ó
.Ð.õ5õ
LôˆE�NŠNô ð ×Ó˜dÔ#÷ð ó $ð÷(Gð G÷Tð õ86.÷rNð Nðb ‡�ôó ðõ õ;õ@õõ õ+õ*õ%AõP+õ	õõ"/õGñ @O÷Iõ0÷0'ð '÷ð ÷&ð õ )õ'õ#õõ$ñ" ˆcƒN€ÙˆT˜TÔ"€Ø
�;˜s A˜vÕ&¨Ð*Õ
+€ôE�8˜W Q¨ U�^ô Eõõ:õ+õ\õ,<õ~h2÷Vð õ(Gõ:,õ^%õõ(õ
DõUõ
>õ+õõ-ñ $óð, �=Š=×ÒÒ(ð  ñ	
ð õ/óÐ õ:õõ'ó& 
õ"ðJ !#Ñ �IÓ "ö	öð ×Òô"ó ñ"ð( ×Ò÷8ó ñ8ñx !5Ñ Ù$Ñ Ù"Ñ ð°5÷ ð ð 7ð ÷	7ð 7ðt ×Ò�QÓô7ó ñ7ôL˜õ Lô^-�Nõ -ö,ô 
˜õ 
ð ×Òô@ó ñ@ô?Ð'õ ?ð ‡�÷ó ñð: ×ÑôJó ñJð ×Ñô)ó ñ)ööI÷ñ ÷(ñ öö öööðð ðð  ð	ð
  $÷ð ð:BØ6:ðBØOT÷Bð B÷Jñ ð2OØAF÷Oð Oð0VØAF÷Vð Vö Oð ×Ò˜QÔô	ó  ñ	ð ×Ò˜QÔõó  ñô=ð0 ×Ò˜QÔöó  ñ÷>ð >÷Bð ÷BEó E÷PJð Jð ˜EŸJšJÕ&�‰)Ô &ð ‡�÷ñ ó ñð( ‡�öó ñð. ‡�ö5(ó ñ5(ðp ‡�ö@ó ñ@ð ×Ñöó ñð ‡�öRó ñR÷:"ð "÷Jð ÷ð ÷Hð H÷&ð &÷ð ÷4=ó =÷@'ð '÷Cô C÷"&ð &÷$ð ÷&/ð /÷(ð (÷Vð ÷
ð 
÷#ð #ð, ×Òö.ó ñ.÷"Fó F÷*	ð 	÷ð ÷B&ð &÷&ð &÷ð ÷ð ð ×Òöó ñ÷ð ð ‡�öDó ñDð ‡�ö26ó ñ26ðj ‡�öó ñ÷ Qð Q÷ð ÷0ð 0÷
ð 
÷(ð (÷#ð #÷Kð K÷ð ö(*�$—)’)õ *÷ 
ð  
÷F	ð 	÷ð ÷#ó #÷L(ð (÷	Qó 	Q÷#ð #÷Jð J÷Lð Lò *=÷
ò 
ò< *=÷ò ÷<Tð T÷ ð  ÷8ð ð  ×Ò÷ò ó ñð ×Òöó ñ÷,!ð !÷ð ÷)ð )ð
 ×Ò�TÓö'ó ñ'ó,10ð "&ñ	10ð
 %)ñ10ð !÷10ó 10÷h.ð .÷2ð 2÷$ð $÷N!Hð !H÷Hð ÷Lð L÷'ð '÷ ð ÷2ð ÷:ð ÷0	:ð 	:÷$ó $÷:ð ÷$*Ið *I÷Zð ÷Bð ÷&/ð /ñ ‰yÙ™Ù‘mÙ™/ÙÑ*ñ ™Ù™:ð
Ñ ñ +?×*DÒ*DÔ*FÔGÑ*F¡$ !˜šÑ*FÒGÑ ð —*“*™YÓ'�÷Hð H÷ð ÷	ð 	÷ð ð ‡�öó ñð& ×Ò÷.ò .ó ñ.ð
 68Ñ Ñ2Ô 7÷ð ñ /9«lÑ ™OÔ :÷)ð )÷#ó #÷"ð ÷ð ñ ‘)Ó
�Ù
‘)Ó
�ö-"�¡©Ð 0Õ1õ -"ñ`  DÕ)óÀD÷ ñ ó *ñ÷ð ö 4§9¢9õ ð ‡�ö4ó ñ4÷2Yð Y÷5ð 5÷!ð !÷H"ð "÷Jð ÷4
ð 
÷ð ÷1ð 1÷ð ñ *Ñ ÷	ð 	÷	ð 	÷ð ÷& 
ó  
÷FRð R÷j&2ð &2÷Rð ÷6ð ÷ð ÷>$ð $ð  ×Òö#ó ñ#ð& ×Ó˜dÔ#÷ò ó $ñð ™3 ˜8Õ$�ØÙÑ.Ð/±Ð@õÑ ÷0ô 0ò *:Ó);Ñ &÷=ð =÷+ð +÷\
ð 
÷
Lð 
L÷#-ð #-ñN óó�÷$%ð %ð
 öó ñ÷"ð "÷
Yð Y÷xð ÷<"õ "ùóc Hs   ï6{ 