kernels/pr_754/en/_app/immutable/nodes/4.DBszlrKN.js

import"../chunks/DsnmJJEf.js";import{i as C,h as b,C as k,H as n,a,E as A,s as B}from"../chunks/CKEMnrQO.js";import{p as N,o as Q,s,f as S,a as w,b as Z,d as m,n as _}from"../chunks/Bb-LL0eD.js";const E='{"title":"Ship Triton autotune configurations","local":"ship-triton-autotune-configurations","sections":[{"title":"Shipping data files with a kernel","local":"shipping-data-files-with-a-kernel","sections":[],"depth":2},{"title":"Configuration file layout","local":"configuration-file-layout","sections":[],"depth":2},{"title":"Looking up configurations at runtime","local":"looking-up-configurations-at-runtime","sections":[],"depth":2},{"title":"Generating the configurations","local":"generating-the-configurations","sections":[],"depth":2},{"title":"Impact","local":"impact","sections":[],"depth":2}],"depth":1}';var G=m('<meta name="hf:doc:metadata"/>'),V=m(`<p></p> <!> <!> <p>Triton kernels typically have parameters, such as tile sizes, number of warps and
number of pipeline stages, whose optimal values depend on the GPU and the
problem shape. Autotuning finds good values for these parameters, but doing
it at runtime (e.g. with the <a href="https://triton-lang.org/main/python-api/generated/triton.autotune.html" rel="nofollow"><code>@triton.autotune</code></a> decorator) re-benchmarks every candidate configuration in each new process.</p> <p>However, one can run the tuner for the GPU models they like and store the
best found configurations as files for using them later. This effectively
reduces the potentially costly tuning time.</p> <p><code>kernel-builder</code> support this by packaging these configurations as JSON files
with the kernel. At runtime, the kernel looks up the configuration for the
current GPU and shape and falls back to sensible defaults when there is no
matching configuration.</p> <p>This is how, for example, the vLLM fused MoE kernel ships on the Hub: the <a href="https://huggingface.co/RedHatAI/moe/tree/main/torch-ext/moe/configs" rel="nofollow"><code>RedHatAI/moe</code></a> repository contains a <code>configs</code> directory with tuned configurations for many
GPUs. This page walks through a small, complete example of the same pattern:
the <a href="https://github.com/huggingface/kernels/tree/main/examples/kernels/gemm-triton-autotune" rel="nofollow"><code>gemm-triton-autotune</code></a> example kernel, a Triton GEMM published as <a href="https://huggingface.co/kernels-test/gemm-triton-autotune" rel="nofollow"><code>kernels-test/gemm-triton-autotune</code></a>.</p> <!> <p>Since autotune files are plain JSON files, we can store them anywhere inside the
kernels main Python sources in <code>torch-ext/&lt;kernel_name&gt;</code>. For this example,
we will use <code>torch-ext/&lt;kernel_name&gt;/configs/</code>. By default, only <code>py</code> and <code>pyi</code> files are picked up from the kenel’s Python source directory, so add <code>json</code> to the <a href="writing-kernels#torch-noarch"><code>pyext</code> option</a> in <code>build.toml</code>:</p> <!> <!> <p>A GEMM computes <code>(M, K) @ (K, N)</code>. For a model, the weight dimensions <code>N</code> and <code>K</code> are known ahead of time, while <code>M</code> (e.g. the number of tokens)
varies at runtime. The example therefore stores one file per <code>(N, K)</code> shape
and GPU, following the same naming convention as the MoE kernel:</p> <!> <p>Each file maps an <code>M</code> value to the best configuration that the autotuner
found for that <code>M</code>:</p> <!> <p>Since the device name is part of the file name, configurations tuned for one
GPU are never applied to another. A configuration that would exceed the
resources of a smaller GPU (e.g. shared memory) is therefore harmless to
ship.</p> <!> <p>At kernel launch, the kernel checks whether a configuration file exists for
the current device and shape. If it does, the configuration with the nearest
tuned <code>M</code> is used; otherwise the kernel falls back to a conservative default
and logs a warning. The lookup is cached, so the file is read at most once
per process:</p> <!> <p>The configuration is then passed to the Triton kernel as its <code>constexpr</code> and launch parameters (see <a href="https://github.com/huggingface/kernels/blob/main/examples/kernels/gemm-triton-autotune/torch-ext/gemm_triton_autotune/gemm.py" rel="nofollow"><code>gemm.py</code></a> in the example):</p> <!> <!> <p>The autotuner itself is ordinary benchmarking code: for each <code>M</code>, benchmark
every candidate configuration with <code>triton.testing.do_bench</code> and keep the
fastest. The example ships the tuner as part of the kernel (the <code>tune_gemm</code> function in <a href="https://github.com/huggingface/kernels/blob/main/examples/kernels/gemm-triton-autotune/torch-ext/gemm_triton_autotune/tuning.py" rel="nofollow"><code>tuning.py</code></a>),
so users can also generate configurations for GPUs that the kernel author
did not tune. Candidates that do not fit the device (Triton raises <code>OutOfResources</code>) are skipped.</p> <p>The example repository contains a small script, <a href="https://github.com/huggingface/kernels/blob/main/examples/kernels/gemm-triton-autotune/tune.py" rel="nofollow"><code>tune.py</code></a>,
that runs the tuner and writes the configuration files to the source tree:</p> <!> <p>Commit the generated files, rebuild, and the configurations ship with the
kernel. When tuning a kernel that is not yet published, build it locally
(see <a href="local-dev">Develop locally</a>) and point <code>LOCAL_KERNELS</code> at the
build:</p> <!> <!> <p>Tuned configurations are cheap to ship and can make a large difference. On
an NVIDIA L4, the tuned configuration for a <code>(1024, 4096) @ (4096, 4096)</code> float16 GEMM is ~1.4× faster than the example’s heuristic default (0.50 ms
vs. 0.71 ms) — on par with cuBLAS for this shape.</p> <!> <p></p>`,1);function q(g,I){N(I,!1),Q(()=>{new URLSearchParams(window.location.search).get("fw")}),C();var l=V();b("w9xnlz",J=>{var T=G();B(T,"content",E),w(J,T)});var t=s(S(l),2);k(t,{containerStyle:"float: right; margin-left: 10px; display: inline-flex; position: relative; z-index: 10;"});var e=s(t,2);n(e,{title:"Ship Triton autotune configurations",local:"ship-triton-autotune-configurations",headingTag:"h1"});var o=s(e,10);n(o,{title:"Shipping data files with a kernel",local:"shipping-data-files-with-a-kernel",headingTag:"h2"});var i=s(o,4);a(i,{code:"JTVCZ2VuZXJhbCU1RCUwQW5hbWUlMjAlM0QlMjAlMjJnZW1tLXRyaXRvbi1hdXRvdHVuZSUyMiUwQXZlcnNpb24lMjAlM0QlMjAxJTBBZWRpdGlvbiUyMCUzRCUyMDUlMEFsaWNlbnNlJTIwJTNEJTIwJTIyQXBhY2hlLTIuMCUyMiUwQWJhY2tlbmRzJTIwJTNEJTIwJTVCJTIyY3VkYSUyMiUyQyUyMCUyMnJvY20lMjIlMkMlMjAlMjJ4cHUlMjIlNUQlMEElMEElNUJnZW5lcmFsLmh1YiU1RCUwQXJlcG8taWQlMjAlM0QlMjAlMjJrZXJuZWxzLXRlc3QlMkZnZW1tLXRyaXRvbi1hdXRvdHVuZSUyMiUwQSUwQSU1QnRvcmNoLW5vYXJjaCU1RCUwQXB5ZXh0JTIwJTNEJTIwJTVCJTIyanNvbiUyMiUyQyUyMCUyMnB5JTIyJTVE",highlighted:`<span class="hljs-section">[general]</span>
<span class="hljs-attr">name</span> = <span class="hljs-string">&quot;gemm-triton-autotune&quot;</span>
<span class="hljs-attr">version</span> = <span class="hljs-number">1</span>
<span class="hljs-attr">edition</span> = <span class="hljs-number">5</span>
<span class="hljs-attr">license</span> = <span class="hljs-string">&quot;Apache-2.0&quot;</span>
<span class="hljs-attr">backends</span> = [<span class="hljs-string">&quot;cuda&quot;</span>, <span class="hljs-string">&quot;rocm&quot;</span>, <span class="hljs-string">&quot;xpu&quot;</span>]

<span class="hljs-section">[general.hub]</span>
<span class="hljs-attr">repo-id</span> = <span class="hljs-string">&quot;kernels-test/gemm-triton-autotune&quot;</span>

<span class="hljs-section">[torch-noarch]</span>
<span class="hljs-attr">pyext</span> = [<span class="hljs-string">&quot;json&quot;</span>, <span class="hljs-string">&quot;py&quot;</span>]`,lang:"toml",wrap:!1});var p=s(i,2);n(p,{title:"Configuration file layout",local:"configuration-file-layout",headingTag:"h2"});var c=s(p,4);a(c,{code:"Y29uZmlncyUyRk4lM0Q0MDk2JTJDSyUzRDQwOTYlMkNkZXZpY2VfbmFtZSUzRE5WSURJQV9MNC5qc29uJTBBY29uZmlncyUyRk4lM0QxNDMzNiUyQ0slM0Q0MDk2JTJDZGV2aWNlX25hbWUlM0ROVklESUFfTDQuanNvbg==",highlighted:`<span class="hljs-attribute">configs</span>/N=<span class="hljs-number">4096</span>,K=<span class="hljs-number">4096</span>,device_name=NVIDIA_L4.json
<span class="hljs-attribute">configs</span>/N=<span class="hljs-number">14336</span>,K=<span class="hljs-number">4096</span>,device_name=NVIDIA_L4.json`,lang:"",wrap:!1});var M=s(c,4);a(M,{code:"JTdCJTBBJTIwJTIwJTIwJTIwJTIyMSUyMiUzQSUyMCU3QiUwQSUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMkJMT0NLX1NJWkVfTSUyMiUzQSUyMDE2JTJDJTBBJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIyQkxPQ0tfU0laRV9OJTIyJTNBJTIwMTI4JTJDJTBBJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIyQkxPQ0tfU0laRV9LJTIyJTNBJTIwMzIlMkMlMEElMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjJHUk9VUF9TSVpFX00lMjIlM0ElMjA4JTJDJTBBJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIybnVtX3dhcnBzJTIyJTNBJTIwNCUyQyUwQSUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMm51bV9zdGFnZXMlMjIlM0ElMjAzJTBBJTIwJTIwJTIwJTIwJTdEJTJDJTBBJTIwJTIwJTIwJTIwJTIyMTAyNCUyMiUzQSUyMCU3QiUwQSUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMkJMT0NLX1NJWkVfTSUyMiUzQSUyMDEyOCUyQyUwQSUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMkJMT0NLX1NJWkVfTiUyMiUzQSUyMDEyOCUyQyUwQSUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMkJMT0NLX1NJWkVfSyUyMiUzQSUyMDMyJTJDJTBBJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIyR1JPVVBfU0laRV9NJTIyJTNBJTIwOCUyQyUwQSUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMm51bV93YXJwcyUyMiUzQSUyMDQlMkMlMEElMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjJudW1fc3RhZ2VzJTIyJTNBJTIwMyUwQSUyMCUyMCUyMCUyMCU3RCUwQSU3RA==",highlighted:`<span class="hljs-punctuation">{</span>
    <span class="hljs-attr">&quot;1&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-punctuation">{</span>
        <span class="hljs-attr">&quot;BLOCK_SIZE_M&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">16</span><span class="hljs-punctuation">,</span>
        <span class="hljs-attr">&quot;BLOCK_SIZE_N&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">128</span><span class="hljs-punctuation">,</span>
        <span class="hljs-attr">&quot;BLOCK_SIZE_K&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">32</span><span class="hljs-punctuation">,</span>
        <span class="hljs-attr">&quot;GROUP_SIZE_M&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">8</span><span class="hljs-punctuation">,</span>
        <span class="hljs-attr">&quot;num_warps&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">4</span><span class="hljs-punctuation">,</span>
        <span class="hljs-attr">&quot;num_stages&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">3</span>
    <span class="hljs-punctuation">}</span><span class="hljs-punctuation">,</span>
    <span class="hljs-attr">&quot;1024&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-punctuation">{</span>
        <span class="hljs-attr">&quot;BLOCK_SIZE_M&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">128</span><span class="hljs-punctuation">,</span>
        <span class="hljs-attr">&quot;BLOCK_SIZE_N&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">128</span><span class="hljs-punctuation">,</span>
        <span class="hljs-attr">&quot;BLOCK_SIZE_K&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">32</span><span class="hljs-punctuation">,</span>
        <span class="hljs-attr">&quot;GROUP_SIZE_M&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">8</span><span class="hljs-punctuation">,</span>
        <span class="hljs-attr">&quot;num_warps&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">4</span><span class="hljs-punctuation">,</span>
        <span class="hljs-attr">&quot;num_stages&quot;</span><span class="hljs-punctuation">:</span> <span class="hljs-number">3</span>
    <span class="hljs-punctuation">}</span>
<span class="hljs-punctuation">}</span>`,lang:"json",wrap:!1});var h=s(M,4);n(h,{title:"Looking up configurations at runtime",local:"looking-up-configurations-at-runtime",headingTag:"h2"});var r=s(h,4);a(r,{code:"JTQwZnVuY3Rvb2xzLmNhY2hlJTBBZGVmJTIwX2xvYWRfdHVuZWRfY29uZmlncyhOJTNBJTIwaW50JTJDJTIwSyUzQSUyMGludCklMjAtJTNFJTIwT3B0aW9uYWwlNUJEaWN0JTVCaW50JTJDJTIwRGljdCU1QnN0ciUyQyUyMGludCU1RCU1RCU1RCUzQSUwQSUyMCUyMCUyMCUyMHBhdGglMjAlM0QlMjBfQ09ORklHU19ESVIlMjAlMkYlMjBjb25maWdfZmlsZV9uYW1lKE4lMkMlMjBLKSUwQSUyMCUyMCUyMCUyMGlmJTIwcGF0aC5leGlzdHMoKSUzQSUwQSUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMGxvZ2dlci5pbmZvKCUyMlVzaW5nJTIwdHVuZWQlMjBHRU1NJTIwY29uZmlndXJhdGlvbnMlMjBmcm9tJTIwJTI1cy4lMjIlMkMlMjBwYXRoKSUwQSUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMHdpdGglMjBvcGVuKHBhdGgpJTIwYXMlMjBmJTNBJTBBJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwcmV0dXJuJTIwJTdCaW50KG0pJTNBJTIwY29uZmlnJTIwZm9yJTIwbSUyQyUyMGNvbmZpZyUyMGluJTIwanNvbi5sb2FkKGYpLml0ZW1zKCklN0QlMEElMjAlMjAlMjAlMjBsb2dnZXIud2FybmluZyglMEElMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjJObyUyMHR1bmVkJTIwR0VNTSUyMGNvbmZpZ3VyYXRpb24lMjBmb3VuZCUyMGZvciUyMHRoaXMlMjBkZXZpY2UlMjBhbmQlMjBzaGFwZSUyMCglMjVzKS4lMjAlMjIlMEElMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjJGYWxsaW5nJTIwYmFjayUyMHRvJTIwaGV1cmlzdGljJTIwZGVmYXVsdHMlMkMlMjBwZXJmb3JtYW5jZSUyMG1heSUyMGJlJTIwc3Vib3B0aW1hbC4lMjAlMjIlMEElMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjJHZW5lcmF0ZSUyMGElMjBjb25maWd1cmF0aW9uJTIwd2l0aCUyMCU2MHR1bmVfZ2VtbShOJTNEJTI1ZCUyQyUyMEslM0QlMjVkKSU2MC4lMjIlMkMlMEElMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjBwYXRoLm5hbWUlMkMlMEElMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjBOJTJDJTBBJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwSyUyQyUwQSUyMCUyMCUyMCUyMCklMEElMjAlMjAlMjAlMjByZXR1cm4lMjBOb25lJTBBJTBBJTBBZGVmJTIwZ2V0X2NvbmZpZyhNJTNBJTIwaW50JTJDJTIwTiUzQSUyMGludCUyQyUyMEslM0ElMjBpbnQpJTIwLSUzRSUyMERpY3QlNUJzdHIlMkMlMjBpbnQlNUQlM0ElMEElMjAlMjAlMjAlMjB0dW5lZCUyMCUzRCUyMF9sb2FkX3R1bmVkX2NvbmZpZ3MoTiUyQyUyMEspJTBBJTIwJTIwJTIwJTIwaWYlMjB0dW5lZCUyMGlzJTIwbm90JTIwTm9uZSUzQSUwQSUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMCUyMyUyMFR1bmVkJTIwTXMlMjBhcmUlMjBzcGFjZWQlMjBsb2dhcml0aG1pY2FsbHklMkMlMjBzbyUyMHBpY2slMjB0aGUlMjBuZWFyZXN0JTIwaW4lMjBsb2clMjBzcGFjZS4lMEElMjAlMjAlMjAlMjAlMjAlMjAlMjAlMjBuZWFyZXN0X20lMjAlM0QlMjBtaW4odHVuZWQlMkMlMjBrZXklM0RsYW1iZGElMjBtJTNBJTIwYWJzKG1hdGgubG9nKE0lMjAlMkYlMjBtKSkpJTBBJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwcmV0dXJuJTIwdHVuZWQlNUJuZWFyZXN0X20lNUQlMEElMjAlMjAlMjAlMjByZXR1cm4lMjBkZWZhdWx0X2NvbmZpZyhNJTJDJTIwTiUyQyUyMEsp",highlighted:`<span class="hljs-meta">@functools.cache</span>
<span class="hljs-keyword">def</span> <span class="hljs-title function_">_load_tuned_configs</span>(<span class="hljs-params">N: <span class="hljs-built_in">int</span>, K: <span class="hljs-built_in">int</span></span>) -&gt; <span class="hljs-type">Optional</span>[<span class="hljs-type">Dict</span>[<span class="hljs-built_in">int</span>, <span class="hljs-type">Dict</span>[<span class="hljs-built_in">str</span>, <span class="hljs-built_in">int</span>]]]:
    path = _CONFIGS_DIR / config_file_name(N, K)
    <span class="hljs-keyword">if</span> path.exists():
        logger.info(<span class="hljs-string">&quot;Using tuned GEMM configurations from %s.&quot;</span>, path)
        <span class="hljs-keyword">with</span> <span class="hljs-built_in">open</span>(path) <span class="hljs-keyword">as</span> f:
            <span class="hljs-keyword">return</span> {<span class="hljs-built_in">int</span>(m): config <span class="hljs-keyword">for</span> m, config <span class="hljs-keyword">in</span> json.load(f).items()}
    logger.warning(
        <span class="hljs-string">&quot;No tuned GEMM configuration found for this device and shape (%s). &quot;</span>
        <span class="hljs-string">&quot;Falling back to heuristic defaults, performance may be suboptimal. &quot;</span>
        <span class="hljs-string">&quot;Generate a configuration with \`tune_gemm(N=%d, K=%d)\`.&quot;</span>,
        path.name,
        N,
        K,
    )
    <span class="hljs-keyword">return</span> <span class="hljs-literal">None</span>


<span class="hljs-keyword">def</span> <span class="hljs-title function_">get_config</span>(<span class="hljs-params">M: <span class="hljs-built_in">int</span>, N: <span class="hljs-built_in">int</span>, K: <span class="hljs-built_in">int</span></span>) -&gt; <span class="hljs-type">Dict</span>[<span class="hljs-built_in">str</span>, <span class="hljs-built_in">int</span>]:
    tuned = _load_tuned_configs(N, K)
    <span class="hljs-keyword">if</span> tuned <span class="hljs-keyword">is</span> <span class="hljs-keyword">not</span> <span class="hljs-literal">None</span>:
        <span class="hljs-comment"># Tuned Ms are spaced logarithmically, so pick the nearest in log space.</span>
        nearest_m = <span class="hljs-built_in">min</span>(tuned, key=<span class="hljs-keyword">lambda</span> m: <span class="hljs-built_in">abs</span>(math.log(M / m)))
        <span class="hljs-keyword">return</span> tuned[nearest_m]
    <span class="hljs-keyword">return</span> default_config(M, N, K)`,lang:"python",wrap:!1});var u=s(r,4);a(u,{code:"ZGVmJTIwbGF1bmNoX2dlbW1fa2VybmVsKGElMkMlMjBiJTJDJTIwb3V0JTJDJTIwY29uZmlnKSUzQSUwQSUyMCUyMCUyMCUyME0lMkMlMjBLJTIwJTNEJTIwYS5zaGFwZSUwQSUyMCUyMCUyMCUyME4lMjAlM0QlMjBiLnNoYXBlJTVCMSU1RCUwQSUyMCUyMCUyMCUyMGdyaWQlMjAlM0QlMjAoJTBBJTIwJTIwJTIwJTIwJTIwJTIwJTIwJTIwdHJpdG9uLmNkaXYoTSUyQyUyMGNvbmZpZyU1QiUyMkJMT0NLX1NJWkVfTSUyMiU1RCklMjAqJTIwdHJpdG9uLmNkaXYoTiUyQyUyMGNvbmZpZyU1QiUyMkJMT0NLX1NJWkVfTiUyMiU1RCklMkMlMEElMjAlMjAlMjAlMjApJTBBJTIwJTIwJTIwJTIwX2dlbW1fa2VybmVsJTVCZ3JpZCU1RChhJTJDJTIwYiUyQyUyMG91dCUyQyUyME0lMkMlMjBOJTJDJTIwSyUyQyUyMC4uLiUyQyUyMCoqY29uZmlnKQ==",highlighted:`<span class="hljs-keyword">def</span> <span class="hljs-title function_">launch_gemm_kernel</span>(<span class="hljs-params">a, b, out, config</span>):
    M, K = a.shape
    N = b.shape[<span class="hljs-number">1</span>]
    grid = (
        triton.cdiv(M, config[<span class="hljs-string">&quot;BLOCK_SIZE_M&quot;</span>]) * triton.cdiv(N, config[<span class="hljs-string">&quot;BLOCK_SIZE_N&quot;</span>]),
    )
    _gemm_kernel[grid](a, b, out, M, N, K, ..., **config)`,lang:"python",wrap:!1});var y=s(u,2);n(y,{title:"Generating the configurations",local:"generating-the-configurations",headingTag:"h2"});var d=s(y,6);a(d,{code:"JTI0JTIwcHl0aG9uJTIwdHVuZS5weSUyMC0tbiUyMDQwOTYlMjAtLWslMjA0MDk2JTBBJTI0JTIwcHl0aG9uJTIwdHVuZS5weSUyMC0tbiUyMDE0MzM2JTIwLS1rJTIwNDA5Ng==",highlighted:`$ python tune.py --n 4096 --k 4096
$ python tune.py --n 14336 --k 4096`,lang:"bash",wrap:!1});var j=s(d,4);a(j,{code:"JTI0JTIwTE9DQUxfS0VSTkVMUyUzRGtlcm5lbHMtdGVzdCUyRmdlbW0tdHJpdG9uLWF1dG90dW5lJTNEYnVpbGQlMjBweXRob24lMjB0dW5lLnB5JTIwLS1uJTIwNDA5NiUyMC0tayUyMDQwOTY=",highlighted:"$ LOCAL_KERNELS=kernels-test/gemm-triton-autotune=build python tune.py --n 4096 --k 4096",lang:"bash",wrap:!1});var U=s(j,2);n(U,{title:"Impact",local:"impact",headingTag:"h2"});var f=s(U,4);A(f,{source:"https://github.com/huggingface/kernels/blob/main/docs/source/builder/triton-autotune.md"}),_(2),w(g,l),Z()}export{q as component};