K iNddlmZddlZddlZddlZddlZddlZddlmZddl m Z m Z m Z m Z ddlmZddlmZmZdd lmZmZdd lmZdd lmZmZdd lmZGd deZGddZ ddZGddeZdZ y)) annotationsN)cached_property)DictTupleListOptional)knobs)KernelInterface JITFunction)OutOfResources PTXASError)driver)get_cache_manager triton_key)get_cache_invalidating_env_varscNeZdZ d d dZedZdZdZdZd dZ dZ y) AutotunerNc( |stidddg_n|_|_i_|_|xs7t j jxrt jj _ g_ |t|_ g_ |t|_ dd_ d_d_d_|r|_ d _n=t#jd kDst#jd kDr dfd }|_ |r|_d _n$t#jd kDr fd }|_d_d _d_| rc| j+dj$_| j+dj&_| j+dj(_|_|_t1j2j.s;j.j,_t1j2j.s;| _ _ _| _  | rAd dl}|j?dt@d| rddl!m" fd_yd dl#  fd_yy)a :param prune_configs_by: a dict of functions that are used to prune configs, fields: 'perf_model': performance model used to predicate running time with different configs, returns running time 'top_k': number of configs to bench 'prune_num_stages_by'(optional): a function used to prune num_stages. It takes configs:List[Config] as its input, and returns pruned configs. r ) num_warps num_stagesnum_ctasNFcyNr)kwargs reset_onlys ^/mnt/ssd/data/python-lab/Trading/venv/lib/python3.12/site-packages/triton/runtime/autotuner.pyz$Autotuner.__init__..0cyrr)r exceptions r!r"z$Autotuner.__init__..1r#r$TrcjD]}||j|s2jDcic]}|||jc}_yycc}wN) reset_to_zerozero_ restore_valueclonerestore_copies)rr nameselfs r! _pre_hookz%Autotuner.__init__.._pre_hook9s^ ..)D4L&&()!RVRdRd*e$41C1C1E+E*eD'"*esAcxjD]#}||jj|%i_yr()r+copy_r-)rr&r.r/s r! _post_hookz&Autotuner.__init__.._post_hookFs? ..BD4L&&t':':4'@AB&(#r$? perf_modeltop_kearly_config_prunez{warmup, rep, and use_cuda_graph parameters are deprecated. See https://github.com/triton-lang/triton/pull/4496 for details.) stacklevelr )do_bench_cudagraphc.||Sd|S)Nd)rep quantilesr) kernel_callr=r9r<s r!r"z$Autotuner.__init__..hs)@R"'A47'Ar$c^jj|nd|Sd|S)Nr;)warmupr<r=)testingdo_bench)r>r=r<tritonrAs r!r"z$Autotuner.__init__..psHFNN> !==002 2~~r$cX ddlm}tjj}|r%t dj j d|jjjz}|rtddj|dt|fij ij   fd} j|d S#t|t f$r:}|rt d |t#d t#d t#d gcYd}~Sd}~wwxYw) Nr )CompileTimeAssertionFailurezAutotuning kernel z with config zConflicting meta-parameters: , z8. Make sure that you don't re-define auto-tuned symbols.cjrjj jjij dy#t$r} j |#xYwd}~wwxYw)N)r&)rOrVrun ExceptionrP)eargsconfigcurrent full_nargsr/s r!r>z%Autotuner._bench..kernel_calls + MM* %   NN:N 6 NN:N;E  s)A!! B +A?>B?BBB )g?g?g?)r=zAutotuning failed with inf)compiler.errorsrjr rJprintrW__name__rGr ValueErrorjoindict all_kwargsnargsrCrrfloat) r/rqrpmetarjverbose conflictsr>rorrrss ``` @@r!_benchzAutotuner._bench~sA""((  &t||'<'<&=]6(S TIIK&--"4"4"66 ===H H ;ZH >/s34%L%,e = = >s CD)//D$D)$D)c |rtd|Dr|yddlm}|j}t |t s|j}t |t st }t|tjjj|jtt|jt|g|Dcgc] }t|c}z}t!j"dj%|j'dj)}t+|} |j,ddd} | j/| } | rt1| d 5} t3j4| d } | Dcic]\}}t7di||} }}t9j:| | j< |j>|<| |_ dddy || jCt3jD||j@jD cgc] \}} |jFr|jH| f"c} }d | dycc}wcc}}w#1swYy xYwcc} }w)Nc34K|]}|jywr()rO).0cfgs r! z-Autotuner.check_disk_cache..s A# AsFr) make_backend-zutf-8z.autotune.jsonrconfigs_timingsrbT)rbr)binaryr)%anytriton.compiler.compilerrrV isinstancer rrrrfget_current_targethash cache_keystrsorteditemshashlibsha256ryencode hexdigestrrwget_fileopenjsonloadrEbuiltinsminrUrHrputdumpsrO__dict__)r/ tuning_keyrFbench_fnrrVenv_varscrrH file_namepathcached_configstimingsrqtimings r!check_disk_cachezAutotuner.check_disk_cachesS A AA J9 WWR-BR-34 L 99; < A A C LL x~~'( )  O  % %SV % & NN388I#6#=#=g#FGQQS !),{{4C()8 ~~i( dC /N))N34EFJQR6+F+V3RR)1g7;;)O :&'.$  /   JJCGCWCWC]C]C_wgmgvgv&//7+w   %  )+ &S / xs0I/I% I#8I%#I18I1I%%I.c ttj_d}t j dkDrij}|j Dcic]\}}|jvs||}}}jDcgc] }||vs||c}|j D]6\} } t| dsjt| j8tjvrCd}jfd} jrj!| }n| j} nj d} | _t$j&j(rC|sAt)dj*j,dd j.d d j"d | j01ij| j3} | j1| j4j6i| j3}d_|Scc}}wcc}w) NTr dtypeFc tj}Dcic]}| jd|i}}tj}||z _tj||j  j <i j j j} j|d| _ ycc}w)NrqrT)r ) timer bench_timerrrUrHr|r{rOr) bench_startrqr bench_endrsrprbrpruned_configsr/s r! benchmarkz Autotuner.run..benchmarks"&))+KaopW]v{t{{D'R'R6'RRpGp $ I&/+&=DO&.ll7 &LDJJsO!YDJJ!Y&!YDJJsO+2D( qsC rzTriton autotuning for function z, with key as z, finished after z.2fzs, best config selected: ;)rzziprIr|rSrFrrGhasattrappendrrtuplerH prune_configsrMr best_configr rJrvrWrwrrOr{rVrm)r/rprused_cached_resultall_argskv_argsrb_argrrqrsretrs``` ` @r!rmz Autotuner.runs%#dnnd34 ! t|| q /$**//H(0(8Pfq!A)>sNT])^&KZZ_F\\!_F!    ! !*< 3DLL4I4I3J/Z]Y^_$$(OOC#88RSWScScRddeg h ?? &HDJJH&HF4E4E4GHJ OOJ 'dggkk   !    SQCs(I?I I I c >|j}|jr(|j|j|jfi|}|jr|j}t |t r'|dkr"tt|j|z}nt |ts tdt||kDr^|Dcic]1}||jdi|j||j3c}tjfdd|}|Scc}w)Nr4zPError while pruning configs, top_k must be either 1) a float <= 1.0 or 2) an intc|Sr(r)x est_timings r!r"z)Autotuner.prune_configs..s TUr$rr) rFr7r|r5rTrr}intrS TypeErrorr{rrG)r/rrr6rqrs @r!rzAutotuner.prune_configss  " "4T44T\\4::XQWXN ??&&E%'ESLC -56s+ rss>"U*#1   ODOO** !++- "( (9?V!WX^Y^!_s>6Dc tt|j||_g}|j |D]=}|j |j j|i||j?d|_|Sr() rzrrIr|rrrVrAr{)r/rprrautotune_configs r!rAzAutotuner.warmups#dnnd34 #11&9 O JJ~tww~~",,.     r$)NNNNNFNF)rczOptional[Dict])rrreturnz List[Config]) rw __module__ __qualname__rdrrCrrrmrrArr$r!rrsKkopt$a#1aF (>T(T.`0 r$rc6eZdZdZd dZdZdZdZdZdZ y) rEa An object that represents a possible kernel configuration for the auto-tuner to try. :ivar kwargs: a dictionary of meta-parameters to pass to the kernel as keyword arguments. :type kwargs: dict[Str, Any] :ivar num_warps: the number of warps to use for the kernel when compiled for GPUs. For example, if `num_warps=8`, then each kernel instance will be automatically parallelized to cooperatively execute using `8 * 32 = 256` threads. :type num_warps: int :ivar num_stages: the number of stages that the compiler should use when software-pipelining loops. Mostly useful for matrix multiplication workloads on SM80+ GPUs. :type num_stages: int :ivar num_ctas: number of blocks in a block cluster. SM90+ only. :type num_ctas: int :type maxnreg: Optional[int] :ivar maxnreg: maximum number of registers one thread can use. Corresponds to ptx .maxnreg directive. Not supported on all platforms. :ivar pre_hook: a function that will be called before the kernel is called. Parameters of this function are args. :ivar ir_override: filename of a user-defined IR (*.{ttgir|llir|ptx|amdgcn}). Ncf||_||_||_||_||_||_||_yr()rrrrmaxnregrO ir_override)r/rrrrrrOrs r!rdzConfig.__init__>s4 "  $   &r$cF|jdi|_|jdd|_|jdd|_|jdd|_|jdd|_|jd d|_|jd d|_y) Nrrrrrrr rrOr)rUrrrrrrOr)r/states r! __setstate__zConfig.__setstate__Gsii"- ;2))L!4 *a0 yyD1  *d3  99]D9r$ci|jd|jfd|jfd|jfd|jfd|j ffDcic] \}}||| c}}Scc}}w)Nrrrrr)rrrrrr)r/rrs r!r{zConfig.all_kwargsPs  kk  !$..1/!4??3 -"D$4$45  Q ]1    sA(cg}|jjD]\}}|j|d||jd|j|jd|j|jd|j |jd|j dj|S)Nz: z num_warps: z num_ctas: z num_stages: z maxnreg: rk)rrrrrrrry)r/resrrs r!__str__zConfig.__str__^sKK%%' $DAq JJ!Bqc{ # $ [ 012 Z /0 \$//!234 Yt||n-.yy~r$cntg|jj|jSr()rr{rrOrhs r!__hash__zConfig.__hash__hs,?doo'--/??@@r$ctg|jj|j}tg|jj|j}||k(Sr()rr{rrO)r/other self_tuple other_tuples r!__eq__z Config.__eq__ksy __  $ $ & MM       % % ' NN  [((r$)rrr NNN) rwrr__doc__rdrr{rrrrr$r!rErE's&,':  A )r$rEc >  f d} | S)a Decorator for auto-tuning a :code:`triton.jit`'d function. .. highlight:: python .. code-block:: python @triton.autotune(configs=[ triton.Config(kwargs={'BLOCK_SIZE': 128}, num_warps=4), triton.Config(kwargs={'BLOCK_SIZE': 1024}, num_warps=8), ], key=['x_size'] # the two above configs will be evaluated anytime # the value of x_size changes ) @triton.jit def kernel(x_ptr, x_size, BLOCK_SIZE: tl.constexpr): ... :note: When all the configurations are evaluated, the kernel will run multiple times. This means that whatever value the kernel updates will be updated multiple times. To avoid this undesired behavior, you can use the `reset_to_zero` argument, which resets the value of the provided tensor to `zero` before running any configuration. If the environment variable :code:`TRITON_PRINT_AUTOTUNING` is set to :code:`"1"`, Triton will print a message to stdout after autotuning each kernel, including the time spent autotuning and the best configuration. :param configs: a list of :code:`triton.Config` objects :type configs: list[triton.Config] :param key: a list of argument names whose change in value will trigger the evaluation of all provided configs. :type key: list[str] :param prune_configs_by: a dict of functions that are used to prune configs, fields: 'perf_model': performance model used to predicate running time with different configs, returns running time 'top_k': number of configs to bench 'early_config_prune'(optional): a function used to do early prune (eg, num_stages). It takes configs:List[Config] as its input, and returns pruned configs. :param reset_to_zero: a list of argument names whose value will be reset to zero before evaluating any configs. :type reset_to_zero: list[str] :param restore_value: a list of argument names whose value will be restored after evaluating any configs. :type restore_value: list[str] :param pre_hook: a function that will be called before the kernel is called. This overrides the default pre_hook used for 'reset_to_zero' and 'restore_value'. 'kwargs': a dict of all arguments passed to the kernel. 'reset_only': a boolean indicating whether the pre_hook is called to reset the values only, without a corresponding post_hook. :type pre_hook: lambda args, reset_only :param post_hook: a function that will be called after the kernel is called. This overrides the default post_hook used for 'restore_value'. 'kwargs': a dict of all arguments passed to the kernel. 'exception': the exception raised by the kernel in case of a compilation or runtime error. :type post_hook: lambda args, exception :param warmup: warmup time (in ms) to pass to benchmarking (deprecated). :type warmup: int :param rep: repetition time (in ms) to pass to benchmarking (deprecated). :type rep: int :param do_bench: a benchmark function to measure the time of each run. :type do_bench: lambda fn, quantiles :param cache_results: whether to cache autotune timings to disk. Defaults to False. "type cache_results: bool cJ t||j   S)N)rOrPrcrAr<r]rCrM)rrI) rVrMrFrCrbrPrOrcr<r)r+r]rAs r! decoratorzautotune..decorators9R\\7C `h#,?OX^dg(6Yfh hr$r) rFrbrcr)r+rOrPrAr<r]rCrMrs ```````````` r!autotunerwsvhh r$ceZdZddZdZy) Heuristicsc.||_||_||_yr()rVvaluesrI)r/rVrIrs r!rdzHeuristics.__init__s "r$c |jjD]1\}}|itt|j||||<3|j j |i|Sr()rrrzrrIrVrm)r/rprrheurs r!rmzHeuristics.runsg{{((* LGAtJS%> ?J6JKF1I Ltww{{D+F++r$N)rNone)rwrrrdrmrr$r!rrs # ,r$rcfd}|S)a Decorator for specifying how the values of certain meta-parameters may be computed. This is useful for cases where auto-tuning is prohibitively expensive, or just not applicable. .. highlight:: python .. code-block:: python # smallest power-of-two >= x_size @triton.heuristics(values={'BLOCK_SIZE': lambda args: triton.next_power_of_2(args['x_size'])}) @triton.jit def kernel(x_ptr, x_size, BLOCK_SIZE: tl.constexpr): ... :param values: a dictionary of meta-parameter names and functions that compute the value of the meta-parameter. each such function takes a list of positional arguments as input. :type values: dict[str, Callable[[dict[str, Any]], Any]] c2t||jSr()rrI)rVrs r!rzheuristics..decorators"bllF33r$r)rrs` r! heuristicsrs$4 r$) NNNNNNNFNF)! __future__rrrrXrr functoolsrtypingrrrrr jitr r errorsrrrrHrrtriton._C.libtritonrrrErrrrr$r!rsp"  %..-.0?QQhM)M)`txW\@F , ,r$