Update docs

Author: GitHub Actions

_sources/deeplearning_operators/elementwise.md.txt

@@ -27,7 +27,7 @@ Please note that this tutorial does not delve deeply into the design principles
 def elementwise_add(N, threads=256, dtype="bfloat16"):
 
     @T.prim_func
-    def main(A: T.Buffer((N), dtype), B: T.Buffer((N), dtype), C: T.Buffer((N), dtype)):
+    def main(A: T.Tensor((N), dtype), B: T.Tensor((N), dtype), C: T.Tensor((N), dtype)):
         with T.Kernel(T.ceildiv(N, threads), threads=threads) as (b_x):
             # vector add.
             for i in T.Parallel(threads):
@@ -67,9 +67,9 @@ def elementwise_add(
 ):
     @T.prim_func
     def main(
-            A: T.Buffer((M, N), in_dtype),
-            B: T.Buffer((M, N), in_dtype),
-            C: T.Buffer((M, N), out_dtype),
+            A: T.Tensor((M, N), in_dtype),
+            B: T.Tensor((M, N), in_dtype),
+            C: T.Tensor((M, N), out_dtype),
     ):
         with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=threads) as (bx, by):
             start_x = bx * block_N
@@ -105,7 +105,7 @@ When compiling the example below, let's set `N` to 2047:
 def elementwise_add(N, num_per_thread=8, threads=256, dtype="bfloat16"):
 
     @T.prim_func
-    def main(A: T.Buffer((N), dtype), B: T.Buffer((N), dtype), C: T.Buffer((N), dtype)):
+    def main(A: T.Tensor((N), dtype), B: T.Tensor((N), dtype), C: T.Tensor((N), dtype)):
         with T.Kernel(T.ceildiv(N, threads * num_per_thread), threads=threads) as (b_x):
             # vector add.
             for i, j in T.Parallel(threads, num_per_thread):
@@ -179,7 +179,7 @@ In such scenarios, explicitly specifying the number of elements computed per thr
 def elementwise_add(N, num_per_thread=8, threads=256, dtype="bfloat16"):
 
     @T.prim_func
-    def main(A: T.Buffer((N), dtype), B: T.Buffer((N), dtype), C: T.Buffer((N), dtype)):
+    def main(A: T.Tensor((N), dtype), B: T.Tensor((N), dtype), C: T.Tensor((N), dtype)):
         with T.Kernel(T.ceildiv(N, threads * num_per_thread), threads=threads) as (b_x):
             # vector add.
             for i, j in T.Parallel(threads, num_per_thread):
@@ -215,7 +215,7 @@ But what happens if we provide additional hints to TileLang? For instance, by ex
 def elementwise_add(N, NUM_ELE_PER_THREAD=8, threads=256, dtype="bfloat16"):
 
     @T.prim_func
-    def main(A: T.Buffer((N), dtype), B: T.Buffer((N), dtype), C: T.Buffer((N), dtype)):
+    def main(A: T.Tensor((N), dtype), B: T.Tensor((N), dtype), C: T.Tensor((N), dtype)):
         with T.Kernel(T.ceildiv(N, threads * NUM_ELE_PER_THREAD), threads=threads) as (b_x):
             A_register = T.alloc_fragment((threads * NUM_ELE_PER_THREAD), dtype)
             B_register = T.alloc_fragment((threads * NUM_ELE_PER_THREAD), dtype)

_sources/deeplearning_operators/matmul.md.txt

@@ -67,9 +67,9 @@ from tilelang.intrinsics import make_mma_swizzle_layout
 def matmul(M, N, K, block_M, block_N, block_K, dtype="float16", accum_dtype="float"):
     @T.prim_func
     def main(
-        A: T.Buffer((M, K), dtype),
-        B: T.Buffer((K, N), dtype),
-        C: T.Buffer((M, N), dtype),
+        A: T.Tensor((M, K), dtype),
+        B: T.Tensor((K, N), dtype),
+        C: T.Tensor((M, N), dtype),
     ):
         # Initialize Kernel Context
         with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):

_sources/tutorials/debug_tools_for_tilelang.md.txt

@@ -66,7 +66,7 @@ For example, consider a case where a simple `T.copy` in 1D causes the lowering p
 
 ```python
 @T.prim_func
-def main(Q: T.Buffer(shape_q, dtype)):
+def main(Q: T.Tensor(shape_q, dtype)):
     # ...existing code...
 ```
 

deeplearning_operators/elementwise.html

@@ -284,9 +284,9 @@ <h3>Basic Structure<a class="headerlink" href="#basic-structure" title="Permalin
 <span class="k">def</span> <span class="nf">matmul</span><span class="p">(</span><span class="n">M</span><span class="p">,</span> <span class="n">N</span><span class="p">,</span> <span class="n">K</span><span class="p">,</span> <span class="n">block_M</span><span class="p">,</span> <span class="n">block_N</span><span class="p">,</span> <span class="n">block_K</span><span class="p">,</span> <span class="n">dtype</span><span class="o">=</span><span class="s2">&quot;float16&quot;</span><span class="p">,</span> <span class="n">accum_dtype</span><span class="o">=</span><span class="s2">&quot;float&quot;</span><span class="p">):</span>
     <span class="nd">@T</span><span class="o">.</span><span class="n">prim_func</span>
     <span class="k">def</span> <span class="nf">main</span><span class="p">(</span>
-        <span class="n">A</span><span class="p">:</span> <span class="n">T</span><span class="o">.</span><span class="n">Buffer</span><span class="p">((</span><span class="n">M</span><span class="p">,</span> <span class="n">K</span><span class="p">),</span> <span class="n">dtype</span><span class="p">),</span>
-        <span class="n">B</span><span class="p">:</span> <span class="n">T</span><span class="o">.</span><span class="n">Buffer</span><span class="p">((</span><span class="n">K</span><span class="p">,</span> <span class="n">N</span><span class="p">),</span> <span class="n">dtype</span><span class="p">),</span>
-        <span class="n">C</span><span class="p">:</span> <span class="n">T</span><span class="o">.</span><span class="n">Buffer</span><span class="p">((</span><span class="n">M</span><span class="p">,</span> <span class="n">N</span><span class="p">),</span> <span class="n">dtype</span><span class="p">),</span>
+        <span class="n">A</span><span class="p">:</span> <span class="n">T</span><span class="o">.</span><span class="n">Tensor</span><span class="p">((</span><span class="n">M</span><span class="p">,</span> <span class="n">K</span><span class="p">),</span> <span class="n">dtype</span><span class="p">),</span>
+        <span class="n">B</span><span class="p">:</span> <span class="n">T</span><span class="o">.</span><span class="n">Tensor</span><span class="p">((</span><span class="n">K</span><span class="p">,</span> <span class="n">N</span><span class="p">),</span> <span class="n">dtype</span><span class="p">),</span>
+        <span class="n">C</span><span class="p">:</span> <span class="n">T</span><span class="o">.</span><span class="n">Tensor</span><span class="p">((</span><span class="n">M</span><span class="p">,</span> <span class="n">N</span><span class="p">),</span> <span class="n">dtype</span><span class="p">),</span>
     <span class="p">):</span>
         <span class="c1"># Initialize Kernel Context</span>
         <span class="k">with</span> <span class="n">T</span><span class="o">.</span><span class="n">Kernel</span><span class="p">(</span><span class="n">T</span><span class="o">.</span><span class="n">ceildiv</span><span class="p">(</span><span class="n">N</span><span class="p">,</span> <span class="n">block_N</span><span class="p">),</span> <span class="n">T</span><span class="o">.</span><span class="n">ceildiv</span><span class="p">(</span><span class="n">M</span><span class="p">,</span> <span class="n">block_M</span><span class="p">),</span> <span class="n">threads</span><span class="o">=</span><span class="mi">128</span><span class="p">)</span> <span class="k">as</span> <span class="p">(</span><span class="n">bx</span><span class="p">,</span> <span class="n">by</span><span class="p">):</span>

deeplearning_operators/matmul.html

@@ -279,7 +279,7 @@ <h2>Debugging Generation Issues<a class="headerlink" href="#debugging-generation
 <p>When the code fails to generate (for instance, a compilation error occurs), you do <strong>not</strong> necessarily need to jump directly into C++ passes to debug. Instead, you can first inspect the intermediate representations (IR) in Python by printing them.</p>
 <p>For example, consider a case where a simple <code class="docutils literal notranslate"><span class="pre">T.copy</span></code> in 1D causes the lowering process to fail. The snippet below illustrates a simplified version of the problem (based on community Issue #35):</p>
 <div class="highlight-python notranslate"><div class="highlight"><pre><span></span><span class="nd">@T</span><span class="o">.</span><span class="n">prim_func</span>
-<span class="k">def</span> <span class="nf">main</span><span class="p">(</span><span class="n">Q</span><span class="p">:</span> <span class="n">T</span><span class="o">.</span><span class="n">Buffer</span><span class="p">(</span><span class="n">shape_q</span><span class="p">,</span> <span class="n">dtype</span><span class="p">)):</span>
+<span class="k">def</span> <span class="nf">main</span><span class="p">(</span><span class="n">Q</span><span class="p">:</span> <span class="n">T</span><span class="o">.</span><span class="n">Tensor</span><span class="p">(</span><span class="n">shape_q</span><span class="p">,</span> <span class="n">dtype</span><span class="p">)):</span>
     <span class="c1"># ...existing code...</span>
 </pre></div>
 </div>