+
    –°³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HtHtHt ^ RIHt ^ RIHt ^ RIHtHt ^ RIHtHt ^ RIt^ RIHt ^ R	IHtHtHt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+H,t,H-t-H.t. ^ RI/H0t0H1t1H2t2 ^ RI3H4t4 ^ RI5H6t6 ^ RI7H8t8H9t9H:t:H;t; ^ RI<H=t= ^ RI>H?t? ^ RI@HAtA ^ RIBHCtCHDtD ^ RIEHFtF ^ RIGHHtH ]'       d   ^ RIIHJtJ ^ RIKHLtL ^ RIMHNtN RsORsPR]QR&   ]P¤                  P§                  ]TR4      tU]P¬                  ! ]T4      tWRsXR]QR&   ]	P²                  ! R4      tZR  R! lt[R" t\R# t]R$ R% lt^R& R' lt_R( R) lt`R* R+ lta]
PÄ                  R,8H  tc]P¬                  ! ]T4      tW]F]0,          ! 4       tdR- R. lteR/ tf ]PÎ                  ! ]fR07       R1 R2 ltiR3 R4 ltj]= ! R5 R64      4       tk ! R7 R84      tlR9 R: ltm]PÜ                  ! ]e4       R#   ]h d     LHi ; i);é    )ÚannotationsN)ÚFutureÚThreadPoolExecutorÚTimeoutError)ÚBrokenProcessPool)Úpartial)ÚtimeÚtime_ns)ÚAnyÚTYPE_CHECKING)Ú get_registered_device_interfaces)ÚcountersÚdynamo_timedÚget_metrics_contextÚset_feature_use)Úconfig)Ú_load_triton_kernel_from_sourceÚ	code_hashÚCodeCacheFutureÚCppCodeCacheÚCppPythonBindingsCodeCacheÚCUDACodeCacheÚHalideCodeCacheÚLambdaFutureÚROCmCodeCacheÚStaticAutotunerFutureÚ	torch_keyÚXPUCodeCache)ÚAnyPoolÚSubprocExceptionÚSubprocPool)ÚTrackedProcessPoolExecutor)Ú_async_compile_initializer)Ú_set_triton_libdevice_pathÚ_set_triton_ptxas_pathÚ"_worker_compile_pycodecache_kernelÚ_worker_compile_triton)Úclear_on_fresh_cache)ÚV)Úlog_triton_builds)Ú	_FaketqdmÚtqdm)Ú
OrderedSet)Úhas_triton_package)ÚCallable)Ú
HalideMeta)ÚCachingAutotunerg        zfloat | NoneÚ_t0Úkernel_codez dict[str, dict[str, Any]] | NoneÚ_triton_kernel_metricszsize_hints=(\{.*?\})c               ó   € V ^8„  d   QhRR/# )é   Úreturnzdict[str, str | None]© )Úformats   "Ú^/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torch/_inductor/async_compile.pyÚ__annotate__r;   Y   s   € ÷ 4ñ 4Ð)>ñ 4ó    c                 ór   € . ROp V  Uu/ uF"  q\         P                  P                  V4      bK$  	  up# u upi )ÚTORCHINDUCTOR_CACHE_DIR)r>   ÚTRITON_CACHE_DIRÚTORCHINDUCTOR_CUTLASS_DIR)ÚosÚenvironÚget)Úenv_varsÚvs     r:   Ú_pycodecache_kernel_compile_envrF   Y   s3   € ò€Hñ
 +3Ó3©( QŒr�z‰z�~‰~˜aÓ Ò ©(Ñ3Ð3ùÒ3s   ‰(4c                 óL   € \        4        ^ RIHp Hp V '       d
   V! 4        R# R# )z?
Setup that must be done prior to forking with a process pool.
©Ú
HAS_TRITONÚ
triton_keyN)Úcaching_device_propertiesÚ%torch._inductor.runtime.triton_compatrI   rJ   rH   s     r:   Úpre_fork_setuprM   b   s   € ô Ô÷ MçÙŽñ r<   c                 ó�   € \        4        F7  w  rVP                  4       '       g   K  VP                  P                  4        K9  	  R # ©N)r   Úis_availableÚWorkerÚget_device_properties)Ú_Údevice_interfaces     r:   rK   rK   r   s4   € Ü?ÖAÑˆØ×(Ñ(×*Ô*Ø×#Ñ#×9Ñ9Ö;ó  Br<   c               ó   € V ^8„  d   QhRR/# ©r6   r7   ÚNoner8   )r9   s   "r:   r;   r;   x   s   € ÷ $ñ $˜ñ $r<   c                 óB   € \         f   \        4       s \        f   / sR # R # rO   )r2   r	   r4   r8   r<   r:   Ú_compile_startrY   x   s   € ä
‚{Ü‹fˆÜÒ%Ø!#Òñ &r<   c               ó   € V ^8„  d   QhRR/# rV   r8   )r9   s   "r:   r;   r;   €   s   € ÷ &ñ &�dñ &r<   c                 ó  a€ \         e%   \        4       p \        V \         ,
          ,          sR s \        '       dR   \	        \        \        P                  4       4      4      o\        P                  P                  RR V3R lR7       R sR # R # )NÚartifactc                 ó   € R RRR/# )ÚnameÚtriton_kernel_infoÚencodingÚjsonr8   r8   r<   r:   Ú<lambda>Ú_compile_end.<locals>.<lambda>Œ   s   € ØÐ,Ø˜Fñ!r<   c                 ó0   <€ \         P                  ! S 4      # rO   )ra   Údumps)Úsorted_infos   €r:   rb   rc   �   s   ø€ œtŸzšz¨+Ô6r<   )Úmetadata_fnÚ
payload_fn)
r2   r	   Ú_cumulative_compile_timer4   ÚdictÚsortedÚitemsÚtorchÚ_loggingÚtrace_structured)Út1rf   s    @r:   Ú_compile_endrq   €   ss   ø€ ä
‚Ü‹VˆÜ  B¬¥HÕ,Ð ØˆçÓäœ6Ô"8×">Ñ">Ó"@ÓAÓBˆÜ�‰×'Ñ'Øñô 7ð 	(ô 	
ð "&Òñ r<   c               ó    € V ^8„  d   QhRRRR/# )r6   Úkernel_nameÚstrÚinfoúdict[str, Any]r8   )r9   s   "r:   r;   r;   •   s   € ÷ 3ñ 3¨ð 3°Nñ 3r<   c                ó,   € \         e   V\         V &   R # R # rO   )r4   )rs   ru   s   &&r:   Ú_add_triton_kernel_inforx   •   s   € ô Ò)Ø.2Ô˜{Ó+ñ *r<   c               ó(   € V ^8„  d   QhRRRRRRRR/# )	r6   Úkernelr1   rs   rt   Ú
elapsed_usÚintr7   rW   r8   )r9   s   "r:   r;   r;   œ   s0   € ÷ ñ Øðàðð ðð 
ñ	r<   c                ó„   € V P                   ;'       g    / pW#R&   \        W4       \        4       P                  RW4       R# )aS  Emit per-kernel ``compile_time_us`` to both the dynamo
``_triton_kernel_metrics`` map and the ``MetricsContext`` top-N.

Note: ``kernel.autotune_cache_info`` is only mutated in place when
non-empty; when ``None`` or ``{}`` the metric still reaches
``_triton_kernel_metrics`` via a throwaway dict, but the kernel
attribute stays untouched.
Úcompile_time_usÚtriton_kernel_compile_times_usN)Úautotune_cache_inforx   r   Ú	add_top_n)rz   rs   r{   ru   s   &&& r:   Ú"_emit_triton_kernel_compile_metricr‚   œ   s>   € ð ×%Ñ%×+Ð+¨€DØ(Ð	ÑÜ˜KÔ.ÜÓ×#Ñ#Ø(¨+ör<   Úwin32c               ó   € V ^8„  d   QhRR/# rV   r8   )r9   s   "r:   r;   r;   ¹   s   € ÷ ñ  $ñ r<   c                 ój   € \          F  p V P                  4        K  	  R\        n        \	        4        R# )z/Shut down all outstanding compile-worker pools.N)Ú	_pool_setÚshutdownÚAsyncCompileÚ_ready_futureÚ
after_fork)Úpools    r:   Úshutdown_compile_workersrŒ   ¹   s"   € çˆØ�‰Žñ à!%„LÔÜ†Lr<   c                 ó€   € \         P                  4        R\        n        \        P                  P                  4        R# )z7Reset pools to initial state without shutting them downN)r†   Úclearrˆ   r‰   Úprocess_poolÚcache_clearr8   r<   r:   rŠ   rŠ   Á   s&   € ä‡O�OÔØ!%„LÔÜ×Ñ×)Ñ)Ö+r<   )Úafter_in_childc               ó   € V ^8„  d   QhRR/# )r6   r7   r|   r8   )r9   s   "r:   r;   r;   Î   s   € ÷ "ñ "˜Sñ "r<   c                 ó‚   € \         P                  f   \         P                  ! 4       \         n        \         P                  # )zq
Temporary for internal rollout. Assign config.compile_threads lazily and return it.
TODO: remove after rollout.
)r   Úcompile_threadsÚdecide_compile_threadsr8   r<   r:   Úget_compile_threadsr–   Î   s-   € ô
 ×ÑÒ%Ü!'×!>Ò!>Ó!@ŒÔÜ×!Ñ!Ð!r<   c               ó   € V ^8„  d   QhRR/# )r6   r7   Úboolr8   )r9   s   "r:   r;   r;   Ø   s   € ÷ ñ ˜tñ r<   c                 ó€   € \         P                  R 8H  ;'       g%    \        P                  ! 4       P                  '       * # )Ú
subprocess)r   Úworker_start_methodÚmultiprocessingÚcurrent_processÚdaemonr8   r<   r:   Ú_process_pool_allowedrŸ   Ø   s6   € ô 	×"Ñ" lÑ2÷ 	8ð 	8Ü×.Ò.Ó0×7Ñ7Ô7ðr<   c                  ó’   € ] tR t^ãt$ Rt/ tR]R&   ]R R l4       t]R R l4       t	]R R	 l4       t
]R
 4       t]R R l4       tRtR# )ÚCompiledTritonKernelsa  
In memory cache for storing compiled triton kernels.

Each triton kernel is keyed by the hash of its source code. Each value stored
in the cache is a return value of AsyncCompile.triton().

Currently, the cache stores Future objects, but it should be generalizable for any kernels.
zdict[str, CodeCacheFuture]Ú_cachec               ó   € V ^8„  d   QhRR/# )r6   Ú
kernel_srcrt   r8   )r9   s   "r:   r;   Ú"CompiledTritonKernels.__annotate__ñ   s   € ÷ 8ñ 8˜ñ 8r<   c                ó,   € \        V \        4       R7      # )zò
Generates a cache key given a triton kernel's full source code.
This source includes the inductor meta, compilation metadata, the kernel itself, etc.
`kernel_src` should be the exact string passed to async_compile.triton()'s first argument.
)Úextra)r   r   )r¤   s   &r:   ÚkeyÚCompiledTritonKernels.keyð   s   € ô ˜¬9«;Ô7Ð7r<   c               ó    € V ^8„  d   QhRRRR/# )r6   r¤   rt   Úfuturer   r8   )r9   s   "r:   r;   r¥   û   s   € ÷ 
3ñ 
3˜ð 
3 oñ 
3r<   c                óV   € \         P                  V 4      pV\         P                  V&   R# )a]  
Saves a compiled triton kernel to the cache.
TODO: We store a LambdaFuture as that's the callable returned by async_compile.triton,
but the real type we want to return here is actually an abstract triton kernel.

TODO: Source code here is not just the kernel's source code, but also includes the inductor preamble, etc.
so it could be less strict.
N©r¡   r¨   r¢   )r¤   r«   r¨   s   && r:   ÚsaveÚCompiledTritonKernels.saveú   s%   € ô $×'Ñ'¨
Ó3ˆØ,2Ô×$Ñ$ SÓ)r<   c               ó    € V ^8„  d   QhRRRR/# )r6   r¤   rt   r7   zCodeCacheFuture | Noner8   )r9   s   "r:   r;   r¥     s   € ÷ ;ñ ;˜ð ;Ð 6ñ ;r<   c                	ól   € \         P                  V 4      p\         P                  P                  VR 4      # rO   )r¡   r¨   r¢   rC   ©r¤   r¨   s   & r:   rC   ÚCompiledTritonKernels.get  s+   € ä#×'Ñ'¨
Ó3ˆÜ$×+Ñ+×/Ñ/°°TÓ:Ð:r<   c                 	ó   € / \         n        R # rO   )r¡   r¢   r8   r<   r:   r�   Ú!CompiledTritonKernels.cache_clear  s   € à')ÔÖ$r<   c               ó    € V ^8„  d   QhRRRR/# )r6   r¤   rt   r7   rW   r8   )r9   s   "r:   r;   r¥     s   € ÷ 2ñ 2 #ð 2¨$ñ 2r<   c                	ó€   € \         P                  V 4      pV\         P                  9   d   \         P                  V R # R # rO   r­   r²   s   & r:   Úremove_futureÚ#CompiledTritonKernels.remove_future  s6   € ä#×'Ñ'¨
Ó3ˆð Ô'×.Ñ.Ô.Ü%×,Ñ,¨SÒ1ñ /r<   r8   N)Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__r¢   Ú__annotations__Ústaticmethodr¨   r®   rC   r�   r¸   Ú__static_attributes__r8   r<   r:   r¡   r¡   ã   sv   ‡ ñð *,€FÐ&Ó+àô8ó ð8ð ô
3ó ð
3ð ô;ó ð;ð ñ*ó ð*ð ô2ó ô2r<   r¡   c                  ó.  € ] tR tRt$ RtRtR]R&   RtR]R&   R R	 lt]	]
P                  ! ^4      R
 R l4       4       t]	R 4       t]	]
P                  ! ^4      R R l4       4       t]R R l4       t]R9R R ll4       t]R R l4       t]R 4       t]R9R R ll4       t]R R l4       tR:R 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* t R<R+ R, llt!R- R. lt"R<R/ R0 llt#R=R1 lt$R2 R3 lt%R4 R5 lt&R6 R7 lt'R8t(R# )>rˆ   i  zS
Utilities to compile in thread pools or subprocess pools (in the case of Triton).
NzFuture[Any] | Noner‰   z'list[tuple[str, str, list[str]]] | NoneÚ_metal_sourcesc               ó   € V ^8„  d   QhRR/# rV   r8   )r9   s   "r:   r;   ÚAsyncCompile.__annotate__!  s   € ÷ ñ ˜$ñ r<   c                	ó   € R # rO   r8   )Úselfs   &r:   Ú__init__ÚAsyncCompile.__init__!  s   € Ùr<   c               ó   € V ^8„  d   QhRR/# )r6   r7   r   r8   )r9   s   "r:   r;   rÅ   &  s   € ÷ 9ñ 9Ð$ñ 9r<   c                 	ór   € \        4       ^8:  d   \        R\        4        24      h\        \        4       4      # )é   ú(expected get_compile_threads() > 1, got )r–   ÚAssertionErrorr   r8   r<   r:   r‹   ÚAsyncCompile.pool$  s<   € ô Ó  AÔ%Ü Ø:Ô;NÓ;PÐ:QÐRóð ô "Ô"5Ó"7Ó8Ð8r<   c                 ó   € R# )z>No-op function to help mark when the subprocess pool is ready.Úreadyr8   r8   r<   r:   Ú
_get_readyÚAsyncCompile._get_ready-  s   € ñ r<   c               ó   € V ^8„  d   QhRR/# )r6   r7   r   r8   )r9   s   "r:   r;   rÅ   4  s   € ÷ -ñ -˜'ñ -r<   c            	     	ó   € \        4       ^8:  d   \        R\        4        24      h\        4       '       g   \        R4      hR\        n        \        P                  R\        P                  \        4       4       \        P                  R8X  d%   \        \        4       \        P                  R7      p M�\        P                  R8X  d   R\        P                  R	&   \        4        \        P                   ! \        P                  4      p\#        \        4       V\%        \&        \        P(                  ! 4       4      R
7      p \        P*                  P-                  RV P.                  \0        P2                  R7       \4        P7                  V 4       V # )rÌ   rÍ   a  Inductor async compile process pools are disabled in daemonic multiprocessing processes. Set torch._inductor.config.worker_start_method = 'subprocess' (or TORCHINDUCTOR_WORKER_START=subprocess) to use the SubprocPool path, which is not affected by the daemon restriction.Nz"Creating '%s' pool with %d workersrš   )ÚquiesceÚspawnÚ0ÚTORCH_WARM_POOL)Ú
mp_contextÚinitializer)Úexitpriority)r–   rÎ   rŸ   ÚRuntimeErrorrˆ   r‰   Úlogru   r   r›   r!   Úquiesce_async_compile_poolrA   rB   rM   rœ   Úget_contextr"   r   r#   ÚgetpidÚutilÚFinalizer‡   ÚsysÚmaxsizer†   Úadd)r‹   Úctxs     r:   r�   ÚAsyncCompile.process_pool2  s.  € ô Ó  AÔ%Ü Ø:Ô;NÓ;PÐ:QÐRóð ô %×&Ò&ÜðUóð ð &*ŒÔ"Ü�‰Ø0Ü×&Ñ&ÜÓ!ô	
ô ×%Ñ%¨Ô5äÜ#Ó%¬v×/PÑ/Pô‰Dô ×)Ñ)¨WÔ4à03”—
‘
Ð,Ñ-ÜÔÜ!×-Ò-¬f×.HÑ.HÓIˆCÜ-Ü#Ó%ØÜ#Ô$>ÄÇ	Â	ÃÓLôˆDô 	×Ñ×%Ñ% d¨D¯M©MÌÏÉÐ%ÔTä�‰�dÔØˆr<   c               ó   € V ^8„  d   QhRR/# rV   r8   )r9   s   "r:   r;   rÅ   d  s   € ÷ ñ ˜$ñ r<   c                	ó�   € \        4       ^8:  g   \        4       '       g   R# \        4        V P                  4        \	        4        R# )rÌ   N)r–   rŸ   rY   r�   rq   ©Úclss   &r:   Ú	warm_poolÚAsyncCompile.warm_poolc  s2   € äÓ  AÔ%Ô-B×-DÒ-DÙÜÔð 	×ÑÔÜŽr<   c               ó   € V ^8„  d   QhRR/# rV   r8   )r9   s   "r:   r;   rÅ   o  s   € ÷ 6ñ 6¨Tñ 6r<   c                	ó~   € V P                  4        V P                  e   V P                  P                  VR7       R # R # )N©Útimeout)Úuse_process_poolr‰   Úresult)rì   rò   s   &&r:   Úwait_pool_readyÚAsyncCompile.wait_pool_readyn  s6   € à×ÑÔØ×ÑÒ(Ø×Ñ×$Ñ$¨WÐ$Ö5ñ )r<   c               ó    € V ^8„  d   QhRRRR/# )r6   ÚtaskzCallable[..., Any]r7   r   r8   )r9   s   "r:   r;   rÅ   u  s   € ÷ 'ñ 'Ð,ð '°ñ 'r<   c                	ól   € \        4       ^8:  d   V! 4       # V P                  4       P                  V4      # )rÌ   )r–   r‹   Úsubmit)rì   rø   s   &&r:   rú   ÚAsyncCompile.submitt  s+   € äÓ  AÔ%Ù“6ˆMØ�x‰x‹z× Ñ  Ó&Ð&r<   c                	ó<  € \        4       ^8:  g   \        4       '       g   R# \        P                  P                  '       d   R# V P
                  '       g/   V P                  4       P                  V P                  4      V n        V P
                  P                  4       # )rÌ   F)
r–   rŸ   r   ÚtritonÚproton_profilingr‰   r�   rú   rÒ   Údonerë   s   &r:   ró   ÚAsyncCompile.use_process_poolz  sq   € äÓ  AÔ%Ô-B×-DÒ-DÙô
 �=‰=×)×)Ð)Ùð
 × × Ð Ø #× 0Ñ 0Ó 2× 9Ñ 9¸#¿.¹.Ó IˆCÔØ× Ñ ×%Ñ%Ó'Ð'r<   c               ó    € V ^8„  d   QhRRRR/# )r6   rò   Úfloatr7   r˜   r8   )r9   s   "r:   r;   rÅ   �  s   € ÷ ñ ¨eð ¸dñ r<   c                óö  € \        4       ^8:  g   \        4       '       g   R# \        P                  P                  '       d   R# V P
                  '       g/   V P                  4       P                  V P                  4      V n         V P
                  P                  VR7       R#   \         d    \        P                  RT4        R# \        \        3 d"   p\        P                  RT4        Rp?R# Rp?ii ; i)a  Block (up to ``timeout`` s) until the process pool is ready, returning
whether it's usable.

Like use_process_pool() but blocking. Use when a backend's serial
fallback is far costlier than the warmup wait -- e.g. NVGEMM subprocess
precompile, where skipping the pool forces ~15x-slower lazy compilation
at benchmark time. (use_process_pool()'s non-blocking readiness check can
race pool warmup when little other compilation precedes the decision.)
On timeout, degrade gracefully (return False -> serial) rather than hang
on a stuck worker.
Frñ   z8Process pool not ready after %ss; falling back to serialz2Process pool unusable (%s); falling back to serialNT)r–   rŸ   r   rý   rþ   r‰   r�   rú   rÒ   rô   ÚFuturesTimeoutErrorrÞ   Úwarningr   rÝ   )rì   rò   Úes   && r:   Úwait_process_pool_readyÚ$AsyncCompile.wait_process_pool_readyŒ  sÄ   € ô Ó  AÔ%Ô-B×-DÒ-DÙÜ�=‰=×)×)Ð)ÙØ× × Ð Ø #× 0Ñ 0Ó 2× 9Ñ 9¸#¿.¹.Ó IˆCÔð	Ø×Ñ×$Ñ$¨WÐ$Ô5ñ øô #ô 	Ü�K‰KØJÈGôò Ü!¤<Ð0ô 	ô �K‰KÐLÈaÔPÝûð	ús$   ÂB# Â#!C8ÃC8ÃC8ÃC3Ã3C8c               ó   € V ^8„  d   QhRR/# rV   r8   )r9   s   "r:   r;   rÅ   ¯  s   € ÷ 	ñ 	�tñ 	r<   c                ó¦   € V P                  4       '       g   R# V P                  4       p\        V\        4      '       d   VP	                  4        R# R# )zY
If using a SubprocPool, signal the sidecar process to start up its
ProcessPoolExecutor.
N)ró   r�   Ú
isinstancer!   Úwakeup)rì   r‹   s   & r:   r  ÚAsyncCompile.wakeup®  s?   € ð ×#Ñ#×%Ò%ÙØ×ÑÓ!ˆÜ�dœK×(Ò(Ø�K‰KŽMñ )r<   c               ó$   € V ^8„  d   QhRRRRRR/# )r6   rs   rt   Úsource_codeÚ
device_strr8   )r9   s   "r:   r;   rÅ   º  s&   € ÷ Y1ñ Y1 #ð Y1°Cð Y1ÀSñ Y1r<   c                ó¼  aaaaaaa€ \         P                  ! \        SS4      oV3R lo\        R,          R;;,          ^,          uu&   \        P                  RS4       \        4        \        P                  P                  RR4      R8X  d>   \        \        P                  P                  P                  P                  S4      S4      # V P!                  4       p\#        RV4       \        P$                  P&                  P)                  4       o\        \*        P,                  R	R
4      o\.        P                  S4      ;pem   \        R,          R;;,          ^,          uu&   \1        V\2        4      '       d   \.        P5                  S4       SVn        V'       d   V# VP9                  4       # V'       Ed§   \;        4        . ROpV Uu/ uF"  qw\        P                  P                  V4      bK$  	  ppR\        P                  P<                  P>                  /p	\A        \        P                  P<                  PB                  4      ^ 8”  dº   \D        PG                  S4      p
V
'       d   V
PI                  ^4      pM\K        R4      pSPM                  R4      ^,          p^ RI'H(p V! W¼4      pV\        P                  P<                  PB                  9   d1   V\        P                  P<                  PB                  V,          /V	R&   V PS                  4       PU                  \V        SVV	4      oR VVVVVV3R llp\Y        VSR7      p\.        P[                  SV4       V# \]        RRRRRR7      ;_uu_ 4        Rp \_        4       p\a        4        \;        4        S! 4       pVPc                  SS4       VPe                  R
\.        Pg                  S4      R7       \_        4       V,
          R,          p\i        VSV4       V\k        VR7       uuRRR4       # u upi   \l         d   p\K        T4      ph Rp?ii ; i  \k        TR7       i ; i  + '       g   i     R# ; i)aŽ  
Async_compile.triton is more complicated than the other backends because
we're trying to optimize compile time as much as possible for this hot callsite.

First of all, the function is cached by CompiledTritonKernels; if there's a kernel
already compiled, we grab it directly from the cache and return.

Otherwise, if we have multiple compile threads, we kick off triton compilations on each
worker process by giving it a kernel and source code to compile. The worker initializes
a CachingAutotuner, runs triton compilation, and pickles the kernel back to us.
We use TritonCompileResult to represent the objects being pickled back to us by each
worker.

Some maybe not obvious things that are pickled back to us:
- Most of the time, we can avoid sending back CachingAutotuner.fn and other metadata
  and do not have to pay the cost of loading the triton kernel on the parent. But certain
  cases, like coordesc tuning and dynamic_scale_rblock, require us to reload the function
  in the parent lazily when we require it.
- The AutotuneCache, if enabled, is constructed on each worker per triton config
  and pickled by to us via `CachingAutotuner.save_cache_hook`.
c                 óv   <€ \        R 4      ;_uu_ 4        S ! 4       uuRRR4       #   + '       g   i     R# ; i)Úreload_kernel_in_parentN)r   )Úload_kernels   €r:   r  Ú4AsyncCompile.triton.<locals>.reload_kernel_in_parentÔ  s#   ø€ äÐ7×8Õ8Ù"“}÷ 9×8×8Ó8ús   •'§8	ÚinductorÚasync_compile_cache_misszTriton Kernel:
%sÚTRITON_INTERPRETrØ   Ú1Úparallel_compile_post_warmupÚis_backwardFNÚasync_compile_cache_hitÚuse_static_triton_launcherz@triton.jit
)Ú%generate_lookup_hash_from_source_codeÚautotune_lookup_tablec               ó   € V ^8„  d   QhRR/# )r6   r7   r1   r8   )r9   s   "r:   r;   Ú)AsyncCompile.triton.<locals>.__annotate__!  s   € ÷ ñ Ð 0ñ r<   c                 óR  <€  SP                  4       w  rT P                  SS4       \        P                  S4       T P                  R R7       T P                  RS\        P                  S4      R7       \        T ST4       T #   \         d   pTP                  S4      ThR p?ii ; i)N)Ú
old_valuesF)Úwarm_cache_onlyÚreload_kernelÚstatic_triton_bundle_key)
rô   r    Ú	with_nameÚset_compile_infor¡   r¸   Úrestore_after_unpickleÚ
precompiler¨   r‚   )	rz   r{   r  Ú
compile_idr  rs   r  r  rø   s	      €€€€€€r:   Ú
get_resultÚ'AsyncCompile.triton.<locals>.get_result!  s¤   ø€ ð:Ø)-¯©«Ñ&�Fð ×'Ñ'¨
°KÔ@Ü%×3Ñ3°KÔ@à×-Ñ-¸Ð-Ô>à×!Ñ!Ø$)Ø"9Ü-B×-FÑ-FÀ{Ó-Sð "ô ô
 3°6¸;È
ÔSØ�øô! (ô :ØŸ+™+ kÓ2¸Ð9ûð:ús   ƒB ÂB&ÂB!Â!B&©r«   zasync_compile.precompileTÚtriton_compile_time_usÚcompile_triton©Úlog_pt2_compile_eventÚdynamo_compile_column_usÚlog_waitcounterÚwaitcounter_name_override)r$  r&  iè  )Úfail)r>   r?   ÚTRITON_LIBDEVICE_PATH)7Ú	functoolsr   r   r   Úkernel_code_logru   rY   rA   rB   rC   Úgetattrrm   Ú	_inductorÚ	codecacheÚPyCodeCacheÚloadró   r   Ú_guardsÚCompileContextÚcurrent_compile_idr)   Úgraphr¡   r  r   r¸   Úreload_kernel_from_srcrô   r$   r   r  Úlenr  Úsize_hints_regexÚsearchÚgrouprt   ÚsplitÚ)torch._inductor.runtime.triton_heuristicsr  r�   rú   r'   r   r®   r   r
   r%   r(  r*  r¨   r‚   r*   Ú	Exception)rÇ   rs   r  r  Úis_parallelr«   rD   rE   Ú	extra_envÚextra_configÚmÚsize_hints_strÚ
triton_srcr  Úfn_hashr,  r6  Ústart_nsrz   r{   r  r+  r  r  r  rø   s   &ff&                 @@@@@r:   rý   ÚAsyncCompile.tritonº  s‰  þ€ ô,  ×'Ò'Ü+¨[¸+ó
ˆõ	%ô
 	�ÕÐ7×8¸AÕ=Ó8ä×ÑÐ1°;Ô?ÜÔä�:‰:�>‰>Ð,¨cÓ2°cÔ9ÜÜ—‘×)Ñ)×5Ñ5×:Ñ:¸;ÓGÈóð ð ×+Ñ+Ó-ˆÜÐ6¸ÔDä—]‘]×1Ñ1×DÑDÓFˆ
ÜœaŸg™g }°eÓ<ˆä+×/Ñ/°Ó<Ð<ˆFÒIÜ�ZÕ Ð!:×;¸qÕ@Ó;ä˜&Ô"7×8Ò8ä%×3Ñ3°KÔ@Ø0G�Ô-ßØ�à—}‘}“Ð&÷ ˆ;ä&Ô(òˆHñ
 8@Ó@±x°!œBŸJ™JŸN™N¨1Ó-Ò-±xˆIÐ@à,¬e¯o©o×.DÑ.D×._Ñ._ðˆLô ”5—?‘?×)Ñ)×?Ñ?Ó@À1ÔDÜ$×+Ñ+¨KÓ8�ßØ%&§W¡W¨Q£Z‘Nä%(¨£Y�Nà(×.Ñ.¨Ó?ÀÕB�
õñ @Ø"ó�ð œeŸo™o×4Ñ4×JÑJÔJà¤§¡×!7Ñ!7×!MÑ!MÈgÕ!Vð=�LÐ!8Ñ9ð ×$Ñ$Ó&×-Ñ-Ü&ØØØó	ˆD÷õ ô* " *°TÔ:ˆFÜ!×&Ñ& {°FÔ;ØˆMäØ*Ø&*Ø)AØ $Ø*:÷ö ð �ð1Ü&›y�HÜ*Ô,Ü.Ô0Ù(›]�FØ×+Ñ+¨J¸ÔDØ×%Ñ%Ø(-Ü1F×1JÑ1JÈ;Ó1Wð &ô ô #*£)¨hÕ"6¸4Õ!?�JÜ6°v¸{ÈJÔWØ!ô
 &¨4Õ0÷3ò ùòu Aøô^ !ô Ü˜q›6�DØûðûô &¨4Ö0ú÷3÷ ð úsC   Ç(PÍ7Q
Í;BPÏ>Q
ÐP6Ð%P1Ð1P6Ð6P9Ð9QÑQ
Ñ
Q	c               ó   € V ^8„  d   QhRR/# ©r6   r7   r   r8   )r9   s   "r:   r;   rÅ   U  s   € ÷ 0ñ 0¨sñ 0r<   c                	ó   € ^ RI Hp V! V/ VB # )r   )ÚMultiKernelCall)Ú$torch._inductor.codegen.multi_kernelrW  )rÇ   ÚargsÚkwargsrW  s   &*, r:   Úmulti_kernelÚAsyncCompile.multi_kernelU  s   € ÝHñ  Ð/¨Ñ/Ð/r<   c               ó   € V ^8„  d   QhRR/# rU  r8   )r9   s   "r:   r;   rÅ   [  s   € ÷ 8ñ 8¸ñ 8r<   c                	ó   € ^ RI Hp V! V/ VB # )r   )ÚSizeHintMultiKernelCall)rX  r_  )rÇ   rY  rZ  r_  s   &*, r:   Úsize_hint_multi_kernelÚ#AsyncCompile.size_hint_multi_kernel[  s   € ÝPá&¨Ð7°Ñ7Ð7r<   c               ó   € V ^8„  d   QhRR/# )r6   r  rt   r8   )r9   s   "r:   r;   rÅ   `  s   € ÷ =ñ =˜sñ =r<   c                	óî   a€ \         P                  R V4       \        4       ^8:  d!   \        P                  ! V4      P
                  # \        P                  ! WP                  R7      o\        V3R l4      # )zCPP Kernel:
%s©Ú	submit_fnc                 ó&   <€ S ! 4       P                   # rO   )rz   )r,  s   €r:   rb   Ú"AsyncCompile.cpp.<locals>.<lambda>f  s   ø€ ©
«×(;Ò(;r<   )	r9  ru   r–   r   r>  rz   Ú
load_asyncrú   r   )rÇ   r  r,  s   &&@r:   ÚcppÚAsyncCompile.cpp`  sV   ø€ Ü×ÑÐ.°Ô<ÜÓ  AÔ%Ü×$Ò$ [Ó1×8Ñ8Ð8ä%×0Ò0°ÏÉÔTˆJÜÔ ;Ó<Ð<r<   c               ó    € V ^8„  d   QhRRRR/# )r6   Úargtypesú	list[str]r  rt   r8   )r9   s   "r:   r;   rÅ   h  s   € ÷ ,ñ , ið ,¸cñ ,r<   c                	óÒ   € \         P                  R V4       \        4       ^8:  d   \        P                  ! W4      # \        P
                  ! WV P                  R7      p\        V4      # )zCPP+Bindings Kernel:
%srd  )r9  ru   r–   r   Úload_pybindingÚload_pybinding_asyncrú   r   )rÇ   rl  r  r,  s   &&& r:   Úcpp_pybindingÚAsyncCompile.cpp_pybindingh  sU   € Ü×ÑÐ7¸ÔEÜÓ  AÔ%Ü-×<Ò<¸XÓSÐSä3×HÒHØ°·±ôˆJô   
Ó+Ð+r<   c                	ó>   aaaa€ VVVV3R  lpV P                  V4      # )c                 ó¢   <€ S'       d/   SP                  SR 4      vrSP                  P                  V 4       SP                  SS4      ^ ,          # )Úo)ÚcompileÚaot_kernels_oÚappendr>  )Úoutput_pathrS   Úaot_compileÚ	cache_clsÚdst_file_extr  s     €€€€r:   rø   Ú"AsyncCompile.cutlass.<locals>.tasks  sF   ø€ ßð #,×"3Ñ"3°KÀÓ"E��Ø×'Ñ'×.Ñ.¨{Ô;Ø—>‘> +¨|Ó<¸QÕ?Ð?r<   )rú   )rÇ   r{  r  r|  rz  rø   s   &ffff r:   ÚcutlassÚAsyncCompile.cutlassr  s    û€ ÷	@ð 	@ð �{‰{˜4Ó Ð r<   c                	ó\   € \         P                  R V4       V P                  \        WV4      # )zCUDA Kernel:
%s)r9  ru   r~  r   ©rÇ   r  r|  rz  s   &&&&r:   ÚcudaÚAsyncCompile.cuda}  s&   € Ü×ÑÐ/°Ô=Ø�|‰|œM¨;ÀkÓRÐRr<   c                	ó\   € \         P                  R V4       V P                  \        WV4      # )zXPU Kernel:
%s)r9  ru   r~  r   r�  s   &&&&r:   ÚxpuÚAsyncCompile.xpu�  s&   € Ü×ÑÐ.°Ô<Ø�|‰|œL¨+À[ÓQÐQr<   c                	óf   aaa€ \         P                  R S4       VVV3R lpV P                  V4      # )zROCm Kernel:
%sc                 ó0  <€ S'       d9   \         P                  ! SR R7      vr\         P                  P                  V 4       \        P
                  P                  '       d   \         P                  ! SRR7      p\         P                  ! SS4      ^ ,          # )ru  )r|  Úexe)r   rv  rw  rx  r   ÚrocmÚgenerate_test_runnerr>  )ry  rS   rz  r|  r  s     €€€r:   rø   ÚAsyncCompile.rocm.<locals>.task�  si   ø€ ßÜ"/×"7Ò"7¸ÐRUÔ"V��Ü×+Ñ+×2Ñ2°;Ô?Ü�{‰{×/×/Ð/Ü!×)Ò)¨+ÀEÔJ�Ü ×%Ò% k°<Ó@ÀÕCÐCr<   )r9  ru   rú   )rÇ   r  r|  rz  rø   s   &fff r:   rŠ  ÚAsyncCompile.rocm…  s-   ú€ ô 	×ÑÐ/°Ô=÷	Dð �{‰{˜4Ó Ð r<   c               ó    € V ^8„  d   QhRRRR/# )r6   Úmetar0   r  rt   r8   )r9   s   "r:   r;   rÅ   —  s   € ÷ ,ñ ,˜:ð ,°Cñ ,r<   c                	óÒ   € \         P                  R W4       \        4       ^8:  d   \        P                  ! W4      # \        P
                  ! WV P                  R7      p\        V4      # )zHalide Kernel:
%r
%srd  )r9  ru   r–   r   Úgenerate_halideÚgenerate_halide_asyncrú   r   )rÇ   r�  r  r,  s   &&& r:   ÚhalideÚAsyncCompile.halide—  sU   € Ü×ÑÐ5°tÔIÜÓ  AÔ%Ü"×2Ò2°4ÓEÐEä(×>Ò>Ø¨T¯[©[ôˆJô   
Ó+Ð+r<   c                óœ   € \         P                  P                  P                  P	                  WE4      pV RV 2pV! \        Wg4      VR7      # )zAReload a kernel module from PyCodeCache and wrap the entry point.rS   ©Úkernel_path)rm   r;  r<  r=  Úload_by_key_pathr:  )rÇ   rs   Úmain_suffixÚwrapper_clsr¨   ÚpathÚmodÚmain_func_names   &&&&&&  r:   Ú_load_kernel_wrapperÚ!AsyncCompile._load_kernel_wrapper¡  sD   € ä�o‰o×'Ñ'×3Ñ3×DÑDÀSÓOˆØ'˜=¨¨+¨Ð7ˆÙœ7 3Ó7ÀTÔJÐJr<   c               ó    € V ^8„  d   QhRRRR/# ©r6   rs   rt   r  r8   )r9   s   "r:   r;   rÅ   §  s   € ÷ >Xñ >X 3ð >X°Sñ >Xr<   c           	     óâ  a aaaa€ ^ RI HoHo \        P	                  RV4       \        4        S P                  4       pV'       dM   \        4       pS P                  4       P                  \        SVSVV4      oR VVVV V3R llp\        VSR7      # \        P                  P                  P                  P!                  V4      w  rx\        P                  P                  P                  P#                  Wx4      p	S RS 2p
\%        Wš4      '       gH   \'        V	4       Uu. uF!  p\)        \+        W›4      4      '       g   K  VNK#  	  pp\-        RV
 RV 24      hS! \+        Wš4      VR	7      # u upi )
aÈ  
Compile CuteDSL (CUTLASS Python DSL) kernels.

Args:
    kernel_name: Name of the kernel to be defined
    source_code: Source code of the CuteDSL kernel, as a string
    precompile_metadata: Optional dict with shapes/dtypes for triggering
        real CuTe DSL compilation in the subprocess worker.

Note:
    CuteDSL currently requires source files to do its compilation, there we
    use the PyCodeCache to write the source code to a file and load it.
)ÚCuteDSLKernelWrapperÚMAIN_SUFFIXzCuteDSL Kernel:
%sc               ó   € V ^8„  d   QhRR/# )r6   r7   r£  r8   )r9   s   "r:   r;   Ú*AsyncCompile.cutedsl.<locals>.__annotate__Ë  s   € ÷ ñ Ð 4ñ r<   c                 óÌ   <€  SP                  4       w  rp\        P	                  RST4       SP                  SSSY4      #   \         d   pTP                  S4      ThR p?ii ; i)Nz0CuteDSL kernel %s compiled in subprocess in %dus©rô   r    r'  rÞ   Údebugrž  )	r¨   r›  r{   r  r£  r¤  rs   rÇ   Úsubprocess_tasks	       €€€€€r:   r,  Ú(AsyncCompile.cutedsl.<locals>.get_resultË  st   ø€ ð:Ø,;×,BÑ,BÓ,DÑ)�C˜zô —	‘	ØFØØôð
 ×0Ñ0Ø Ð.BÀCóð øô (ô :ØŸ+™+ kÓ2¸Ð9ûð:úó   ƒA ÁA#ÁAÁA#r.  rS   z-Could not find CuteDSL main kernel function 'ú'. Available callables: r–  )Ú.torch._inductor.codegen.cutedsl.cutedsl_kernelr£  r¤  r9  ru   rY   ró   rF   r�   rú   r&   r   rm   r;  r<  r=  Úwriter˜  ÚhasattrÚdirÚcallabler:  rÝ   )rÇ   rs   r  Úprecompile_metadatarK  rL  r,  r¨   r›  rœ  r�  r^   Ú	availabler£  r¤  rª  s   ff&&         @@@r:   ÚcutedslÚAsyncCompile.cutedsl§  s=  ü€ ÷	
ô
 	×ÑÐ2°KÔ@ÜÔà×+Ñ+Ó-ˆçÜ7Ó9ˆIà"×/Ñ/Ó1×8Ñ8Ü2ØØØØØ#óˆO÷ô ô   
°?ÔCÐCäŸ™×1Ñ1×=Ñ=×CÑCÀKÓP‰IˆCÜ—/‘/×+Ñ+×7Ñ7×HÑHÈÓSˆCà +˜}¨A¨k¨]Ð;ˆNÜ˜3×/Ò/Ü.1°#¬hÓW©h d¼(Ä7È3ÓCU×:VŸT˜T©h�	ÐWÜ"ØCÀNÐCSÐSkÐluÐkvÐwóð ñ (¬°Ó(DÐRVÔWÐWùò Xs   ÄE,Å E,c               ó    € V ^8„  d   QhRRRR/# r¡  r8   )r9   s   "r:   r;   rÅ   ç  s   € ÷ "9ñ "9 #ð "9°Cñ "9r<   c                óÆ   aaaaa€ ^ RI HoHo \        P	                  RS4       VVVV3R lp\        4       ^8:  d   V! 4       # V P                  V4      o\        V3R l4      # )a8  
Compile Pallas (JAX experimental) kernels.

Args:
    kernel_name: Name of the kernel to be defined
    source_code: Source code of the Pallas kernel, as a string

Note:
    Pallas kernels are Python code that uses JAX and Pallas APIs.
    We use the PyCodeCache to write the source code to a file and load it.
)r¤  ÚPallasKernelWrapperzPallas Kernel:
%sc            	     óÂ  <€ \         P                  P                  P                  P	                  S	4      w  r\         P                  P                  P                  P                  W4      pS R S 2p\        W#4      '       gH   \        V4       Uu. uF!  p\        \        W$4      4      '       g   K  VNK#  	  pp\        RV RV 24      hS! \        W#4      VR7      # u upi )rS   z,Could not find Pallas main kernel function 'r­  r–  )rm   r;  r<  r=  r¯  r˜  r°  r±  r²  r:  rÝ   )
r¨   r›  rœ  r�  r^   r´  r¤  r¹  rs   r  s
         €€€€r:   rø   Ú!AsyncCompile.pallas.<locals>.task÷  sÄ   ø€ ÜŸ™×1Ñ1×=Ñ=×CÑCÀKÓP‰IˆCÜ—/‘/×+Ñ+×7Ñ7×HÑHÈÓSˆCð !,˜}¨A¨k¨]Ð;ˆNÜ˜3×/Ò/Ü.1°#¬hÓW©h d¼(Ä7È3ÓCU×:VŸT˜T©h�	ÐWÜ"ØBÀ>ÐBRÐRjÐktÐjuÐvóð ñ '¤w¨sÓ'CÐQUÔVÐVùò Xs   ÂCÂ0Cc                 ó$   <€ S P                  4       # rO   )rô   r.  s   €r:   rb   Ú%AsyncCompile.pallas.<locals>.<lambda>	  s   ø€ ¨¯©¬r<   )Útorch._inductor.codegen.pallasr¤  r¹  r9  ru   r–   rú   r   )rÇ   rs   r  rø   r¤  r¹  r«   s   &ff @@@r:   ÚpallasÚAsyncCompile.pallasç  sT   ü€ ÷ 	Tä×ÑÐ1°;Ô?÷	Wð 	Wô Ó  AÔ%Ù“6ˆMà—[‘[ Ó&ˆFÜÔ 7Ó8Ð8r<   c               ó    € V ^8„  d   QhRRRR/# r¡  r8   )r9   s   "r:   r;   rÅ     s    € ÷ Eñ EØðEØ-0ñEr<   c           	     óê  a aaaa€ ^ RI Ho ^ RIHo \        P                  RV4       \        4        S P                  4       pV'       dM   \        4       pS P                  4       P                  \        SVSVV4      oR VVVV V3R llp\        VSR7      # \        P                  P                  P                   P#                  V4      w  rx\        P                  P                  P                   P%                  Wx4      p	S RS 2p
\'        Wš4      '       gH   \)        V	4       Uu. uF!  p\+        \-        W›4      4      '       g   K  VNK#  	  pp\/        RV
 R	V 24      hS! \-        Wš4      VR
7      # u upi )aÌ  
Compile NVIDIA Universal GEMM kernels.

Args:
    kernel_name: Name of the kernel to be defined
    source_code: Source code of the kernel, as a string
    precompile_metadata: Optional dict with shapes/dtypes for triggering
        real CuTe DSL compilation in the subprocess worker.

Note:
    NVIDIA Universal GEMM kernels are Python code that calls the cutlass.operators library.
    We use the PyCodeCache to write the source code to a file and load it.
)ÚNVUniversalGemmKernelWrapper)r¤  z NVIDIA Universal GEMM Kernel:
%sc               ó   € V ^8„  d   QhRR/# )r6   r7   rÃ  r8   )r9   s   "r:   r;   Ú4AsyncCompile.nv_universal_gemm.<locals>.__annotate__3  s   € ÷ ñ Ð <ñ r<   c                 óÌ   <€  SP                  4       w  rp\        P	                  RST4       SP                  SSSY4      #   \         d   pTP                  S4      ThR p?ii ; i)Nz:NV Universal GEMM kernel %s compiled in subprocess in %dusr¨  )	r¨   r›  r{   r  r¤  rÃ  rs   rÇ   rª  s	       €€€€€r:   r,  Ú2AsyncCompile.nv_universal_gemm.<locals>.get_result3  st   ø€ ð:Ø,;×,BÑ,BÓ,DÑ)�C˜zô —	‘	ØPØØôð
 ×0Ñ0Ø Ð.JÈCóð øô (ô :ØŸ+™+ kÓ2¸Ð9ûð:úr¬  r.  rS   z;Could not find NVIDIA Universal GEMM main kernel function 'r­  r–  )ÚBtorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_kernelrÃ  ÚFtorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_schedulingr¤  r9  ru   rY   ró   rF   r�   rú   r&   r   rm   r;  r<  r=  r¯  r˜  r°  r±  r²  r:  rÝ   )rÇ   rs   r  r³  rK  rL  r,  r¨   r›  rœ  r�  r^   r´  r¤  rÃ  rª  s   ff&&         @@@r:   Únv_universal_gemmÚAsyncCompile.nv_universal_gemm  sJ  ü€ õ 	
õ	
ô 	×ÑÐ@À+ÔNÜÔà×+Ñ+Ó-ˆçÜ7Ó9ˆIà"×/Ñ/Ó1×8Ñ8Ü2ØØØØØ#óˆO÷ô ô   
°?ÔCÐCäŸ™×1Ñ1×=Ñ=×CÑCÀKÓP‰IˆCÜ—/‘/×+Ñ+×7Ñ7×HÑHÈÓSˆCà +˜}¨A¨k¨]Ð;ˆNÜ˜3×/Ò/Ü.1°#¬hÓW©h d¼(Ä7È3ÓCU×:VŸT˜T©h�	ÐWÜ"ðØ&Ð'Ð'?À	¸{ðLóð ñ
 0Ü˜Ó,¸$ôð ùò Xs   Ä#E0ÅE0c                ó|   € ^ RI Hp \        4       pV P                  4       P	                  VVVVVVVVVVV	V
VVV4      # )a3  Submit NVGEMM kernel precompilation to the subprocess pool.

Compiles the CuTeDSL kernel artifact in a subprocess worker so the
thread-unsafe kernel.compile() call is process-isolated. The compiled
artifact is saved to the disk cache; the main process loads it on the
next cache lookup during benchmarking.
)Ú$_worker_nvgemm_autotuning_precompile)rÈ  rÍ  rF   r�   rú   )rÇ   rs   Úvariant_nameÚaccumulator_typeÚinput_tensor_metaÚoutput_tensor_metaÚcuda_ctxÚscale_type_aÚscale_type_bÚswizzle_type_aÚswizzle_type_bÚhas_bias_epilogueÚswap_abÚmetadatarÍ  rL  s   &&&&&&&&&&&&&&  r:   Únvgemm_precompileÚAsyncCompile.nvgemm_precompileR  sZ   € õ.	
ô 4Ó5ˆ	à× Ñ Ó"×)Ñ)Ø0ØØØØØØØØØØØØØØó
ð 	
r<   c               ó(   € V ^8„  d   QhRRRRRRRR/# )r6   rs   rt   ÚsourceÚheadersrm  r7   rW   r8   )r9   s   "r:   r;   rÅ   �  s.   € ÷ Cñ C ð C¨cð C¸Ið CÈ$ñ Cr<   c                ój   € V P                   f   . V n         V P                   P                  WV34       R# )zVRegister a Metal kernel body; wait() compiles all registered kernels into one library.N)rÃ   rx  )rÇ   rs   rÝ  rÞ  s   &&&&r:   ÚmetalÚAsyncCompile.metal�  s/   € à×ÑÒ&Ø"$ˆDÔØ×Ñ×"Ñ" K¸Ð#AÖBr<   c               ó    € V ^8„  d   QhRRRR/# ©r6   Úscoperv   r7   rW   r8   )r9   s   "r:   r;   rÅ   ‡  s   € ÷ ñ ˜.ð ¨Tñ r<   c                	óh  € \        4       ^8”  d3   \        RRRRRR7      ;_uu_ 4        V P                  V4       RRR4       V P                  '       dB   ^ RIHp VP                  V! V P                  4      4       V P                  P                  4        \        4        R#   + '       g   i     Lo; i)rÌ   zasync_compile.waitTr/  r0  r1  N)Úcompile_mps_shaders)	r–   r   Ú_wait_futuresrÃ   Ú%torch._inductor.runtime.runtime_utilsræ  ÚupdaterŽ   rq   )rÇ   rä  ræ  s   && r:   ÚwaitÚAsyncCompile.wait‡  s‰   € ÜÓ  1Ô$ÜØ$Ø&*Ø)AØ $Ø*:÷ö ð ×"Ñ" 5Ô)÷ð ××ÐÝQà�L‰LÑ,¨T×-@Ñ-@ÓAÔBØ×Ñ×%Ñ%Ô'äŽ÷÷ ús   ¨B!Â!B1	c               ó    € V ^8„  d   QhRRRR/# rã  r8   )r9   s   "r:   r;   rÅ   š  s   € ÷ &ñ & >ð &°dñ &r<   c           	     	ó’  € VP                  4        UUu/ uF%  w  r#\        V\        \        34      '       g   K#  W#bK'  	  ppp\	        \        V4      R \        P                  ^ R7      p\        P                  ;'       g    RpVP                  4        Fj  w  r'\        P                  '       d(   \        V\        4      '       g   VP                  V4        VP                  VR7      pW�V&   TP!                  ^4       Kl  	  R# u uppi   \         d   p	\        RT: RT R24      T	hRp	?	i\         d   p	\        R4      T	hRp	?	ii ; i)zInductor Compilation)ÚtotalÚdescÚdisableÚdelayNrñ   z#Inductor compile-worker future for z did not complete within zEs. Override with TORCHINDUCTOR_COMPILE_WORKER_WAIT_TIMEOUT=<seconds>.zÌA compilation subprocess exited unexpectedly. This is likely due to a crash. To facilitate debugging, you can re-run with TORCHINDUCTOR_COMPILE_THREADS=1 to cause compilation to occur in the main process.)rl   r  r   r   r,   rD  r   Údisable_progressÚcompile_worker_wait_timeoutÚverbose_progressr+   Úset_postfix_strrô   r  rÝ   r   ré  )
rÇ   rä  r¨   ÚvalueÚkernelsÚpbarÚwait_timeoutrô   rz   r  s
   &&        r:   rç  ÚAsyncCompile._wait_futuresš  sC  € ð $Ÿk™kœmô
á+‘
�Ü˜%¤&¬/Ð!:×;ô ˆCŠJÙ+ð 	ñ 
ô
 Ü�g“,Ø'Ü×+Ñ+Øô	
ˆô ×9Ñ9×AÐA¸TˆØ"Ÿ=™=ž?‰KˆCÜ×&×&Ð&¬z¸$Ä	×/JÒ/JØ×$Ñ$ SÔ)ðØŸ™¨|˜Ó<�Ø#�c‘
ð" �K‰K˜ŽNó- +ùó
øô* 'ô ô #Ø9¸#¹ð A'Ø'3 nð 5KðKóð ð	ûô
 %ô Ü"ðIóð
 ðûðús4   ” D¹DÃDÄEÄD'Ä'EÄ4EÄ5EÅE)rÃ   )éx   )r‚  )FrO   )NNNNFFN))rº   r»   r¼   r½   r¾   r‰   r¿   rÃ   rÈ   rÀ   r8  Ú	lru_cacher‹   rÒ   r�   Úclassmethodrí   rõ   rú   ró   r  r  rý   r[  r`  ri  rq  r~  r‚  r…  rŠ  r“  rž  rµ  r¿  rÊ  rÚ  rà  rê  rç  rÁ   r8   r<   r:   rˆ   rˆ     sk  ‡ ñð )-€MÐ%Ó,Ø>B€NÐ;ÓBõð Ø×Ò˜Óô9ó ó ð9ð ñó ðð Ø×Ò˜Óô-ó ó ð-ð^ ôó ðð ö6ó ð6ð
 ô'ó ð'ð
 ñ(ó ð(ð" öó ððB ô	ó ð	÷Y1õv0õ8õ
=õ,ô	!ôSôRô!õ$,òK÷>Xõ@"9÷HEôN-
õ^Cõ÷&&ñ &r<   rˆ   c               ó   € V ^8„  d   QhRR/# rV   r8   )r9   s   "r:   r;   r;   Ã  s   € ÷ ñ ˜ñ r<   c                 óD  € \         P                  P                  R R4      R8X  gQ   \         P                  P                  RR4      R8w  g,   \        4       '       d   \        P
                  ! 4       '       d   R# \        P                  4        \        P                  4        R# )ÚTORCH_TNT_IN_USErØ   r  rÙ   N)	rA   rB   rC   r.   r   Ú	is_fbcoderˆ   rí   r  r8   r<   r:   Úmaybe_warm_poolr  Ã  si   € ä
�
‰
�‰Ð)¨3Ó/°3Ô6Ü�:‰:�>‰>Ð+¨SÓ1°SÔ8ä!×#Ò#ô ×Ò×Òáä×ÑÔô ×ÑÖr<   )oÚ__conditional_annotations__Ú
__future__r   Úatexitr8  ra   Úloggingrœ   rA   Úrerä   Úconcurrent.futuresr   r   r   r  Úconcurrent.futures.processr   r   r	   r
   Útypingr   r   rm   Útorch._dynamo.device_interfacer   Útorch._dynamo.utilsr   r   r   r   Útorch._inductorr   Útorch._inductor.codecacher   r   r   r   r   r   r   r   r   r   r   r   Ú+torch._inductor.compile_worker.subproc_poolr   r    r!   Ú3torch._inductor.compile_worker.tracked_process_poolr"   Ú$torch._inductor.compile_worker.utilsr#   Ú%torch._inductor.runtime.compile_tasksr$   r%   r&   r'   Útorch._inductor.utilsr(   Útorch._inductor.virtualizedr)   Útorch._utils_internalr*   Ú	torch.hubr+   r,   Útorch.utils._ordered_setr-   Útorch.utils._tritonr.   Úcollections.abcr/   Útorch._inductor.runtime.hintsr0   rI  r1   ri   r2   r¿   rn   ÚgetArtifactLoggerrº   r9  Ú	getLoggerrÞ   r4   rv  rE  rF   rM   rK   rY   rq   rx   r‚   ÚplatformÚ_IS_WINDOWSr†   rŒ   rŠ   Úregister_at_forkÚAttributeErrorr–   rŸ   r¡   rˆ   r  Úregister)r  s   @r:   Ú<module>r"     sÅ  øðç "Ð "ã Û Û Û Û Û 	Û 	Û 
÷ñ õ
 9Ý ß ß %ã Ý K÷ó õ #÷÷ ÷ ó ÷ñ õ
õ L÷ó õ 7Ý )Ý 3ß %Ý /Ý 2÷ Ý(å8ÝJð Ð Ø€€\Ó à—.‘.×2Ñ2°8¸]ÓK€à×Ò˜Ó!€à;?Ð Ð8Ó ?à—:’:ØóÐ õ
4òò <õ$õ&õ*3õð* �l‰l˜gÑ%€à×Ò˜Ó!€ð �wÖÓ!€	õò,ð	Ø×Ò zÕ2õ
"õð ÷22ð 22ó ð22÷jg
ñ g
õTð0 ‡‚Ð(Ö )øðc ô 	Ùð	ús   ÆG* Ç*G4Ç3G4