ó
    Eñi¸n  ã                  ó˜  • % S SK Jr  S SKrS SKrS SKrS SKrS SKrS SKrS SKrS SK	r	S SK
JrJr  S SKJr  S SKJr  S SKJrJr  S SKJrJrJr  S SKrS SKJr  S S	KJrJrJrJr  S S
KJr  S SK J!r!J"r"J#r#J$r$J%r%J&r&J'r'J(r(J)r)J*r*J+r+  S SK,J-r-J.r.J/r/  S SK0J1r1  S SK2J3r3  S SK4J5r5J6r6  S SK7J8r8  S SK9J:r:  S SK;J<r<  S SK=J>r>J?r?  S SK@JArA  S SKBJCrC  \(       a  S SKDJErE  S SKFJGrG  S SKHJIrI  SqJSqKS\LS'   \Rš                  R�                  \OS5      rP\R¢                  " \O5      rRSqSS\LS'   \R¨                  " S5      rUS  rVS! rWS/S" jrXS/S# jrYS0S$ jrZ\	R¶                  S%:H  r\\R¢                  " \O5      rR\A\-   " 5       r]S/S& jr^S' r_ \RÀ                  " \_S(9  S1S) jrb\8 " S* S+5      5       rc " S, S-5      rdS/S. jre\RÌ                  " \^5        g! \a a     N>f = f)2é    )ÚannotationsN)ÚFutureÚThreadPoolExecutor)ÚBrokenProcessPool)Úpartial)ÚtimeÚtime_ns)ÚAnyÚOptionalÚ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)ÚAnyPoolÚSubprocExceptionÚSubprocPool)ÚTrackedProcessPoolExecutor)Ú_async_compile_initializer)Ú_set_triton_ptxas_pathÚ_worker_compile_triton)Úclear_on_fresh_cache)ÚV)Úlog_triton_builds)Ú	_FaketqdmÚtqdm)Ú
OrderedSet)Úhas_triton_package)ÚCallable)Ú
HalideMeta)ÚCachingAutotunerg        zOptional[float]Ú_t0Úkernel_codez#Optional[dict[str, dict[str, Any]]]Ú_triton_kernel_metricszsize_hints=(\{.*?\})c                 óF   • [        5         SSKJn Jn  U (       a  U" 5         gg)z?
Setup that must be done prior to forking with a process pool.
r   ©Ú
HAS_TRITONÚ
triton_keyN)Úcaching_device_propertiesÚ%torch._inductor.runtime.triton_compatr4   r5   r3   s     ÚZ/home/mande/repo/quber/.venv/lib/python3.13/site-packages/torch/_inductor/async_compile.pyÚpre_fork_setupr9   R   s   € ô Ô÷ MæÙ�ð ó    c                 óŒ   • [        5        H6  u  pUR                  5       (       d  M  UR                  R                  5         M8     g ©N)r   Úis_availableÚWorkerÚget_device_properties)Ú_Údevice_interfaces     r8   r6   r6   b   s4   € Ü?ÖAÑˆØ×(Ñ(×*Ó*Ø×#Ñ#×9Ñ9Ö;ò  Br:   c                 ó:   • [         c
  [        5       q [        c  0 qg g r<   )r/   r   r1   © r:   r8   Ú_compile_startrD   h   s   € ä
�{Ü‹fˆÜÑ%Ø!#Ñð &r:   c                 óü   ^• [         b  [        5       n [        U [         -
  -  qS q [        (       aM  [	        [        [        R                  5       5      5      m[        R                  R                  SS U4S jS9  S qg g )NÚartifactc                 ó   • SSS.$ )NÚtriton_kernel_infoÚjson)ÚnameÚencodingrC   rC   r:   r8   Ú<lambda>Ú_compile_end.<locals>.<lambda>|   s   € Ø,Ø"ò!r:   c                 ó0   >• [         R                  " T 5      $ r<   )rI   Údumps)Úsorted_infos   €r8   rL   rM   €   s   ø€ œtŸzšz¨+Ô6r:   )Úmetadata_fnÚ
payload_fn)
r/   r   Ú_cumulative_compile_timer1   ÚdictÚsortedÚitemsÚtorchÚ_loggingÚtrace_structured)Út1rP   s    @r8   Ú_compile_endr[   p   ss   ø€ ä
�Ü‹VˆÜ  B¬¡HÑ,Ð ØˆçÒäœ6Ô"8×">Ñ">Ó"@ÓAÓBˆÜ�‰×'Ñ'Øñô 7ð 	(ñ 	
ð "&Ñð r:   c                ó&   • [         b
  U[         U '   g g r<   )r1   )Úkernel_nameÚinfos     r8   Ú_add_triton_kernel_infor_   …   s   € ô Ñ)Ø.2Ô˜{Ò+ð *r:   Úwin32c                 óh   • [          H  n U R                  5         M     S[        l        [	        5         g)z/Shut down all outstanding compile-worker pools.N)Ú	_pool_setÚshutdownÚAsyncCompileÚ_ready_futureÚ
after_fork)Úpools    r8   Úshutdown_compile_workersrh   ”   s"   € çˆØ�‰Žñ à!%„LÔÜ…Lr:   c                 óh   • [         R                  5         [        R                  R	                  5         g)z7Reset pools to initial state without shutting them downN)rb   Úclearrd   Úprocess_poolÚcache_clearrC   r:   r8   rf   rf   œ   s   € ä‡O�OÔÜ×Ñ×)Ñ)Õ+r:   )Úafter_in_childc                 ó€   • [         R                  c  [         R                  " 5       [         l        [         R                  $ )zq
Temporary for internal rollout. Assign config.compile_threads lazily and return it.
TODO: remove after rollout.
)r   Úcompile_threadsÚdecide_compile_threadsrC   r:   r8   Úget_compile_threadsrq   ¨   s-   € ô
 ×ÑÑ%Ü!'×!>Ò!>Ó!@ŒÔÜ×!Ñ!Ð!r:   c                  óˆ   • \ rS rSr% Sr0 rS\S'   \SS j5       r\SS j5       r	\SS j5       r
\S 5       r\SS	 j5       rS
rg)ÚCompiledTritonKernelsé²   a  
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                ó&   • [        U [        5       S9$ )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   )Ú
kernel_srcs    r8   ÚkeyÚCompiledTritonKernels.key¿   s   € ô ˜¬9«;Ñ7Ð7r:   c                óT   • [         R                  U 5      nU[         R                  U'   g)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©rs   ry   ru   )rx   Úfuturery   s      r8   ÚsaveÚCompiledTritonKernels.saveÉ   s%   € ô $×'Ñ'¨
Ó3ˆØ,2Ô×$Ñ$ SÒ)r:   c                ól   • [         R                  U 5      n[         R                  R                  US 5      $ r<   )rs   ry   ru   Úget©rx   ry   s     r8   r�   ÚCompiledTritonKernels.getÖ   s+   € ä#×'Ñ'¨
Ó3ˆÜ$×+Ñ+×/Ñ/°°TÓ:Ð:r:   c                 ó   • 0 [         l        g r<   )rs   ru   rC   r:   r8   rl   Ú!CompiledTritonKernels.cache_clearÛ   s   € à')ÔÕ$r:   c                óz   • [         R                  U 5      nU[         R                  ;   a  [         R                  U	 g g r<   r|   r‚   s     r8   Úremove_futureÚ#CompiledTritonKernels.remove_futureß   s6   € ä#×'Ñ'¨
Ó3ˆð Ô'×.Ñ.Ó.Ü%×,Ñ,¨SÑ1ð /r:   rC   N)rx   Ústr)rx   r‰   r}   r   )rx   r‰   ÚreturnzOptional[CodeCacheFuture])rx   r‰   rŠ   ÚNone)Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__ru   Ú__annotations__Ústaticmethodry   r~   r�   rl   r‡   Ú__static_attributes__rC   r:   r8   rs   rs   ²   sv   ‡ ñð *,€FÐ&Ó+àó8ó ð8ð ó
3ó ð
3ð ó;ó ð;ð ñ*ó ð*ð ó2ó ó2r:   rs   c                  óª  • \ rS rSr% SrSrS\S'   SS jr\\	R                  " S5      SS j5       5       r\S	 5       r\\	R                  " S5      S S
 j5       5       r\SS j5       r\S!SS jj5       r\S"S j5       r\S 5       r\SS j5       rS#S$S jjrS%S jrS%S jrS&S jrS'S jrS(S jr S(S jrS)S jrS*S jrS*S jrS*S jrS+S jrS+S jr Sr!g),rd   éè   zS
Utilities to compile in thread pools or subprocess pools (in the case of Triton).
NzOptional[Future[Any]]re   c                ó   • g r<   rC   )Úselfs    r8   Ú__init__ÚAsyncCompile.__init__ï   s   € Ør:   é   c                 óH   • [        5       S:”  d   e[        [        5       5      $ ©Nrš   )rq   r   rC   r:   r8   rg   ÚAsyncCompile.poolò   s$   € ô #Ó$ qÓ(Ð(Ð(Ü!Ô"5Ó"7Ó8Ð8r:   c                 ó   • g)z>No-op function to help mark when the subprocess pool is ready.ÚreadyrC   rC   r:   r8   Ú
_get_readyÚAsyncCompile._get_readyø   s   € ð r:   c            	     óª  • [        5       S:”  d   eS [        l        [        R	                  S[
        R                  [        5       5        [
        R                  S:X  a!  [        [        5       [
        R                  S9n OÀ[
        R                  S:X  a  S[        R                  S'   [        5         [        R                  " [
        R                  5      n[        [        5       U[        [         [        R"                  " 5       5      S9n [        R$                  R'                  S U R(                  [*        R,                  S	9  [.        R1                  U 5        U $ )
Nrš   z"Creating '%s' pool with %d workersÚ
subprocess)ÚquiesceÚspawnÚ0ÚTORCH_WARM_POOL)Ú
mp_contextÚinitializer)Úexitpriority)rq   rd   re   Úlogr^   r   Úworker_start_methodr    Úquiesce_async_compile_poolÚosÚenvironr9   ÚmultiprocessingÚget_contextr!   r   r"   ÚgetpidÚutilÚFinalizerc   ÚsysÚmaxsizerb   Úadd)rg   Úctxs     r8   rk   ÚAsyncCompile.process_poolý   sø   € ô #Ó$ qÓ(Ð(Ð(Ø%)ŒÔ"Ü�‰Ø0Ü×&Ñ&ÜÓ!ô	
ô ×%Ñ%¨Ó5äÜ#Ó%¬v×/PÑ/Pñ‰Dô ×)Ñ)¨WÓ4à03”—
‘
Ð,Ñ-ÜÔÜ!×-Ò-¬f×.HÑ.HÓIˆCÜ-Ü#Ó%ØÜ#Ô$>ÄÇ	Â	ÃÓLñˆDô × Ñ ×)Ñ)¨$°·±ÌCÏKÉKÐ)ÑXä�‰�dÔØˆr:   c                ój   • [        5       S::  a  g [        5         U R                  5         [        5         g rœ   )rq   rD   rk   r[   ©Úclss    r8   Ú	warm_poolÚAsyncCompile.warm_pool"  s)   € äÓ  AÓ%ØÜÔð 	×ÑÔÜ�r:   c                ór   • U R                  5         U R                  b  U R                  R                  US9  g g )N)Útimeout)Úuse_process_poolre   Úresult)r¼   rÀ   s     r8   Úwait_pool_readyÚAsyncCompile.wait_pool_ready-  s6   € à×ÑÔØ×ÑÑ(Ø×Ñ×$Ñ$¨WÐ$Ò5ð )r:   c                ój   • [        5       S::  a  U" 5       $ U R                  5       R                  U5      $ rœ   )rq   rg   Úsubmit)r¼   Útasks     r8   rÆ   ÚAsyncCompile.submit3  s+   € äÓ  AÓ%Ù“6ˆMØ�x‰x‹z× Ñ  Ó&Ð&r:   c                ó  • [        5       S::  a  g[        R                  R                  (       a  gU R                  (       d.  U R                  5       R                  U R                  5      U l        U R                  R                  5       $ )Nrš   F)	rq   r   ÚtritonÚproton_profilingre   rk   rÆ   r    Údoner»   s    r8   rÁ   ÚAsyncCompile.use_process_pool9  sb   € äÓ  AÓ%Øô
 �=‰=×)×)Øð
 × × Ø #× 0Ñ 0Ó 2× 9Ñ 9¸#¿.¹.Ó IˆCÔØ× Ñ ×%Ñ%Ó'Ð'r:   c                óœ   • U R                  5       (       d  gU R                  5       n[        U[        5      (       a  UR	                  5         gg)zY
If using a SubprocPool, signal the sidecar process to start up its
ProcessPoolExecutor.
N)rÁ   rk   Ú
isinstancer    Úwakeup)r¼   rg   s     r8   rÐ   ÚAsyncCompile.wakeupK  s?   € ð ×#Ñ#×%Ñ%ØØ×ÑÓ!ˆÜ�dœK×(Ñ(Ø�K‰K�Mð )r:   c                óx  ^^^^^^^• [         R                  " [        TT5      mU4S jm[        S   S==   S-  ss'   [        R                  ST5        [        5         [        R                  R                  SS5      S:X  a=  [        [        R                  R                  R                  R                  T5      T5      $ U R!                  5       n[#        S	U5        [        R$                  R&                  R)                  5       m[        [*        R,                  S
S5      m[.        R                  T5      =nb^  [        S   S==   S-  ss'   [1        U[2        5      (       a  [.        R5                  T5        TUl        U(       a  U$ UR9                  5       $ U(       Ga˜  SS/nU Vs0 s H,  ow[        R                  ;   d  M  U[        R                  U   _M.     nnS[        R                  R:                  R<                  0n	[?        [        R                  R:                  R@                  5      S:”  a¯  [B        RE                  T5      n
U
(       a  U
RG                  S5      nO[I        S5      nTRK                  S5      S   nSSK&J'n  U" X¼5      nU[        R                  R:                  R@                  ;   a,  U[        R                  R:                  R@                  U   0U	S'   U RQ                  5       RS                  [T        TUU	5      mS!UUUUUU4S jjn[W        UTS9n[.        RY                  TU5        U$ [[        SSSSSS9   Sn []        5       n[_        5         T" 5       nURa                  TT5        URc                  S[.        Re                  T5      S9  []        5       U-
  S-  n[g        5       Ri                  STU5        URj                  =(       d    0 nUUS'   [m        TU5        U[o        US 9  sSSS5        $ s  snf ! [p         a  n[I        U5      ne SnAff = f! [o        US 9  f = f! , (       d  f       g= f)"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                 ó^   >• [        S5         T " 5       sS S S 5        $ ! , (       d  f       g = f)NÚreload_kernel_in_parent)r   )Úload_kernels   €r8   rÔ   Ú4AsyncCompile.triton.<locals>.reload_kernel_in_parentq  s   ø€ äÐ7Õ8Ù"“}÷ 9×8×8ús   �ž
,ÚinductorÚasync_compile_cache_missrš   zTriton Kernel:
%sÚTRITON_INTERPRETr¦   Ú1Úparallel_compile_post_warmupÚis_backwardFNÚasync_compile_cache_hitÚTORCHINDUCTOR_CACHE_DIRÚTRITON_CACHE_DIRÚuse_static_triton_launcherr   z@triton.jit
)Ú%generate_lookup_hash_from_source_codeÚautotune_lookup_tablec                 óª  >•  T	R                  5       u  pU R                  TT5        [        R                  T5        U R                  S S9  U R                  ST[        R                  T5      S9  U R                  =(       d    0 nXS'   [        TU5        [        5       R                  STU5        U $ ! [         a  nUR                  T5      UeS nAff = f)N)Ú
old_valuesF)Úwarm_cache_onlyÚreload_kernelÚstatic_triton_bundle_keyÚcompile_time_usÚtriton_kernel_compile_times_us)rÂ   r   Ú	with_nameÚset_compile_infors   r‡   Úrestore_after_unpickleÚ
precompilery   Úautotune_cache_infor_   r   Ú	add_top_n)
ÚkernelÚ
elapsed_usÚer^   Ú
compile_idrÜ   r]   rÔ   Úsource_coderÇ   s
       €€€€€€r8   Ú
get_resultÚ'AsyncCompile.triton.<locals>.get_result¸  sÕ   ø€ ð:Ø)-¯©«Ñ&�Fð ×'Ñ'¨
°KÔ@Ü%×3Ñ3°KÔ@à×-Ñ-¸Ð-Ñ>à×!Ñ!Ø$)Ø"9Ü-B×-FÑ-FÀ{Ó-Sð "ñ ð
 ×1Ñ1×7°R�Ø*4Ð&Ñ'Ü'¨°TÔ:Ü#Ó%×/Ñ/Ø4°kÀ:ôð �øô+ (ó :ØŸ+™+ kÓ2¸Ð9ûð:ús   ƒB1 Â1
CÂ;CÃC©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è  ré   rè   )Úfail)rŠ   r.   )9Ú	functoolsr   r   r   Úkernel_code_logr^   rD   r®   r¯   r�   ÚgetattrrW   Ú	_inductorÚ	codecacheÚPyCodeCacheÚloadrÁ   r   Ú_guardsÚCompileContextÚcurrent_compile_idr&   Úgraphrs   rÏ   r   r‡   Úreload_kernel_from_srcrÂ   r   rà   Úlenrâ   Úsize_hints_regexÚsearchÚgroupr‰   ÚsplitÚ)torch._inductor.runtime.triton_heuristicsrá   rk   rÆ   r$   r   r~   r   r	   r#   rë   rí   ry   r   rï   rî   r_   r'   Ú	Exception)r—   r]   rô   Ú
device_strÚis_parallelr}   Úenv_varsÚvÚ	extra_envÚextra_configÚmÚsize_hints_strÚ
triton_srcrá   Úfn_hashrõ   rÿ   Ústart_nsrð   rñ   r^   rò   ró   rÜ   rÕ   rÔ   rÇ   s    ``                   @@@@@r8   rÊ   ÚAsyncCompile.tritonW  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�Ô-ÞØ�à—}‘}“Ð&÷ ð 2Ð3EÐFˆHÙ3;ÓO²8¨aÄBÇJÁJ¹Ó)˜œBŸJ™J q™MÒ)±8ˆIÐOà,¬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÷ô ô4 " *°TÑ:ˆFÜ!×&Ñ& {°FÔ;ØˆMäØ*Ø&*Ø)AØ $Ø*:óð �ð1Ü&›y�HÜ*Ô,Ù(›]�FØ×+Ñ+¨J¸ÔDØ×%Ñ%Ø(-Ü1F×1JÑ1JÈ;Ó1Wð &ñ ô #*£)¨hÑ"6¸4Ñ!?�JÜ'Ó)×3Ñ3Ø8¸+Àzôð "×5Ñ5×;¸�DØ.8�DÐ*Ñ+Ü+¨K¸Ô>Ø!ô
 &¨4Ò0÷;ñ ùò Pøôp !ó Ü˜q›6�DØûðûô &¨4Ó0ú÷;õ úsI   Æ$O:Æ?O:ÍP+ÍB"O?Ï'	P+Ï?
PÐ	PÐPÐPÐP(Ð(P+Ð+
P9c                ó   • SSK Jn  U" U0 UD6$ )Nr   )ÚMultiKernelCall)Ú$torch._inductor.codegen.multi_kernelr   )r—   ÚargsÚkwargsr   s       r8   Úmulti_kernelÚAsyncCompile.multi_kernelõ  s   € ÝHñ  Ð/¨Ñ/Ð/r:   c                ó   • SSK Jn  U" U0 UD6$ )Nr   )ÚSizeHintMultiKernelCall)r!  r'  )r—   r"  r#  r'  s       r8   Úsize_hint_multi_kernelÚ#AsyncCompile.size_hint_multi_kernelû  s   € ÝPá&¨Ð7°Ñ7Ð7r:   c                óæ   ^• [         R                  SU5        [        5       S::  a   [        R                  " U5      R
                  $ [        R                  " XR                  S9m[        U4S j5      $ )NzCPP Kernel:
%srš   ©Ú	submit_fnc                 ó&   >• T " 5       R                   $ r<   )rð   )rõ   s   €r8   rL   Ú"AsyncCompile.cpp.<locals>.<lambda>  s   ø€ ©
«×(;Ò(;r:   )	r  r^   rq   r   r  rð   Ú
load_asyncrÆ   r   )r—   rô   rõ   s     @r8   ÚcppÚAsyncCompile.cpp   sV   ø€ Ü×ÑÐ.°Ô<ÜÓ  AÓ%Ü×$Ò$ [Ó1×8Ñ8Ð8ä%×0Ò0°ÏÉÑTˆJÜÔ ;Ó<Ð<r:   c                óÊ   • [         R                  SU5        [        5       S::  a  [        R                  " X5      $ [        R
                  " XU R                  S9n[        U5      $ )NzCPP+Bindings Kernel:
%srš   r+  )r  r^   rq   r   Úload_pybindingÚload_pybinding_asyncrÆ   r   )r—   Úargtypesrô   rõ   s       r8   Úcpp_pybindingÚAsyncCompile.cpp_pybinding  sU   € Ü×ÑÐ7¸ÔEÜÓ  AÓ%Ü-×<Ò<¸XÓSÐSä3×HÒHØ°·±ñˆJô   
Ó+Ð+r:   c                óf   ^^^• [         R                  ST5        UUU4S jnU R                  U5      $ )NzCUDA Kernel:
%sc                 ó´   >• T(       a7  [         R                  " TS5      tp[         R                  R                  U 5        [         R                  " TT5      S   $ )NÚor   )r   ÚcompileÚaot_kernels_oÚappendr  ©Úoutput_pathr@   Úaot_compileÚdst_file_extrô   s     €€€r8   rÇ   ÚAsyncCompile.cuda.<locals>.task  sH   ø€ Þô #0×"7Ò"7¸ÀSÓ"I��Ü×+Ñ+×2Ñ2°;Ô?Ü ×%Ò% k°<Ó@ÀÑCÐCr:   ©r  r^   rÆ   ©r—   rô   rA  r@  rÇ   s    ``` r8   ÚcudaÚAsyncCompile.cuda  s+   ú€ Ü×ÑÐ/°Ô=÷	Dð �{‰{˜4Ó Ð r:   c                óf   ^^^• [         R                  ST5        UUU4S jnU R                  U5      $ )NzROCm Kernel:
%sc                 ó  >• T(       a5  [         R                  " TSS9tp[         R                  R                  U 5        [        R
                  R                  (       a  [         R                  " TSS9n[         R                  " TT5      S   $ )Nr:  )rA  Úexer   )r   r;  r<  r=  r   ÚrocmÚgenerate_test_runnerr  r>  s     €€€r8   rÇ   ÚAsyncCompile.rocm.<locals>.task'  sf   ø€ ÞÜ"/×"7Ò"7¸ÐRUÑ"V��Ü×+Ñ+×2Ñ2°;Ô?Ü�{‰{×/×/Ü!×)Ò)¨+ÀEÑJ�Ü ×%Ò% k°<Ó@ÀÑCÐCr:   rC  rD  s    ``` r8   rJ  ÚAsyncCompile.rocm  s-   ú€ ô 	×ÑÐ/°Ô=÷	Dð �{‰{˜4Ó Ð r:   c                óÊ   • [         R                  SX5        [        5       S::  a  [        R                  " X5      $ [        R
                  " XU R                  S9n[        U5      $ )NzHalide Kernel:
%r
%srš   r+  )r  r^   rq   r   Úgenerate_halideÚgenerate_halide_asyncrÆ   r   )r—   Úmetarô   rõ   s       r8   ÚhalideÚAsyncCompile.halide1  sU   € Ü×ÑÐ5°tÔIÜÓ  AÓ%Ü"×2Ò2°4ÓEÐEä(×>Ò>Ø¨T¯[©[ñˆJô   
Ó+Ð+r:   c                óÄ   ^^^^^• SSK JmJm  [        R	                  ST5        UUUU4S jn[        5       S::  a  U" 5       $ U R                  U5      m[        U4S j5      $ )aC  
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

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.
r   )ÚCuteDSLKernelWrapperÚMAIN_SUFFIXzCuteDSL Kernel:
%sc            	     óº  >• [         R                  R                  R                  R	                  T	5      u  p[         R                  R                  R                  R                  X5      nT ST 3n[        X#5      (       dG  [        U5       Vs/ s H   n[        [        X$5      5      (       d  M  UPM"     nn[        SU SU 35      eT" [        X#5      US9$ s  snf )Nr@   z-Could not find CuteDSL main kernel function 'ú'. Available callables: ©Úkernel_path©rW   r  r  r  ÚwriteÚload_by_key_pathÚhasattrÚdirÚcallabler  ÚRuntimeError)
ry   ÚpathÚmodÚmain_func_namerJ   Ú	availablerU  rV  r]   rô   s
         €€€€r8   rÇ   Ú"AsyncCompile.cutedsl.<locals>.taskN  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©h�	ÐWÜ"ØCÀNÐCSÐSkÐluÐkvÐwóð ñ (¬°Ó(DÐRVÑWÐWùò Xó   ÂCÂ/Crš   c                 ó$   >• T R                  5       $ r<   ©rÂ   r÷   s   €r8   rL   Ú&AsyncCompile.cutedsl.<locals>.<lambda>`  ó   ø€ ¨¯©¬r:   )Ú.torch._inductor.codegen.cutedsl.cutedsl_kernelrU  rV  r  r^   rq   rÆ   r   )r—   r]   rô   rÇ   rU  rV  r}   s    `` @@@r8   ÚcutedslÚAsyncCompile.cutedsl;  sU   ü€ ÷	
ô
 	×ÑÐ2°KÔ@÷	Xð 	Xô Ó  AÓ%Ù“6ˆMà—[‘[ Ó&ˆFÜÔ 7Ó8Ð8r:   c                óÄ   ^^^^^• SSK JmJm  [        R	                  ST5        UUUU4S jn[        5       S::  a  U" 5       $ U R                  U5      m[        U4S j5      $ )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   )rV  ÚPallasKernelWrapperzPallas Kernel:
%sc            	     óº  >• [         R                  R                  R                  R	                  T	5      u  p[         R                  R                  R                  R                  X5      nT ST 3n[        X#5      (       dG  [        U5       Vs/ s H   n[        [        X$5      5      (       d  M  UPM"     nn[        SU SU 35      eT" [        X#5      US9$ s  snf )Nr@   z,Could not find Pallas main kernel function 'rX  rY  r[  )
ry   rb  rc  rd  rJ   re  rV  rp  r]   rô   s
         €€€€r8   rÇ   Ú!AsyncCompile.pallas.<locals>.taskr  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©h�	ÐWÜ"ØBÀ>ÐBRÐRjÐktÐjuÐvóð ñ '¤w¨sÓ'CÐQUÑVÐVùò Xrg  rš   c                 ó$   >• T R                  5       $ r<   ri  r÷   s   €r8   rL   Ú%AsyncCompile.pallas.<locals>.<lambda>„  rk  r:   )Útorch._inductor.codegen.pallasrV  rp  r  r^   rq   rÆ   r   )r—   r]   rô   rÇ   rV  rp  r}   s    `` @@@r8   ÚpallasÚAsyncCompile.pallasb  sT   ü€ ÷ 	Tä×ÑÐ1°;Ô?÷	Wð 	Wô Ó  AÓ%Ù“6ˆMà—[‘[ Ó&ˆFÜÔ 7Ó8Ð8r:   c                óÌ   ^^^^^• SSK Jm  SSKJm  [        R                  ST5        UUUU4S jn[        5       S::  a  U" 5       $ U R                  U5      m[        U4S j5      $ )aA  
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

Note:
    NVIDIA Universal GEMM kernels are Python code that calls the cutlass_api library.
    We use the PyCodeCache to write the source code to a file and load it.
r   )ÚNVUniversalGemmKernelWrapper)rV  z NVIDIA Universal GEMM Kernel:
%sc            	     óº  >• [         R                  R                  R                  R	                  T	5      u  p[         R                  R                  R                  R                  X5      nT ST 3n[        X#5      (       dG  [        U5       Vs/ s H   n[        [        X$5      5      (       d  M  UPM"     nn[        SU SU 35      eT" [        X#5      US9$ s  snf )Nr@   z;Could not find NVIDIA Universal GEMM main kernel function 'rX  rY  r[  )
ry   rb  rc  rd  rJ   re  rV  ry  r]   rô   s
         €€€€r8   rÇ   Ú,AsyncCompile.nv_universal_gemm.<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©h�	ÐWÜ"ðØ&Ð'Ð'?À	¸{ðLóð ñ
 0Ü˜Ó,¸$ñð ùò Xrg  rš   c                 ó$   >• T R                  5       $ r<   ri  r÷   s   €r8   rL   Ú0AsyncCompile.nv_universal_gemm.<locals>.<lambda>°  rk  r:   )	ÚBtorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_kernelry  ÚFtorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_schedulingrV  r  r^   rq   rÆ   r   )r—   r]   rô   rÇ   rV  ry  r}   s    `` @@@r8   Únv_universal_gemmÚAsyncCompile.nv_universal_gemm†  sX   ü€ õ	
õ	
ô 	×ÑÐ@À+ÔN÷	ð 	ô" Ó  AÓ%Ù“6ˆMà—[‘[ Ó&ˆFÜÔ 7Ó8Ð8r:   c                ó¤   • [        5       S:”  a'  [        SSSSSS9   U R                  U5        S S S 5        [        5         g ! , (       d  f       N= f)Nrš   zasync_compile.waitTrø   rù   rú   )rq   r   Ú_wait_futuresr[   )r—   Úscopes     r8   ÚwaitÚAsyncCompile.wait²  sK   € ÜÓ  1Ó$ÜØ$Ø&*Ø)AØ $Ø*:óð ×"Ñ" 5Ô)÷ô 	�÷õ ús   œAÁ
Ac           	     ó  • UR                  5        VVs0 s H$  u  p#[        U[        [        45      (       d  M"  X#_M&     nnn[	        [        U5      S[        R                  SS9nUR                  5        Hf  u  p&[        R                  (       a&  [        U[        5      (       d  UR                  U5         UR                  5       nXqU'   UR                  S5        Mh     g s  snnf ! [         a  n[        S5      UeS nAff = f)NzInductor Compilationr   )ÚtotalÚdescÚdisableÚdelayzÌ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.rš   )rV   rÏ   r   r   r)   r  r   Údisable_progressÚverbose_progressr(   Úset_postfix_strrÂ   r   ra  Úupdate)	r—   r„  ry   ÚvalueÚkernelsÚpbarrÂ   rð   rò   s	            r8   rƒ  ÚAsyncCompile._wait_futures¿  së   € ð $Ÿk™kœmô
â+‘
�Ü˜%¤&¬/Ð!:×;ó ˆCŠJÙ+ð 	ñ 
ô
 Ü�g“,Ø'Ü×+Ñ+Øñ	
ˆð #Ÿ=™=ž?‰KˆCÜ×&×&¬z¸$Ä	×/JÑ/JØ×$Ñ$ SÔ)ð	ØŸ™›�Ø#�c‘
ð �K‰K˜ŽNò +ùó
øô" %ó Ü"ðIóð
 ðûðús"   ”!C ¹C Â6C&Ã&
DÃ0C<Ã<DrC   ©rŠ   r‹   )rŠ   r   )rŠ   r   )éx   )rÇ   zCallable[..., Any]rŠ   r
   )rE  )r]   r‰   rô   r‰   r  r‰   )rŠ   r
   )rô   r‰   )r5  z	list[str]rô   r‰   )F)rQ  r-   rô   r‰   )r]   r‰   rô   r‰   )r„  údict[str, Any]rŠ   r‹   )"rŒ   r�   rŽ   r�   r�   re   r‘   r˜   r’   r   Ú	lru_cacherg   r    rk   Úclassmethodr½   rÃ   rÆ   rÁ   rÐ   rÊ   r$  r(  r0  r6  rE  rJ  rR  rm  rv  r€  r…  rƒ  r“   rC   r:   r8   rd   rd   è   s+  ‡ ñð ,0€MÐ(Ó/ôð Ø×Ò˜Óó9ó ó ð9ð ñó ðð Ø×Ò˜Óó!ó ó ð!ðF óó ðð õ6ó ð6ð
 ó'ó ð'ð
 ñ(ó ð(ð" ó	ó ð	ö\1ô|0ô8ô
=ô,ô!ð" ô	!ô$,ô%9ôN"9ôH*9ôX÷r:   rd   c                 ó8  • [         R                  R                  SS5      S:X  dM  [         R                  R                  SS5      S:w  d)  [        5       (       a  [        R
                  " 5       (       a  g [        R                  5         [        R                  5         g )NÚTORCH_TNT_IN_USEr¦   rÚ   r§   )	r®   r¯   r�   r+   r   Ú	is_fbcoderd   r½   rÐ   rC   r:   r8   Úmaybe_warm_poolrœ  Û  si   € ä
�
‰
�‰Ð)¨3Ó/°3Ó6Ü�:‰:�>‰>Ð+¨SÓ1°SÓ8ä!×#Ñ#ô ×Ò×Ñàä×ÑÔô ×ÑÕr:   r”  )r]   r‰   r^   r–  )rŠ   Úint)gÚ
__future__r   Úatexitr   rI   Úloggingr°   r®   Úrerµ   Úconcurrent.futuresr   r   Úconcurrent.futures.processr   r   r   r	   Útypingr
   r   r   rW   Ú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   Ú+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$   Útorch._inductor.utilsr%   Útorch._inductor.virtualizedr&   Útorch._utils_internalr'   Ú	torch.hubr(   r)   Útorch.utils._ordered_setr*   Útorch.utils._tritonr+   Úcollections.abcr,   Útorch._inductor.runtime.hintsr-   r  r.   rS   r/   r‘   rX   ÚgetArtifactLoggerrŒ   r  Ú	getLoggerr«   r1   r;  r  r9   r6   rD   r[   r_   ÚplatformÚ_IS_WINDOWSrb   rh   rf   Úregister_at_forkÚAttributeErrorrq   rs   rd   rœ  ÚregisterrC   r:   r8   Ú<module>r¼     s§  ðæ "ã Û Û Û Û Û 	Û 	Û 
ß 9Ý 8Ý ß ß /Ñ /ã Ý K÷ó õ #÷÷ ÷ ñ ÷ñ õ
õ L÷õ 7Ý )Ý 3ß %Ý /Ý 2ö Ý(å8ÝJð Ð Ø€€_Ó à—.‘.×2Ñ2°8¸]ÓK€à×Ò˜Ó!€à>BÐ Ð;Ó Bà—:’:ØóÐ ò
ò <ô$ô&ô*3ð �l‰l˜gÑ%€à×Ò˜Ó!€ð �wÒÓ!€	ôò,ð	Ø×Ò zÒ2ô
"ð ÷22ð 22ó ð22÷jpñ pôfð0 ‡‚Ð(Õ )øð_ ó 	Ùð	ús   Å:G  Ç G	ÇG	