Skip to content

Commit ccbf299

Browse files
github-actions[bot]github-actions[bot]
committed
Update docs
1 parent f64d1eb commit ccbf299

File tree

5 files changed

+90
-4
lines changed

5 files changed

+90
-4
lines changed

‎_sources/autoapi/tilelang/jit/adapter/wrapper/index.rst.txt‎

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,7 @@ Attributes
1818
tilelang.jit.adapter.wrapper.L2_PERSISTENT_MAP_INIT_FUNC
1919
tilelang.jit.adapter.wrapper.L2_PERSISTENT_MAP_RESET_HANDLE
2020
tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC
21+
tilelang.jit.adapter.wrapper.TMA_IM2COL_DESC_INIT_FUNC
2122
tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC_PY
2223
tilelang.jit.adapter.wrapper.KERNEL_LAUNCH_FUNC_PY
2324
tilelang.jit.adapter.wrapper.logger
@@ -272,6 +273,50 @@ Module Contents
272273

273274

274275

276+
.. py:data:: TMA_IM2COL_DESC_INIT_FUNC
277+
:value: Multiline-String
278+
279+
.. raw:: html
280+
281+
<details><summary>Show Value</summary>
282+
283+
.. code-block:: python
284+
285+
"""
286+
CUtensorMap {0};
287+
CUtensorMapDataType {0}_type= (CUtensorMapDataType){1};
288+
cuuint32_t {0}_tensorRank= {2};
289+
void *{0}_globalAddress= {3};
290+
cuuint64_t {0}_globalDim[{2}]= {{{4}}};
291+
cuuint64_t {0}_globalStride[{2}]= {{{5}}};
292+
cuuint32_t {0}_elementStrides[{2}]= {{{6}}};
293+
int {0}_lowerCorner[{2} - 2]= {{{7}}};
294+
int {0}_upperCorner[{2} - 2]= {{{8}}};
295+
cuuint32_t {0}_channelsPerPixel= {9};
296+
cuuint32_t {0}_pixelsPerColumn= {10};
297+
CUtensorMapInterleave {0}_interleave= (CUtensorMapInterleave){11};
298+
CUtensorMapSwizzle {0}_swizzle= (CUtensorMapSwizzle){12};
299+
CUtensorMapL2promotion {0}_l2Promotion= (CUtensorMapL2promotion){13};
300+
CUtensorMapFloatOOBfill {0}_oobFill= (CUtensorMapFloatOOBfill){14};
301+
302+
CUresult {0}_result = CUTLASS_CUDA_DRIVER_WRAPPER_CALL(cuTensorMapEncodeIm2col)(
303+
&{0}, {0}_type, {0}_tensorRank, {0}_globalAddress, {0}_globalDim, {0}_globalStride + 1,
304+
{0}_lowerCorner, {0}_upperCorner, {0}_channelsPerPixel, {0}_pixelsPerColumn, {0}_elementStrides, {0}_interleave, {0}_swizzle, {0}_l2Promotion, {0}_oobFill);
305+
306+
if ({0}_result != CUDA_SUCCESS) {{
307+
std::stringstream ss;
308+
ss << "Error: Failed to initialize the TMA descriptor {0}";
309+
snprintf(error_buf, ERROR_BUF_SIZE, "%s", ss.str().c_str());
310+
return -1;
311+
}}
312+
"""
313+
314+
.. raw:: html
315+
316+
</details>
317+
318+
319+
275320
.. py:data:: TMA_DESC_INIT_FUNC_PY
276321
:value: Multiline-String
277322

‎autoapi/tilelang/jit/adapter/wrapper/index.html‎

Lines changed: 42 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -513,13 +513,16 @@ <h2>Attributes<a class="headerlink" href="#attributes" title="Link to this headi
513513
<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC" title="tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC"><code class="xref py py-obj docutils literal notranslate"><span class="pre">TMA_DESC_INIT_FUNC</span></code></a></p></td>
514514
<td><p></p></td>
515515
</tr>
516-
<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC_PY" title="tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC_PY"><code class="xref py py-obj docutils literal notranslate"><span class="pre">TMA_DESC_INIT_FUNC_PY</span></code></a></p></td>
516+
<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TMA_IM2COL_DESC_INIT_FUNC" title="tilelang.jit.adapter.wrapper.TMA_IM2COL_DESC_INIT_FUNC"><code class="xref py py-obj docutils literal notranslate"><span class="pre">TMA_IM2COL_DESC_INIT_FUNC</span></code></a></p></td>
517517
<td><p></p></td>
518518
</tr>
519-
<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.jit.adapter.wrapper.KERNEL_LAUNCH_FUNC_PY" title="tilelang.jit.adapter.wrapper.KERNEL_LAUNCH_FUNC_PY"><code class="xref py py-obj docutils literal notranslate"><span class="pre">KERNEL_LAUNCH_FUNC_PY</span></code></a></p></td>
519+
<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC_PY" title="tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC_PY"><code class="xref py py-obj docutils literal notranslate"><span class="pre">TMA_DESC_INIT_FUNC_PY</span></code></a></p></td>
520520
<td><p></p></td>
521521
</tr>
522-
<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.jit.adapter.wrapper.logger" title="tilelang.jit.adapter.wrapper.logger"><code class="xref py py-obj docutils literal notranslate"><span class="pre">logger</span></code></a></p></td>
522+
<tr class="row-even"><td><p><a class="reference internal" href="#tilelang.jit.adapter.wrapper.KERNEL_LAUNCH_FUNC_PY" title="tilelang.jit.adapter.wrapper.KERNEL_LAUNCH_FUNC_PY"><code class="xref py py-obj docutils literal notranslate"><span class="pre">KERNEL_LAUNCH_FUNC_PY</span></code></a></p></td>
523+
<td><p></p></td>
524+
</tr>
525+
<tr class="row-odd"><td><p><a class="reference internal" href="#tilelang.jit.adapter.wrapper.logger" title="tilelang.jit.adapter.wrapper.logger"><code class="xref py py-obj docutils literal notranslate"><span class="pre">logger</span></code></a></p></td>
523526
<td><p></p></td>
524527
</tr>
525528
</tbody>
@@ -711,6 +714,41 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
711714
</div>
712715
</details></dd></dl>
713716

717+
<dl class="py data">
718+
<dt class="sig sig-object py" id="tilelang.jit.adapter.wrapper.TMA_IM2COL_DESC_INIT_FUNC">
719+
<span class="sig-prename descclassname"><span class="pre">tilelang.jit.adapter.wrapper.</span></span><span class="sig-name descname"><span class="pre">TMA_IM2COL_DESC_INIT_FUNC</span></span><em class="property"><span class="w"> </span><span class="p"><span class="pre">=</span></span><span class="w"> </span><span class="pre">Multiline-String</span></em><a class="headerlink" href="#tilelang.jit.adapter.wrapper.TMA_IM2COL_DESC_INIT_FUNC" title="Link to this definition"></a></dt>
720+
<dd><details><summary>Show Value</summary><div class="highlight-python notranslate"><div class="highlight"><pre><span></span><span class="sd">&quot;&quot;&quot;</span>
721+
<span class="sd"> CUtensorMap {0};</span>
722+
<span class="sd"> CUtensorMapDataType {0}_type= (CUtensorMapDataType){1};</span>
723+
<span class="sd"> cuuint32_t {0}_tensorRank= {2};</span>
724+
<span class="sd"> void *{0}_globalAddress= {3};</span>
725+
<span class="sd"> cuuint64_t {0}_globalDim[{2}]= {{{4}}};</span>
726+
<span class="sd"> cuuint64_t {0}_globalStride[{2}]= {{{5}}};</span>
727+
<span class="sd"> cuuint32_t {0}_elementStrides[{2}]= {{{6}}};</span>
728+
<span class="sd"> int {0}_lowerCorner[{2} - 2]= {{{7}}};</span>
729+
<span class="sd"> int {0}_upperCorner[{2} - 2]= {{{8}}};</span>
730+
<span class="sd"> cuuint32_t {0}_channelsPerPixel= {9};</span>
731+
<span class="sd"> cuuint32_t {0}_pixelsPerColumn= {10};</span>
732+
<span class="sd"> CUtensorMapInterleave {0}_interleave= (CUtensorMapInterleave){11};</span>
733+
<span class="sd"> CUtensorMapSwizzle {0}_swizzle= (CUtensorMapSwizzle){12};</span>
734+
<span class="sd"> CUtensorMapL2promotion {0}_l2Promotion= (CUtensorMapL2promotion){13};</span>
735+
<span class="sd"> CUtensorMapFloatOOBfill {0}_oobFill= (CUtensorMapFloatOOBfill){14};</span>
736+
737+
<span class="sd"> CUresult {0}_result = CUTLASS_CUDA_DRIVER_WRAPPER_CALL(cuTensorMapEncodeIm2col)(</span>
738+
<span class="sd"> &amp;{0}, {0}_type, {0}_tensorRank, {0}_globalAddress, {0}_globalDim, {0}_globalStride + 1,</span>
739+
<span class="sd"> {0}_lowerCorner, {0}_upperCorner, {0}_channelsPerPixel, {0}_pixelsPerColumn, {0}_elementStrides, {0}_interleave, {0}_swizzle, {0}_l2Promotion, {0}_oobFill);</span>
740+
741+
<span class="sd"> if ({0}_result != CUDA_SUCCESS) {{</span>
742+
<span class="sd"> std::stringstream ss;</span>
743+
<span class="sd"> ss &lt;&lt; &quot;Error: Failed to initialize the TMA descriptor {0}&quot;;</span>
744+
<span class="sd"> snprintf(error_buf, ERROR_BUF_SIZE, &quot;%s&quot;, ss.str().c_str());</span>
745+
<span class="sd"> return -1;</span>
746+
<span class="sd"> }}</span>
747+
<span class="sd">&quot;&quot;&quot;</span>
748+
</pre></div>
749+
</div>
750+
</details></dd></dl>
751+
714752
<dl class="py data">
715753
<dt class="sig sig-object py" id="tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC_PY">
716754
<span class="sig-prename descclassname"><span class="pre">tilelang.jit.adapter.wrapper.</span></span><span class="sig-name descname"><span class="pre">TMA_DESC_INIT_FUNC_PY</span></span><em class="property"><span class="w"> </span><span class="p"><span class="pre">=</span></span><span class="w"> </span><span class="pre">Multiline-String</span></em><a class="headerlink" href="#tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC_PY" title="Link to this definition"></a></dt>
@@ -1457,6 +1495,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
14571495
<li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.L2_PERSISTENT_MAP_INIT_FUNC"><code class="docutils literal notranslate"><span class="pre">L2_PERSISTENT_MAP_INIT_FUNC</span></code></a></li>
14581496
<li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.L2_PERSISTENT_MAP_RESET_HANDLE"><code class="docutils literal notranslate"><span class="pre">L2_PERSISTENT_MAP_RESET_HANDLE</span></code></a></li>
14591497
<li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC"><code class="docutils literal notranslate"><span class="pre">TMA_DESC_INIT_FUNC</span></code></a></li>
1498+
<li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TMA_IM2COL_DESC_INIT_FUNC"><code class="docutils literal notranslate"><span class="pre">TMA_IM2COL_DESC_INIT_FUNC</span></code></a></li>
14601499
<li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC_PY"><code class="docutils literal notranslate"><span class="pre">TMA_DESC_INIT_FUNC_PY</span></code></a></li>
14611500
<li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.KERNEL_LAUNCH_FUNC_PY"><code class="docutils literal notranslate"><span class="pre">KERNEL_LAUNCH_FUNC_PY</span></code></a></li>
14621501
<li><a class="reference internal" href="#tilelang.jit.adapter.wrapper.BaseWrapper"><code class="docutils literal notranslate"><span class="pre">BaseWrapper</span></code></a><ul>

‎genindex.html‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6050,6 +6050,8 @@ <h2>T</h2>
60506050
<li><a href="autoapi/tilelang/jit/adapter/wrapper/index.html#tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC_PY">TMA_DESC_INIT_FUNC_PY (in module tilelang.jit.adapter.wrapper)</a>
60516051
</li>
60526052
<li><a href="autoapi/tilelang/jit/adapter/wrapper/index.html#tilelang.jit.adapter.wrapper.TLCUDASourceWrapper.tma_descriptor_args">tma_descriptor_args (tilelang.jit.adapter.wrapper.TLCUDASourceWrapper attribute)</a>
6053+
</li>
6054+
<li><a href="autoapi/tilelang/jit/adapter/wrapper/index.html#tilelang.jit.adapter.wrapper.TMA_IM2COL_DESC_INIT_FUNC">TMA_IM2COL_DESC_INIT_FUNC (in module tilelang.jit.adapter.wrapper)</a>
60536055
</li>
60546056
<li><a href="autoapi/tilelang/language/builtin/index.html#tilelang.language.builtin.tma_load">tma_load() (in module tilelang.language.builtin)</a>
60556057
</li>

‎objects.inv‎

11 Bytes
Binary file not shown.

‎searchindex.js‎

Lines changed: 1 addition & 1 deletion
Some generated files are not rendered by default. Learn more about customizing how changed files appear on GitHub.

0 commit comments

Comments
 (0)