Update docs

Author: github-actions[bot]

_sources/autoapi/tilelang/intrinsics/wgmma_macro_generator/index.rst.txt

@@ -84,6 +84,14 @@ Module Contents
       :type:  str
 
 
+   .. py:attribute:: wgmma_inst_m
+      :type:  int
+
+
+   .. py:attribute:: wgmma_inst_n
+      :type:  int
+
+
    .. py:attribute:: a_shared_layout
       :type:  tilelang.layout.Layout
       :value: None
@@ -96,10 +104,10 @@ Module Contents
 
 
 
-   .. py:method:: wgmma(A_buf, B_buf, C_local_buf, clear_accum = False)
+   .. py:method:: wgmma(A_buf, B_buf, C_local_buf, clear_accum = False, wg_wait = 0)
 
 
-   .. py:method:: wgmma_rs(A_buf, B_buf, C_local_buf, clear_accum = False)
+   .. py:method:: wgmma_rs(A_buf, B_buf, C_local_buf, clear_accum = False, wg_wait = 0)
 
 
    .. py:method:: make_mma_load_layout(local_buf, matrix = 'A')

autoapi/tilelang/intrinsics/wgmma_macro_generator/index.html

@@ -642,6 +642,16 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <span class="sig-name descname"><span class="pre">wgmma_prefix</span></span><em class="property"><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="pre">str</span></em><a class="headerlink" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_prefix" title="Link to this definition">¶</a></dt>
 <dd></dd></dl>
 
+<dl class="py attribute">
+<dt class="sig sig-object py" id="tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_inst_m">
+<span class="sig-name descname"><span class="pre">wgmma_inst_m</span></span><em class="property"><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="pre">int</span></em><a class="headerlink" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_inst_m" title="Link to this definition">¶</a></dt>
+<dd></dd></dl>
+
+<dl class="py attribute">
+<dt class="sig sig-object py" id="tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_inst_n">
+<span class="sig-name descname"><span class="pre">wgmma_inst_n</span></span><em class="property"><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="pre">int</span></em><a class="headerlink" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_inst_n" title="Link to this definition">¶</a></dt>
+<dd></dd></dl>
+
 <dl class="py attribute">
 <dt class="sig sig-object py" id="tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.a_shared_layout">
 <span class="sig-name descname"><span class="pre">a_shared_layout</span></span><em class="property"><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="pre">tilelang.layout.Layout</span></em><em class="property"><span class="w"> </span><span class="p"><span class="pre">=</span></span><span class="w"> </span><span class="pre">None</span></em><a class="headerlink" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.a_shared_layout" title="Link to this definition">¶</a></dt>
@@ -654,29 +664,31 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 
 <dl class="py method">
 <dt class="sig sig-object py" id="tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma">
-<span class="sig-name descname"><span class="pre">wgmma</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">A_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">B_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">C_local_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">clear_accum</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">False</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma" title="Link to this definition">¶</a></dt>
+<span class="sig-name descname"><span class="pre">wgmma</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">A_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">B_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">C_local_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">clear_accum</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">False</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">wg_wait</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">0</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma" title="Link to this definition">¶</a></dt>
 <dd><dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
 <dd class="field-odd"><ul class="simple">
 <li><p><strong>A_buf</strong> (<em>tvm.tir.Buffer</em>)</p></li>
 <li><p><strong>B_buf</strong> (<em>tvm.tir.Buffer</em>)</p></li>
 <li><p><strong>C_local_buf</strong> (<em>tvm.tir.Buffer</em>)</p></li>
 <li><p><strong>clear_accum</strong> (<em>tvm.tir.PrimExpr</em>)</p></li>
+<li><p><strong>wg_wait</strong> (<em>int</em>)</p></li>
 </ul>
 </dd>
 </dl>
 </dd></dl>
 
 <dl class="py method">
 <dt class="sig sig-object py" id="tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_rs">
-<span class="sig-name descname"><span class="pre">wgmma_rs</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">A_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">B_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">C_local_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">clear_accum</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">False</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_rs" title="Link to this definition">¶</a></dt>
+<span class="sig-name descname"><span class="pre">wgmma_rs</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">A_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">B_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">C_local_buf</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">clear_accum</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">False</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">wg_wait</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">0</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_rs" title="Link to this definition">¶</a></dt>
 <dd><dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
 <dd class="field-odd"><ul class="simple">
 <li><p><strong>A_buf</strong> (<em>tvm.tir.Buffer</em>)</p></li>
 <li><p><strong>B_buf</strong> (<em>tvm.tir.Buffer</em>)</p></li>
 <li><p><strong>C_local_buf</strong> (<em>tvm.tir.Buffer</em>)</p></li>
 <li><p><strong>clear_accum</strong> (<em>tvm.tir.PrimExpr</em>)</p></li>
+<li><p><strong>wg_wait</strong> (<em>int</em>)</p></li>
 </ul>
 </dd>
 </dl>
@@ -811,6 +823,8 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 </li>
 <li><a class="reference internal" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter"><code class="docutils literal notranslate"><span class="pre">TensorCoreIntrinEmitter</span></code></a><ul>
 <li><a class="reference internal" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_prefix"><code class="docutils literal notranslate"><span class="pre">TensorCoreIntrinEmitter.wgmma_prefix</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_inst_m"><code class="docutils literal notranslate"><span class="pre">TensorCoreIntrinEmitter.wgmma_inst_m</span></code></a></li>
+<li><a class="reference internal" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_inst_n"><code class="docutils literal notranslate"><span class="pre">TensorCoreIntrinEmitter.wgmma_inst_n</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.a_shared_layout"><code class="docutils literal notranslate"><span class="pre">TensorCoreIntrinEmitter.a_shared_layout</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.b_shared_layout"><code class="docutils literal notranslate"><span class="pre">TensorCoreIntrinEmitter.b_shared_layout</span></code></a></li>
 <li><a class="reference internal" href="#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma"><code class="docutils literal notranslate"><span class="pre">TensorCoreIntrinEmitter.wgmma()</span></code></a></li>

genindex.html

@@ -7131,10 +7131,10 @@ <h2>W</h2>
 </li>
         <li><a href="autoapi/tilelang/language/warpgroup/index.html#tilelang.language.warpgroup.WarpSpecialize">WarpSpecialize() (in module tilelang.language.warpgroup)</a>
 </li>
-    </ul></td>
-    <td style="width: 33%; vertical-align: top;"><ul>
         <li><a href="autoapi/tilelang/transform/index.html#tilelang.transform.WarpSpecialized">WarpSpecialized() (in module tilelang.transform)</a>
 </li>
+    </ul></td>
+    <td style="width: 33%; vertical-align: top;"><ul>
         <li><a href="autoapi/tilelang/transform/index.html#tilelang.transform.WarpSpecializedPipeline">WarpSpecializedPipeline() (in module tilelang.transform)</a>
 </li>
         <li><a href="autoapi/tilelang/language/warpgroup/index.html#tilelang.language.warpgroup.WarpSpecializeFrame">WarpSpecializeFrame (class in tilelang.language.warpgroup)</a>
@@ -7150,6 +7150,10 @@ <h2>W</h2>
         <li><a href="autoapi/tilelang/tileop/gemm/index.html#tilelang.tileop.gemm.GemmInst.WGMMA">WGMMA (tilelang.tileop.gemm.GemmInst attribute)</a>
 </li>
         <li><a href="autoapi/tilelang/intrinsics/wgmma_macro_generator/index.html#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma">wgmma() (tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter method)</a>
+</li>
+        <li><a href="autoapi/tilelang/intrinsics/wgmma_macro_generator/index.html#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_inst_m">wgmma_inst_m (tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter attribute)</a>
+</li>
+        <li><a href="autoapi/tilelang/intrinsics/wgmma_macro_generator/index.html#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_inst_n">wgmma_inst_n (tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter attribute)</a>
 </li>
         <li><a href="autoapi/tilelang/intrinsics/wgmma_macro_generator/index.html#tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter.wgmma_prefix">wgmma_prefix (tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter attribute)</a>
 </li>