Ë
    ÃÓ^j€†  ã                  ó²  — U d dl mZ d dlZd dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
mZmZmZ d dlmZ d dlmZ d dlmZmZ d dlmZmZ d dlZd dlmZ d d	lmZmZmZmZ d d
lm Z  d dl!m"Z"m#Z#m$Z$m%Z%m&Z&m'Z'm(Z(m)Z)m*Z*m+Z+m,Z,m-Z- d dl.m/Z/m0Z0m1Z1 d dl2m3Z3 d dl4m5Z5 d dl6m7Z7m8Z8m9Z9m:Z: d dl;m<Z< d dl=m>Z> d dl?m@Z@ d dlAmBZBmCZC d dlDmEZE d dlFmGZG erd dlHmIZI d dlJmKZK d dlLmMZM daNdaOdePd<   ej¢                  j¥                  eSd«      ZT ejª                  eS«      ZVdaWdePd<    ej°                  d«      ZYd „ ZZd!„ Z[d1d"„Z\d1d#„Z]d2d$„Z^	 	 	 	 	 	 	 	 d3d%„Z_e	jÀ                  d&k(  Za ejª                  eS«      ZV eEe/   «       Zbd1d'„Zcd(„ Zd	  ejÊ                  ed¬)«       d4d*„Zgd5d+„Zhe< G d,„ d-«      «       Zi G d.„ d/«      Zjd1d0„Zk ejØ                  ec«       y# ef$ r Y Œ?w xY w)6é    )Ú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                 ó<   — t        «        ddlm} m} | r |«        yy)zG
    Setup that must be done prior to forking with a process pool.
    r   ©Ú
HAS_TRITONÚ
triton_keyN)Úcaching_device_propertiesÚ%torch._inductor.runtime.triton_compatr7   r8   r6   s     úh/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_inductor/async_compile.pyÚpre_fork_setupr<   Y   s   € ô Ô÷ MáÙ�ð ó    c                 ó~   — t        «       D ]0  \  } }|j                  «       sŒ|j                  j                  «        Œ2 y ©N)r   Úis_availableÚWorkerÚget_device_properties)Ú_Údevice_interfaces     r;   r9   r9   i   s8   € Ü?ÓAò <ÑˆÐØ×(Ñ(Õ*Ø×#Ñ#×9Ñ9Õ;ñ<r=   c                 ó6   — t         €
t        «       a t        €i ay y r?   )r2   r	   r4   © r=   r;   Ú_compile_startrG   o   s   € ä
€{Ü‹fˆÜÐ%Ø!#Ñð &r=   c                 óô   ‡— t         �t        «       } t        | t         z
  z  ad a t        rOt	        t        t        j                  «       «      «      Št        j                  j                  dd„ ˆfd„¬«       d ay y )NÚartifactc                 ó   — dddœS )NÚtriton_kernel_infoÚjson)ÚnameÚencodingrF   rF   r=   r;   ú<lambda>z_compile_end.<locals>.<lambda>ƒ   s   € Ø,Ø"ñ!€ r=   c                 ó.   •— t        j                  ‰ «      S r?   )rL   Údumps)Úsorted_infos   €r;   rO   z_compile_end.<locals>.<lambda>‡   s   ø€ œtŸz™z¨+Ó6€ r=   )Úmetadata_fnÚ
payload_fn)
r2   r	   Ú_cumulative_compile_timer4   ÚdictÚsortedÚitemsÚtorchÚ_loggingÚtrace_structured)Út1rR   s    @r;   Ú_compile_endr]   w   sp   ø€ ä
€Ü‹VˆÜ  B¬¡HÑ,Ð Øˆåäœ6Ô"8×">Ñ">Ó"@ÓAÓBˆÜ�‰×'Ñ'Øñó 7ð 	(ô 	
ð "&Ñð r=   c                ó$   — t         �
|t         | <   y y r?   )r4   )Úkernel_nameÚinfos     r;   Ú_add_triton_kernel_infora   Œ   s   € ô Ð)Ø.2Ô˜{Ò+ð *r=   c                ó|   — | j                   xs i }||d<   t        ||«       t        «       j                  d||«       y)ak  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_infora   r   Ú	add_top_n)Úkernelr_   Ú
elapsed_usr`   s       r;   Ú"_emit_triton_kernel_compile_metricri   “   sA   € ð ×%Ñ%Ò+¨€DØ(€DÐ	ÑÜ˜K¨Ô.ÜÓ×#Ñ#Ø(¨+°zõr=   Úwin32c                 ód   — t         D ]  } | j                  «        Œ dt        _        t	        «        y)z/Shut down all outstanding compile-worker pools.N)Ú	_pool_setÚshutdownÚAsyncCompileÚ_ready_futureÚ
after_fork)Úpools    r;   Úshutdown_compile_workersrr   °   s'   € äò ˆØ�‰�ðà!%„LÔÜ…Lr=   c                 ó~   — t         j                  «        dt        _        t        j                  j                  «        y)z7Reset pools to initial state without shutting them downN)rl   Úclearrn   ro   Úprocess_poolÚcache_clearrF   r=   r;   rp   rp   ¸   s&   € ä‡O�OÔØ!%„LÔÜ×Ñ×)Ñ)Õ+r=   )Úafter_in_childc                 ó|   — t         j                  €t        j                  «       t         _        t         j                  S )z}
    Temporary for internal rollout. Assign config.compile_threads lazily and return it.
    TODO: remove after rollout.
    )r   Úcompile_threadsÚdecide_compile_threadsrF   r=   r;   Úget_compile_threadsr{   Å   s-   € ô
 ×ÑÐ%Ü!'×!>Ñ!>Ó!@ŒÔÜ×!Ñ!Ð!r=   c                 ój   — t         j                  dk(  xs t        j                  «       j                   S )NÚ
subprocess)r   Úworker_start_methodÚmultiprocessingÚcurrent_processÚdaemonrF   r=   r;   Ú_process_pool_allowedr‚   Ï   s1   € ô 	×"Ñ" lÑ2ò 	8Ü×.Ñ.Ó0×7Ñ7Ð7ðr=   c                  óx   — e Zd ZU dZi Zded<   ed
d„«       Zedd„«       Zedd„«       Z	ed„ «       Z
edd„«       Zy	)Ú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                ó,   — t        | t        «       ¬«      S )a  
        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    r;   ÚkeyzCompiledTritonKernels.keyç   s   € ô ˜¬9«;Ô7Ð7r=   c                óT   — t         j                  | «      }|t         j                  |<   y)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ˆ   Úfuturer‰   s      r;   ÚsavezCompiledTritonKernels.saveñ   s%   € ô $×'Ñ'¨
Ó3ˆØ,2Ô×$Ñ$ SÒ)r=   c                ól   — t         j                  | «      }t         j                  j                  |d «      S r?   )r„   r‰   r…   Úget©rˆ   r‰   s     r;   r�   zCompiledTritonKernels.getþ   s+   € ä#×'Ñ'¨
Ó3ˆÜ$×+Ñ+×/Ñ/°°TÓ:Ð:r=   c                 ó   — i t         _        y r?   )r„   r…   rF   r=   r;   rv   z!CompiledTritonKernels.cache_clear  s   € à')ÔÕ$r=   c                óv   — t         j                  | «      }|t         j                  v rt         j                  |= y y r?   r‹   r�   s     r;   Úremove_futurez#CompiledTritonKernels.remove_future  s6   € ä#×'Ñ'¨
Ó3ˆð Ô'×.Ñ.Ñ.Ü%×,Ñ,¨SÑ1ð /r=   N)rˆ   Ústr)rˆ   r”   rŒ   r   )rˆ   r”   ÚreturnzCodeCacheFuture | None)rˆ   r”   r•   ÚNone)Ú__name__Ú
__module__Ú__qualname__Ú__doc__r…   Ú__annotations__Ústaticmethodr‰   r�   r�   rv   r“   rF   r=   r;   r„   r„   Ú   sv   … ñð *,€FÐ&Ó+àò8ó ð8ð ò
3ó ð
3ð ò;ó ð;ð ñ*ó ð*ð ò2ó ñ2r=   r„   c                  ó®  — e Zd ZU dZdZded<   dZded<   d"d„Ze e	j                  d«      d#d	„«       «       Zed
„ «       Ze e	j                  d«      d$d„«       «       Zed"d„«       Zed%d"d„«       Zed&d„«       Zed„ «       Zed"d„«       Zd'd(d„Zd)d„Zd)d„Zd*d„Zd+d„Zd,d„Zd,d„Zd,d„Z	 d,d„Zd-d„Zd„ Zd.d/d„Zd/d„Z 	 d.	 	 	 d/d„Z!d0d„Z"d1d „Z#d1d!„Z$y)2rn   z[
    Utilities to compile in thread pools or subprocess pools (in the case of Triton).
    NzFuture[Any] | Nonero   z'list[tuple[str, str, list[str]]] | NoneÚ_metal_sourcesc                 ó   — y r?   rF   )Úselfs    r;   Ú__init__zAsyncCompile.__init__  s   € Ør=   é   c                 óF   — t        «       dkD  sJ ‚t        t        «       «      S ©Nr¢   )r{   r   rF   r=   r;   rq   zAsyncCompile.pool  s$   € ô #Ó$ qÒ(Ð(Ð(Ü!Ô"5Ó"7Ó8Ð8r=   c                  ó   — y)z>No-op function to help mark when the subprocess pool is ready.ÚreadyrF   rF   r=   r;   Ú
_get_readyzAsyncCompile._get_ready!  s   € ð r=   c            	     óÜ  — t        «       dkD  sJ ‚t        «       st        d«      ‚d t        _        t
        j                  dt        j                  t        «       «       t        j                  dk(  r$t        t        «       t        j                  ¬«      } n‰t        j                  dk(  rdt        j                  d<   t        «        t        j                  t        j                  «      }t!        t        «       |t#        t$        t        j&                  «       «      ¬	«      } t        j(                  j+                  d | j,                  t.        j0                  ¬
«       t2        j5                  | «       | S )Nr¢   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.z"Creating '%s' pool with %d workersr}   )ÚquiesceÚspawnÚ0ÚTORCH_WARM_POOL)Ú
mp_contextÚinitializer)Úexitpriority)r{   r‚   ÚRuntimeErrorrn   ro   Úlogr`   r   r~   r!   Úquiesce_async_compile_poolÚosÚenvironr<   r   Úget_contextr"   r   r#   ÚgetpidÚutilÚFinalizerm   ÚsysÚmaxsizerl   Úadd)rq   Úctxs     r;   ru   zAsyncCompile.process_pool&  s  € ô #Ó$ qÒ(Ð(Ð(Ü$Ô&Üð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                ó|   — t        «       dk  s
t        «       sy t        «        | j                  «        t	        «        y r¤   )r{   r‚   rG   ru   r]   ©Úclss    r;   Ú	warm_poolzAsyncCompile.warm_poolT  s/   € äÓ  AÒ%Ô-BÔ-DØÜÔð 	×ÑÔÜ�r=   c                óv   — | j                  «        | j                  �| j                  j                  |¬«       y y )N©Útimeout)Úuse_process_poolro   Úresult)r¿   rÃ   s     r;   Úwait_pool_readyzAsyncCompile.wait_pool_ready_  s6   € à×ÑÔØ×ÑÐ(Ø×Ñ×$Ñ$¨WÐ$Õ5ð )r=   c                óh   — t        «       dk  r |«       S | j                  «       j                  |«      S r¤   )r{   rq   Úsubmit)r¿   Útasks     r;   rÈ   zAsyncCompile.submite  s+   € äÓ  AÒ%Ù“6ˆMØ�x‰x‹z× Ñ  Ó&Ð&r=   c                ó  — t        «       dk  s
t        «       syt        j                  j                  ry| j
                  s.| j                  «       j                  | j                  «      | _        | j
                  j                  «       S )Nr¢   F)
r{   r‚   r   ÚtritonÚproton_profilingro   ru   rÈ   r§   Údoner¾   s    r;   rÄ   zAsyncCompile.use_process_poolk  sh   € äÓ  AÒ%Ô-BÔ-DØô
 �=‰=×)Ò)Øð
 × Ò Ø #× 0Ñ 0Ó 2× 9Ñ 9¸#¿.¹.Ó IˆCÔØ× Ñ ×%Ñ%Ó'Ð'r=   c                óˆ   — | j                  «       sy| j                  «       }t        |t        «      r|j	                  «        yy)zq
        If using a SubprocPool, signal the sidecar process to start up its
        ProcessPoolExecutor.
        N)rÄ   ru   Ú
isinstancer!   Úwakeup)r¿   rq   s     r;   rÐ   zAsyncCompile.wakeup}  s9   € ð ×#Ñ#Ô%ØØ×ÑÓ!ˆÜ�dœKÔ(Ø�K‰K�Mð )r=   c                ó  ‡‡‡‡‡‡‡— t        j                  t        ‰‰«      Šˆfd„Št        d   dxx   dz  cc<   t        j                  d‰«       t        «        t        j                  j                  dd«      dk(  r=t        t        j                  j                  j                  j                  ‰«      ‰«      S | j!                  «       }t#        d	|«       t        j$                  j&                  j)                  «       Št        t*        j,                  d
d«      Št.        j                  ‰«      x}�Tt        d   dxx   dz  cc<   t1        |t2        «      rt.        j5                  ‰«       ‰|_        |r|S |j9                  «       S |�r—t;        «        g d¢}|D �ci c])  }|t        j                  v sŒ|t        j                  |   “Œ+ }}dt        j                  j<                  j>                  i}	tA        t        j                  j<                  jB                  «      dkD  r©tD        jG                  ‰«      }
|
r|
jI                  d«      }ntK        d«      }‰jM                  d«      d   }ddl'm(}  |||«      }|t        j                  j<                  jB                  v r,|t        j                  j<                  jB                  |   i|	d<   | jS                  «       jU                  tV        ‰||	«      Šdˆˆˆˆˆˆfd„}tY        |‰¬«      }t.        j[                  ‰|«       |S t]        ddddd¬«      5  d}	 t_        «       }ta        «        t;        «         ‰«       }|jc                  ‰‰«       |je                  dt.        jg                  ‰«      ¬«       t_        «       |z
  dz  }ti        |‰|«       |tk        |¬«       cddd«       S c c}w # tl        $ r}tK        |«      }‚ d}~ww xY w# tk        |¬«       w xY w# 1 sw Y   yxY w)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                 óT   •— t        d«      5   ‰ «       cd d d «       S # 1 sw Y   y xY w)NÚreload_kernel_in_parent)r   )Úload_kernels   €r;   rÓ   z4AsyncCompile.triton.<locals>.reload_kernel_in_parent£  s&   ø€ äÐ7Ó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ÚTRITON_LIBDEVICE_PATHÚuse_static_triton_launcherr   z@triton.jit
)Ú%generate_lookup_hash_from_source_codeÚautotune_lookup_tablec                 óP  •— 	 ‰j                  «       \  } }| j                  ‰‰«       t        j                  ‰«       | j                  d ¬«       | j                  d‰t        j                  ‰«      ¬«       t        | ‰|«       | S # t        $ r}|j                  ‰«      |‚d }~ww xY w)N)Ú
old_valuesF)Úwarm_cache_onlyÚreload_kernelÚstatic_triton_bundle_key)
rÅ   r    Ú	with_nameÚset_compile_infor„   r“   Úrestore_after_unpickleÚ
precompiler‰   ri   )	rg   rh   ÚeÚ
compile_idrÚ   r_   rÓ   Úsource_coderÉ   s	      €€€€€€r;   Ú
get_resultz'AsyncCompile.triton.<locals>.get_resultð  s¦   ø€ ð:Ø)-¯©«Ñ&�F˜Jð ×'Ñ'¨
°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•   r1   )7Ú	functoolsr   r   r   Úkernel_code_logr`   rG   r³   r´   r�   ÚgetattrrY   Ú	_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Úgroupr”   ÚsplitÚ)torch._inductor.runtime.triton_heuristicsrà   ru   rÈ   r'   r   r�   r   r
   r%   rè   rê   r‰   ri   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_nsrg   rh   rë   rì   rÚ   rÔ   rÓ   rÉ   s    ``                  @@@@@r;   rË   z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ä%×3Ñ3°KÔ@Ø0G�Ô-ÙØ�à—}‘}“Ð&ò ä&Ô(òˆHð
 4<ÖO¨a¸qÄBÇJÁJº˜œBŸJ™J q™MÑ)ÐOˆIÐOà,¬e¯o©o×.DÑ.D×._Ñ._ðˆLô ”5—?‘?×)Ñ)×?Ñ?Ó@À1ÒDÜ$×+Ñ+¨KÓ8�ÙØ%&§W¡W¨Q£Z‘Nä%(¨£Y�Nà(×.Ñ.¨Ó?ÀÑB�
õñ @Ø" Jó�ð œeŸo™o×4Ñ4×JÑJÑJà¤§¡×!7Ñ!7×!MÑ!MÈgÑ!Vð=�LÐ!8Ñ9ð ×$Ñ$Ó&×-Ñ-Ü&ØØØó	ˆD÷ò ô* " *°TÔ:ˆFÜ!×&Ñ& {°FÔ;ØˆMäØ*Ø&*Ø)AØ $Ø*:ôñ 1ð �ð1Ü&›y�HÜ*Ô,Ü.Ô0Ù(›]�FØ×+Ñ+¨J¸ÔDØ×%Ñ%Ø(-Ü1F×1JÑ1JÈ;Ó1Wð &ô ô #*£)¨hÑ"6¸4Ñ!?�JÜ6°v¸{ÈJÔWØ!ô
 &¨4Õ0÷31ñ 1ùòu Pøô^ !ò Ü˜q›6�DØûðûô &¨4Ö0ú÷31ð 1úsI   ÆOÆ3OÌ0O8Ì4A;O
Î/O8Ï
	O$ÏOÏO$Ï$O'Ï'O5Ï5O8Ï8Pc                ó   — ddl m}  ||i |¤ŽS )Nr   )ÚMultiKernelCall)Ú$torch._inductor.codegen.multi_kernelr  )r    ÚargsÚkwargsr  s       r;   Úmulti_kernelzAsyncCompile.multi_kernel$  s   € ÝHñ  Ð/¨Ñ/Ð/r=   c                ó   — ddl m}  ||i |¤ŽS )Nr   )ÚSizeHintMultiKernelCall)r  r  )r    r  r  r  s       r;   Úsize_hint_multi_kernelz#AsyncCompile.size_hint_multi_kernel*  s   € ÝPá&¨Ð7°Ñ7Ð7r=   c                óæ   ‡— t         j                  d|«       t        «       dk  rt        j                  |«      j
                  S t        j                  || j                  ¬«      Št        ˆfd„«      S )NzCPP Kernel:
%sr¢   ©Ú	submit_fnc                 ó&   •—  ‰ «       j                   S r?   )rg   )rî   s   €r;   rO   z"AsyncCompile.cpp.<locals>.<lambda>5  s   ø€ ©
«×(;Ñ(;€ r=   )	rù   r`   r{   r   rþ   rg   Ú
load_asyncrÈ   r   )r    rí   rî   s     @r;   ÚcppzAsyncCompile.cpp/  sX   ø€ Ü×ÑÐ.°Ô<ÜÓ  AÒ%Ü×$Ñ$ [Ó1×8Ñ8Ð8ä%×0Ñ0°ÈÏÉÔTˆJÜÓ ;Ó<Ð<r=   c                óÎ   — t         j                  d|«       t        «       dk  rt        j                  ||«      S t        j
                  ||| j                  ¬«      }t        |«      S )NzCPP+Bindings Kernel:
%sr¢   r   )rù   r`   r{   r   Úload_pybindingÚload_pybinding_asyncrÈ   r   )r    Úargtypesrí   rî   s       r;   Úcpp_pybindingzAsyncCompile.cpp_pybinding7  sY   € Ü×ÑÐ7¸ÔEÜÓ  AÒ%Ü-×<Ñ<¸XÀ{ÓSÐSä3×HÑHØ˜+°·±ôˆJô   
Ó+Ð+r=   c                ó<   ‡‡‡‡— ˆˆˆˆfd„}| j                  |«      S )Nc                 ó�   •— ‰r/‰j                  ‰d«      ^} }‰j                  j                  | «       ‰j                  ‰‰«      d   S )NÚor   )ÚcompileÚaot_kernels_oÚappendrþ   )Úoutput_pathrC   Úaot_compileÚ	cache_clsÚdst_file_extrí   s     €€€€r;   rÉ   z"AsyncCompile.cutlass.<locals>.taskB  sH   ø€ Ùð #,×"3Ñ"3°KÀÓ"E��˜aØ×'Ñ'×.Ñ.¨{Ô;Ø—>‘> +¨|Ó<¸QÑ?Ð?r=   )rÈ   )r    r2  rí   r3  r1  rÉ   s    ```` r;   ÚcutlasszAsyncCompile.cutlassA  s   û€ ÷	@ð �{‰{˜4Ó Ð r=   c                ó^   — t         j                  d|«       | j                  t        |||«      S )NzCUDA Kernel:
%s)rù   r`   r4  r   ©r    rí   r3  r1  s       r;   ÚcudazAsyncCompile.cudaL  s(   € Ü×ÑÐ/°Ô=Ø�|‰|œM¨;¸ÀkÓRÐRr=   c                ó^   — t         j                  d|«       | j                  t        |||«      S )NzXPU Kernel:
%s)rù   r`   r4  r   r6  s       r;   ÚxpuzAsyncCompile.xpuP  s(   € Ü×ÑÐ.°Ô<Ø�|‰|œL¨+°|À[ÓQÐQr=   c                ód   ‡‡‡— t         j                  d‰«       ˆˆˆfd„}| j                  |«      S )NzROCm Kernel:
%sc                 ó  •— ‰r8t        j                  ‰d¬«      ^} }t         j                  j                  | «       t        j
                  j                  rt        j                  ‰d¬«      }t        j                  ‰‰«      d   S )Nr,  )r3  Úexer   )r   r-  r.  r/  r   ÚrocmÚgenerate_test_runnerrþ   )r0  rC   r1  r3  rí   s     €€€r;   rÉ   zAsyncCompile.rocm.<locals>.task\  sh   ø€ ÙÜ"/×"7Ñ"7¸ÐRUÔ"V��˜aÜ×+Ñ+×2Ñ2°;Ô?Ü�{‰{×/Ò/Ü!×)Ñ)¨+ÀEÔJ�Ü ×%Ñ% k°<Ó@ÀÑCÐCr=   )rù   r`   rÈ   )r    rí   r3  r1  rÉ   s    ``` r;   r=  zAsyncCompile.rocmT  s-   ú€ ô 	×ÑÐ/°Ô=ö	Dð �{‰{˜4Ó Ð r=   c                óÐ   — t         j                  d||«       t        «       dk  rt        j                  ||«      S t        j
                  ||| j                  ¬«      }t        |«      S )NzHalide Kernel:
%r
%sr¢   r   )rù   r`   r{   r   Úgenerate_halideÚgenerate_halide_asyncrÈ   r   )r    Úmetarí   rî   s       r;   ÚhalidezAsyncCompile.halidef  s[   € Ü×ÑÐ5°t¸[ÔIÜÓ  AÒ%Ü"×2Ñ2°4¸ÓEÐEä(×>Ñ>Ø�k¨T¯[©[ôˆJô   
Ó+Ð+r=   c                ó    — t         j                  j                  j                  j	                  ||«      }|› d|› �} |t        ||«      |¬«      S )zAReload a kernel module from PyCodeCache and wrap the entry point.rC   ©Úkernel_path)rY   rû   rü   rý   Úload_by_key_pathrú   )r    r_   Úmain_suffixÚwrapper_clsr‰   ÚpathÚmodÚmain_func_names           r;   Ú_load_kernel_wrapperz!AsyncCompile._load_kernel_wrapperp  sH   € ä�o‰o×'Ñ'×3Ñ3×DÑDÀSÈ$ÓOˆØ'˜=¨¨+¨Ð7ˆÙœ7 3¨Ó7ÀTÔJÐJr=   c           	     ó"  ‡ ‡‡‡‡— ddl mŠmŠ t        j	                  d|«       t        «        ‰ j                  «       }|rwddg}|D �ci c])  }|t        j                  v sŒ|t        j                  |   “Œ+ }}‰ j                  «       j                  t        ‰|‰||«      Šdˆˆˆˆ ˆfd„}t        |‰¬«      S t        j                  j                  j                   j#                  |«      \  }	}
t        j                  j                  j                   j%                  |	|
«      }‰› d‰› �}t'        ||«      s?t)        |«      D �cg c]  }t+        t-        ||«      «      sŒ|‘Œ }}t/        d	|› d
|› �«      ‚ ‰t-        ||«      |
¬«      S c c}w c c}w )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.
        r   )ÚCuteDSLKernelWrapperÚMAIN_SUFFIXzCuteDSL Kernel:
%srÜ   ÚTORCHINDUCTOR_CUTLASS_DIRc                 óÌ   •— 	 ‰j                  «       \  } }}t        j	                  d‰|«       ‰j                  ‰‰‰| |«      S # t        $ r}|j                  ‰«      |‚d }~ww xY w)Nz0CuteDSL kernel %s compiled in subprocess in %dus©rÅ   r    rç   r±   ÚdebugrM  )	r‰   rJ  rh   rë   rO  rP  r_   r    Úsubprocess_tasks	       €€€€€r;   rî   z(AsyncCompile.cutedsl.<locals>.get_result›  sx   ø€ ð:Ø,;×,BÑ,BÓ,DÑ)�C˜˜zô —	‘	ØFØØôð
 ×0Ñ0Ø Ð.BÀCÈóð øô (ò :ØŸ+™+ kÓ2¸Ð9ûð:úó   ƒA Á	A#ÁAÁA#rï   rC   z-Could not find CuteDSL main kernel function 'ú'. Available callables: rE  )r•   rO  )Ú.torch._inductor.codegen.cutedsl.cutedsl_kernelrO  rP  rù   r`   rG   rÄ   r³   r´   ru   rÈ   r&   r   rY   rû   rü   rý   ÚwriterG  ÚhasattrÚdirÚcallablerú   r°   )r    r_   rí   Úprecompile_metadatar  r  r  r  rî   r‰   rJ  rK  rL  rM   Ú	availablerO  rP  rU  s   ``             @@@r;   ÚcutedslzAsyncCompile.cutedslv  sw  ü€ ÷	
ô
 	×ÑÐ2°KÔ@ÜÔà×+Ñ+Ó-ˆáØ1Ð3NÐOˆHØ3;ÖO¨a¸qÄBÇJÁJº˜œBŸJ™J q™MÑ)ÐOˆIÐOà"×/Ñ/Ó1×8Ñ8Ü2ØØØØØ#óˆO÷ñ ô   
°?ÔCÐCäŸ™×1Ñ1×=Ñ=×CÑCÀKÓP‰IˆC�Ü—/‘/×+Ñ+×7Ñ7×HÑHÈÈdÓSˆCà +˜}¨A¨k¨]Ð;ˆNÜ˜3 Ô/Ü.1°#«hÖW d¼(Ä7È3ÐPTÓCUÕ:VšTÐW�	ÐWÜ"ØCÀNÐCSÐSkÐluÐkvÐwóð ñ (¬°°^Ó(DÐRVÔWÐWùòK Pùò@ Xs   ÁFÁ FÅFÅFc                ó¾   ‡‡‡‡‡— ddl mŠmŠ t        j	                  d‰«       ˆˆˆˆfd„}t        «       dk  r |«       S | j                  |«      Št        ˆfd„«      S )ax  
        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   )rP  ÚPallasKernelWrapperzPallas Kernel:
%sc            	     ó®  •— t         j                  j                  j                  j	                  ‰	«      \  } }t         j                  j                  j                  j                  | |«      }‰› d‰› �}t        ||«      s?t        |«      D �cg c]  }t        t        ||«      «      sŒ|‘Œ }}t        d|› d|› �«      ‚ ‰t        ||«      |¬«      S c c}w )NrC   z,Could not find Pallas main kernel function 'rW  rE  )rY   rû   rü   rý   rY  rG  rZ  r[  r\  rú   r°   )
r‰   rJ  rK  rL  rM   r^  rP  ra  r_   rí   s
         €€€€r;   rÉ   z!AsyncCompile.pallas.<locals>.taskÇ  sÉ   ø€ ÜŸ™×1Ñ1×=Ñ=×CÑCÀKÓP‰IˆC�Ü—/‘/×+Ñ+×7Ñ7×HÑHÈÈdÓSˆCð !,˜}¨A¨k¨]Ð;ˆNÜ˜3 Ô/Ü.1°#«hÖW d¼(Ä7È3ÐPTÓCUÕ:VšTÐW�	ÐWÜ"ØBÀ>ÐBRÐRjÐktÐjuÐvóð ñ '¤w¨s°NÓ'CÐQUÔVÐVùò Xs   ÂCÂ'Cr¢   c                 ó$   •— ‰ j                  «       S r?   )rÅ   rï   s   €r;   rO   z%AsyncCompile.pallas.<locals>.<lambda>Ù  s   ø€ ¨¯©«€ r=   )Útorch._inductor.codegen.pallasrP  ra  rù   r`   r{   rÈ   r   )r    r_   rí   rÉ   rP  ra  rŒ   s    `` @@@r;   ÚpallaszAsyncCompile.pallas·  sN   ü€ ÷ 	Tä×ÑÐ1°;Ô?÷	Wô Ó  AÒ%Ù“6ˆMà—[‘[ Ó&ˆFÜÓ 7Ó8Ð8r=   c           	     ó*  ‡ ‡‡‡‡— ddl mŠ ddlmŠ t        j                  d|«       t        «        ‰ j                  «       }|rwddg}|D �ci c])  }|t        j                  v sŒ|t        j                  |   “Œ+ }}‰ j                  «       j                  t        ‰|‰||«      Šdˆˆˆˆ ˆfd„}t        |‰¬«      S t        j                  j                   j"                  j%                  |«      \  }	}
t        j                  j                   j"                  j'                  |	|
«      }‰› d	‰› �}t)        ||«      s?t+        |«      D �cg c]  }t-        t/        ||«      «      sŒ|‘Œ }}t1        d
|› d|› �«      ‚ ‰t/        ||«      |
¬«      S c c}w c c}w )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_api library.
            We use the PyCodeCache to write the source code to a file and load it.
        r   )ÚNVUniversalGemmKernelWrapper)rP  z NVIDIA Universal GEMM Kernel:
%srÜ   rQ  c                 óÌ   •— 	 ‰j                  «       \  } }}t        j	                  d‰|«       ‰j                  ‰‰‰| |«      S # t        $ r}|j                  ‰«      |‚d }~ww xY w)Nz:NV Universal GEMM kernel %s compiled in subprocess in %dusrS  )	r‰   rJ  rh   rë   rP  rg  r_   r    rU  s	       €€€€€r;   rî   z2AsyncCompile.nv_universal_gemm.<locals>.get_result  sy   ø€ ð:Ø,;×,BÑ,BÓ,DÑ)�C˜˜zô —	‘	ØPØØôð
 ×0Ñ0Ø Ð.JÈCÐQUóð øô (ò :ØŸ+™+ kÓ2¸Ð9ûð:úrV  rï   rC   z;Could not find NVIDIA Universal GEMM main kernel function 'rW  rE  )r•   rg  )ÚBtorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_kernelrg  ÚFtorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_schedulingrP  rù   r`   rG   rÄ   r³   r´   ru   rÈ   r&   r   rY   rû   rü   rý   rY  rG  rZ  r[  r\  rú   r°   )r    r_   rí   r]  r  r  r  r  rî   r‰   rJ  rK  rL  rM   r^  rP  rg  rU  s   ``             @@@r;   Únv_universal_gemmzAsyncCompile.nv_universal_gemmÛ  s„  ü€ õ 	
õ	
ô 	×ÑÐ@À+ÔNÜÔà×+Ñ+Ó-ˆáØ1Ð3NÐOˆHØ3;ÖO¨a¸qÄBÇJÁJº˜œBŸJ™J q™MÑ)ÐOˆIÐOà"×/Ñ/Ó1×8Ñ8Ü2ØØØØØ#óˆO÷ñ ô   
°?ÔCÐCäŸ™×1Ñ1×=Ñ=×CÑCÀKÓP‰IˆC�Ü—/‘/×+Ñ+×7Ñ7×HÑHÈÈdÓSˆCà +˜}¨A¨k¨]Ð;ˆNÜ˜3 Ô/Ü.1°#«hÖW d¼(Ä7È3ÐPTÓCUÕ:VšTÐW�	ÐWÜ"ðØ&Ð'Ð'?À	¸{ðLóð ñ
 0Ü˜˜^Ó,¸$ôð ùòM Pùò@ Xs   ÁFÁ$FÅFÅ Fc                óf   — | j                   €g | _         | j                   j                  |||f«       y)zVRegister a Metal kernel body; wait() compiles all registered kernels into one library.N)rž   r/  )r    r_   ÚsourceÚheaderss       r;   ÚmetalzAsyncCompile.metal#  s1   € à×ÑÐ&Ø"$ˆDÔØ×Ñ×"Ñ" K°¸Ð#AÕBr=   c                ó8  — t        «       dkD  r*t        ddddd¬«      5  | j                  |«       d d d «       | j                  rAddlm} |j                   || j                  «      «       | j                  j                  «        t        «        y # 1 sw Y   ŒaxY w)	Nr¢   zasync_compile.waitTrð   rñ   rò   r   )Úcompile_mps_shaders)	r{   r   Ú_wait_futuresrž   Ú%torch._inductor.runtime.runtime_utilsrq  Úupdatert   r]   )r    Úscoperq  s      r;   ÚwaitzAsyncCompile.wait)  s†   € ÜÓ  1Ò$ÜØ$Ø&*Ø)AØ $Ø*:ôñ *ð ×"Ñ" 5Ô)÷*ð ×ÒÝQà�L‰LÑ,¨T×-@Ñ-@ÓAÔBØ×Ñ×%Ñ%Ô'ä�÷*ð *ús   žBÂBc           	     óR  — |j                  «       D ��ci c]  \  }}t        |t        t        f«      r||“Œ  }}}t	        t        |«      dt        j                  d¬«      }t        j                  xs d }|j                  «       D ]_  \  }}t        j                  r!t        |t        «      s|j                  |«       	 |j                  |¬«      }|||<   |j!                  d	«       Œa y c c}}w # t        $ r}	t        d|›d|› d�«      |	‚d }	~	wt        $ r}	t        d«      |	‚d }	~	ww xY w)
NzInductor Compilationr   )ÚtotalÚdescÚdisableÚdelayrÂ   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.r¢   )rX   rÏ   r   r   r,   r  r   Údisable_progressÚcompile_worker_wait_timeoutÚverbose_progressr+   Úset_postfix_strrÅ   ÚFuturesTimeoutErrorr°   r   rt  )
r    ru  r‰   ÚvalueÚkernelsÚpbarÚwait_timeoutrÅ   rg   rë   s
             r;   rr  zAsyncCompile._wait_futures<  sB  € ð $Ÿk™k›m÷
á��UÜ˜%¤&¬/Ð!:Ô;ð �‰Jð
ˆñ 
ô
 Ü�g“,Ø'Ü×+Ñ+Øô	
ˆô ×9Ñ9ÒA¸TˆØ"Ÿ=™=›?ò 	‰KˆC�Ü×&Ò&¬z¸$Ä	Ô/JØ×$Ñ$ SÔ)ðØŸ™¨|˜Ó<�Ø#��c‘
ð" �K‰K˜�Nñ-	ùó
øô* 'ò ô #Ø9¸#¸ð A'Ø'3 nð 5KðKóð ð	ûô
 %ò Ü"ðIóð
 ðûðús)   ”#C'Â<C-Ã-	D&Ã6D	Ä	D&ÄD!Ä!D&©r•   r–   )r•   r   )r•   r   )éx   )rÉ   zCallable[..., Any]r•   r   )r7  )r_   r”   rí   r”   r  r”   )r•   r   )rí   r”   )r(  ú	list[str]rí   r”   )F)rB  r0   rí   r”   r?   )r_   r”   rí   r”   )r_   r”   rm  r”   rn  r‡  r•   r–   )ru  údict[str, Any]r•   r–   )%r—   r˜   r™   rš   ro   r›   rž   r¡   rœ   rø   Ú	lru_cacherq   r§   ru   ÚclassmethodrÀ   rÆ   rÈ   rÄ   rÐ   rË   r  r  r$  r)  r4  r7  r9  r=  rC  rM  r_  re  rk  ro  rv  rr  rF   r=   r;   rn   rn     sm  … ñð )-€MÐ%Ó,Ø>B€NÐ;ÓBóð Ø€Y×Ñ˜Óò9ó ó ð9ð ñó ðð Ø€Y×Ñ˜Óò*ó ó ð*ðX òó ðð ó6ó ð6ð
 ò'ó ð'ð
 ñ(ó ð(ð" ò	ó ð	ôY1óv0ó8ó
=ó,ó	!óSóRð ó	!ó$,òKô?XóB"9ðJ GKðFØðFØ-0óFóPCóô&&r=   rn   c                 ó  — t         j                  j                  dd«      dk(  sAt         j                  j                  dd«      dk7  st        «       rt	        j
                  «       ry t        j                  «        t        j                  «        y )NÚTORCH_TNT_IN_USEr«   rØ   r¬   )	r³   r´   r�   r.   r   Ú	is_fbcodern   rÀ   rÐ   rF   r=   r;   Úmaybe_warm_poolrŽ  e  sc   € ä
�
‰
�‰Ð)¨3Ó/°3Ò6Ü�:‰:�>‰>Ð+¨SÓ1°SÒ8ä!Ô#ô ×ÑÔàä×ÑÔô ×ÑÕr=   r…  )r_   r”   r`   rˆ  )rg   r1   r_   r”   rh   Úintr•   r–   )r•   r�  )r•   Úbool)mÚ
__future__r   Úatexitrø   rL   Úloggingr   r³   Úrer¹   Úconcurrent.futuresr   r   r   r€  Úconcurrent.futures.processr   r   r	   r
   Útypingr   r   rY   Ú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   r	  r1   rU   r2   r›   rZ   ÚgetArtifactLoggerr—   rù   Ú	getLoggerr±   r4   r-  r  r<   r9   rG   r]   ra   ri   ÚplatformÚ_IS_WINDOWSrl   rr   rp   Úregister_at_forkÚAttributeErrorr{   r‚   r„   rn   rŽ  ÚregisterrF   r=   r;   ú<module>r¯     sì  ðæ "ã Û Û Û Û Û 	Û 	Û 
÷ñ õ
 9Ý ß ß %ã Ý K÷ó õ #÷÷ ÷ ó ÷ñ õ
õ L÷ó õ 7Ý )Ý 3ß %Ý /Ý 2ñ Ý(å8ÝJð Ð Ø€€\Ó à—.‘.×2Ñ2°8¸]ÓK€à€g×Ñ˜Ó!€à;?Ð Ð8Ó ?à�2—:‘:ØóÐ ò
ò <ó$ó&ó*3ðØðàðð ðð 
ó	ð* �l‰l˜gÑ%€à€g×Ñ˜Ó!€ð  ˆJ�wÑÓ!€	óò,ð	Ø€B×Ñ zÕ2ó
"óð ÷22ð 22ó ð22÷jR	ñ R	ójð0 €‡�Ð(Õ )øðy ò 	Ùð	ús   ÆG ÇGÇG