Update docs

Author: github-actions[bot]

_sources/autoapi/tilelang/contrib/nvcc/index.rst.txt

@@ -30,6 +30,7 @@ Functions
    tilelang.contrib.nvcc.have_bf16
    tilelang.contrib.nvcc.have_fp8
    tilelang.contrib.nvcc.have_tma
+   tilelang.contrib.nvcc.is_hopper
    tilelang.contrib.nvcc.get_nvcc_compiler
 
 
@@ -176,6 +177,8 @@ Module Contents
    :type target: tvm.target.Target
 
 
+.. py:function:: is_hopper(target)
+
 .. py:function:: get_nvcc_compiler()
 
    Get the path to the nvcc compiler

_sources/autoapi/tilelang/language/allocate/index.rst.txt

@@ -32,6 +32,7 @@ Functions
    tilelang.language.allocate.alloc_fragment
    tilelang.language.allocate.alloc_var
    tilelang.language.allocate.alloc_barrier
+   tilelang.language.allocate.alloc_tmem
    tilelang.language.allocate.alloc_reducer
 
 
@@ -107,6 +108,32 @@ Module Contents
    :rtype: T.Buffer
 
 
+.. py:function:: alloc_tmem(shape, dtype)
+
+   Allocate a Tensor Memory (TMEM) buffer for use with 5th generation Tensor Core operations (e.g., UMMA).
+
+   TMEM is a dedicated on-chip memory introduced in Hopper GPUs, designed to reduce register pressure and enable asynchronous, single-threaded MMA operations. It is organized as a 2D array of 512 columns by 128 rows (lanes), with each cell being 32 bits. Allocation is performed in units of columns, and every lane of a column is allocated together.
+
+   Key properties and requirements:
+       - The number of columns allocated must be a power of 2 and at least 32.
+       - TMEM allocations are dynamic and must be explicitly deallocated.
+       - Both allocation and deallocation must be performed by the same warp.
+       - The base address of the TMEM allocation is stored in shared memory and used as the offset for UMMA accumulator tensors.
+       - Only UMMA and specific TMEM load/store instructions can access TMEM; all pre-processing must occur before data is loaded into TMEM, and all post-processing after data is retrieved.
+       - The number of columns allocated should not increase between any two allocations in the execution order within the CTA.
+
+   :param num_cols: Number of columns to allocate in TMEM. Must be a power of 2 and >= 32 but less than or equal to 512.
+   :type num_cols: int
+
+   :returns: A TVM buffer object allocated in TMEM scope, suitable for use as an accumulator or operand in UMMA operations.
+   :rtype: T.Buffer
+
+   .. note::
+
+      - TMEM is only available on supported architectures (e.g., Hopper and later).
+      - The buffer returned should be used according to TMEM access restrictions and deallocated appropriately.
+
+
 .. py:function:: alloc_reducer(shape, dtype, op='sum', replication=None)
 
    Allocate a reducer buffer.

_sources/autoapi/tilelang/language/fastmath/index.rst.txt

@@ -0,0 +1,6 @@
+tilelang.language.fastmath
+==========================
+
+.. py:module:: tilelang.language.fastmath
+
+

_sources/autoapi/tilelang/language/gemm/index.rst.txt

@@ -21,7 +21,7 @@ Functions
 Module Contents
 ---------------
 
-.. py:function:: gemm(A, B, C, transpose_A = False, transpose_B = False, policy = GemmWarpPolicy.Square, clear_accum = False, k_pack = 1, wg_wait = 0)
+.. py:function:: gemm(A, B, C, transpose_A = False, transpose_B = False, policy = GemmWarpPolicy.Square, clear_accum = False, k_pack = 1, wg_wait = 0, mbar = None)
 
    Perform a General Matrix Multiplication (GEMM) operation.
 
@@ -45,7 +45,11 @@ Module Contents
    :param k_pack: Number of k dimensions packed into a single warp. Defaults to 1.
    :type k_pack: int, optional
    :param wg_wait: Warp group wait count. Defaults to 0.
+                   On hopper it is equivalent to `wgmma.wait_group.sync.aligned <wg_wait>` if wg_wait is not -1
+                   On sm100 (datacenter blackwell), `wg_wait` can only be 0 or -1. `mbarrier_wait(UTCMMA barrier)` will be appended if wg_wait is 0.
    :type wg_wait: int, optional
+   :param mbar: mbarrier for UTCMMA synchronization
+   :type mbar: tir.Buffer, optional
 
    :returns: A handle to the GEMM operation
    :rtype: tir.Call

_sources/autoapi/tilelang/language/index.rst.txt

@@ -21,11 +21,13 @@ Submodules
    /autoapi/tilelang/language/copy/index
    /autoapi/tilelang/language/customize/index
    /autoapi/tilelang/language/experimental/index
+   /autoapi/tilelang/language/fastmath/index
    /autoapi/tilelang/language/fill/index
    /autoapi/tilelang/language/frame/index
    /autoapi/tilelang/language/gemm/index
    /autoapi/tilelang/language/kernel/index
    /autoapi/tilelang/language/logical/index
+   /autoapi/tilelang/language/math_intrinsics/index
    /autoapi/tilelang/language/parallel/index
    /autoapi/tilelang/language/persistent/index
    /autoapi/tilelang/language/pipeline/index

_sources/autoapi/tilelang/language/math_intrinsics/index.rst.txt

@@ -0,0 +1,140 @@
+tilelang.language.math_intrinsics
+=================================
+
+.. py:module:: tilelang.language.math_intrinsics
+
+
+Functions
+---------
+
+.. autoapisummary::
+
+   tilelang.language.math_intrinsics.ieee_add
+   tilelang.language.math_intrinsics.ieee_sub
+   tilelang.language.math_intrinsics.ieee_mul
+   tilelang.language.math_intrinsics.ieee_fmaf
+   tilelang.language.math_intrinsics.ieee_frcp
+   tilelang.language.math_intrinsics.ieee_fsqrt
+   tilelang.language.math_intrinsics.ieee_frsqrt
+   tilelang.language.math_intrinsics.ieee_fdiv
+
+
+Module Contents
+---------------
+
+.. py:function:: ieee_add(x, y, rounding_mode='rn')
+
+   IEEE-compliant addition with specified rounding mode
+
+   :param x: First operand.
+   :type x: PrimExpr
+   :param y: Second operand.
+   :type y: PrimExpr
+   :param rounding_mode: Rounding mode: 'rn' (round to nearest), 'rz' (round toward zero),
+                         'ru' (round toward positive infinity), 'rd' (round toward negative infinity).
+                         Default is 'rn'.
+   :type rounding_mode: str, optional
+
+   :returns: **result** -- The result.
+   :rtype: PrimExpr
+
+
+.. py:function:: ieee_sub(x, y, rounding_mode='rn')
+
+   IEEE-compliant subtraction with specified rounding mode
+
+   :param x: First operand.
+   :type x: PrimExpr
+   :param y: Second operand.
+   :type y: PrimExpr
+   :param rounding_mode: Rounding mode: 'rn', 'rz', 'ru', 'rd'. Default is 'rn'.
+   :type rounding_mode: str, optional
+
+   :returns: **result** -- The result.
+   :rtype: PrimExpr
+
+
+.. py:function:: ieee_mul(x, y, rounding_mode='rn')
+
+   IEEE-compliant multiplication with specified rounding mode
+
+   :param x: First operand.
+   :type x: PrimExpr
+   :param y: Second operand.
+   :type y: PrimExpr
+   :param rounding_mode: Rounding mode: 'rn', 'rz', 'ru', 'rd'. Default is 'rn'.
+   :type rounding_mode: str, optional
+
+   :returns: **result** -- The result.
+   :rtype: PrimExpr
+
+
+.. py:function:: ieee_fmaf(x, y, z, rounding_mode='rn')
+
+   IEEE-compliant fused multiply-add with specified rounding mode
+
+   :param x: First operand.
+   :type x: PrimExpr
+   :param y: Second operand.
+   :type y: PrimExpr
+   :param z: Third operand (addend).
+   :type z: PrimExpr
+   :param rounding_mode: Rounding mode: 'rn', 'rz', 'ru', 'rd'. Default is 'rn'.
+   :type rounding_mode: str, optional
+
+   :returns: **result** -- The result of x * y + z.
+   :rtype: PrimExpr
+
+
+.. py:function:: ieee_frcp(x, rounding_mode='rn')
+
+   IEEE-compliant reciprocal with specified rounding mode
+
+   :param x: Input operand.
+   :type x: PrimExpr
+   :param rounding_mode: Rounding mode: 'rn', 'rz', 'ru', 'rd'. Default is 'rn'.
+   :type rounding_mode: str, optional
+
+   :returns: **result** -- The result of 1/x.
+   :rtype: PrimExpr
+
+
+.. py:function:: ieee_fsqrt(x, rounding_mode='rn')
+
+   IEEE-compliant square root with specified rounding mode
+
+   :param x: Input operand.
+   :type x: PrimExpr
+   :param rounding_mode: Rounding mode: 'rn', 'rz', 'ru', 'rd'. Default is 'rn'.
+   :type rounding_mode: str, optional
+
+   :returns: **result** -- The result of sqrt(x).
+   :rtype: PrimExpr
+
+
+.. py:function:: ieee_frsqrt(x)
+
+   IEEE-compliant reciprocal square root (round to nearest only)
+
+   :param x: Input operand.
+   :type x: PrimExpr
+
+   :returns: **result** -- The result of 1/sqrt(x).
+   :rtype: PrimExpr
+
+
+.. py:function:: ieee_fdiv(x, y, rounding_mode='rn')
+
+   IEEE-compliant division with specified rounding mode
+
+   :param x: Dividend.
+   :type x: PrimExpr
+   :param y: Divisor.
+   :type y: PrimExpr
+   :param rounding_mode: Rounding mode: 'rn', 'rz', 'ru', 'rd'. Default is 'rn'.
+   :type rounding_mode: str, optional
+
+   :returns: **result** -- The result of x/y.
+   :rtype: PrimExpr
+
+

_sources/autoapi/tilelang/transform/index.rst.txt

@@ -31,6 +31,7 @@ Functions
    tilelang.transform.LayoutInference
    tilelang.transform.LowerTileOp
    tilelang.transform.InjectSoftwarePipeline
+   tilelang.transform.FrontendLegalize
    tilelang.transform.InjectAssumes
    tilelang.transform.LowerHopperIntrin
    tilelang.transform.WarpSpecializedPipeline
@@ -64,6 +65,7 @@ Functions
    tilelang.transform.LowerOpaqueBlock
    tilelang.transform.LowerThreadAllreduce
    tilelang.transform.LowerDeviceKernelLaunch
+   tilelang.transform.LowerSharedTmem
    tilelang.transform.LayoutReducer
 
 
@@ -115,6 +117,14 @@ Package Contents
    :rtype: tvm.transform.Pass
 
 
+.. py:function:: FrontendLegalize()
+
+   FrontendLegalize
+
+   :returns: **fpass** -- The result pass
+   :rtype: tvm.transform.Pass
+
+
 .. py:function:: InjectAssumes()
 
    Inject Assumes
@@ -384,6 +394,12 @@ Package Contents
    :rtype: tvm.transform.Pass
 
 
+.. py:function:: LowerSharedTmem()
+
+   LowerSharedTmem
+
+
+
 .. py:function:: LayoutReducer()
 
    Return a TVM transform pass that performs layout reduction/normalization.

_sources/autoapi/tilelang/transform/pass_config/index.rst.txt

@@ -123,6 +123,24 @@ Module Contents
       :type: Disable safe memory access optimization. Default
 
 
+   .. py:attribute:: TL_DISABLE_VECTORIZE_256
+      :value: 'tl.disable_vectorize_256'
+
+
+      False
+
+      :type: Disable usage of LDG/STG 256. Default
+
+
+   .. py:attribute:: TL_DISABLE_WGMMA
+      :value: 'tl.disable_wgmma'
+
+
+      False
+
+      :type: Disable usage of Hopper WGMMA. Default
+
+
    .. py:attribute:: TL_DEBUG_MERGE_SHARED_MEMORY_ALLOCATIONS
       :value: 'tl.debug_merge_shared_memory_allocations'
 

autoapi/index.html

@@ -340,11 +340,13 @@
 <li class="toctree-l4"><a class="reference internal" href="tilelang/language/experimental/gemm_sp/index.html">tilelang.language.experimental.gemm_sp</a></li>
 </ul>
 </li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang/language/fastmath/index.html">tilelang.language.fastmath</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/fill/index.html">tilelang.language.fill</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/frame/index.html">tilelang.language.frame</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/gemm/index.html">tilelang.language.gemm</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/kernel/index.html">tilelang.language.kernel</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/logical/index.html">tilelang.language.logical</a></li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang/language/math_intrinsics/index.html">tilelang.language.math_intrinsics</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/parallel/index.html">tilelang.language.parallel</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/persistent/index.html">tilelang.language.persistent</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/pipeline/index.html">tilelang.language.pipeline</a></li>
@@ -588,11 +590,13 @@ <h1>API Reference<a class="headerlink" href="#api-reference" title="Link to this
 <li class="toctree-l4"><a class="reference internal" href="tilelang/language/experimental/gemm_sp/index.html">tilelang.language.experimental.gemm_sp</a></li>
 </ul>
 </li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang/language/fastmath/index.html">tilelang.language.fastmath</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/fill/index.html">tilelang.language.fill</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/frame/index.html">tilelang.language.frame</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/gemm/index.html">tilelang.language.gemm</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/kernel/index.html">tilelang.language.kernel</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/logical/index.html">tilelang.language.logical</a></li>
+<li class="toctree-l3"><a class="reference internal" href="tilelang/language/math_intrinsics/index.html">tilelang.language.math_intrinsics</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/parallel/index.html">tilelang.language.parallel</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/persistent/index.html">tilelang.language.persistent</a></li>
 <li class="toctree-l3"><a class="reference internal" href="tilelang/language/pipeline/index.html">tilelang.language.pipeline</a></li>

autoapi/tilelang/autotuner/capture/index.html

@@ -340,11 +340,13 @@
 <li class="toctree-l4"><a class="reference internal" href="../../language/experimental/gemm_sp/index.html">tilelang.language.experimental.gemm_sp</a></li>
 </ul>
 </li>
+<li class="toctree-l3"><a class="reference internal" href="../../language/fastmath/index.html">tilelang.language.fastmath</a></li>
 <li class="toctree-l3"><a class="reference internal" href="../../language/fill/index.html">tilelang.language.fill</a></li>
 <li class="toctree-l3"><a class="reference internal" href="../../language/frame/index.html">tilelang.language.frame</a></li>
 <li class="toctree-l3"><a class="reference internal" href="../../language/gemm/index.html">tilelang.language.gemm</a></li>
 <li class="toctree-l3"><a class="reference internal" href="../../language/kernel/index.html">tilelang.language.kernel</a></li>
 <li class="toctree-l3"><a class="reference internal" href="../../language/logical/index.html">tilelang.language.logical</a></li>
+<li class="toctree-l3"><a class="reference internal" href="../../language/math_intrinsics/index.html">tilelang.language.math_intrinsics</a></li>
 <li class="toctree-l3"><a class="reference internal" href="../../language/parallel/index.html">tilelang.language.parallel</a></li>
 <li class="toctree-l3"><a class="reference internal" href="../../language/persistent/index.html">tilelang.language.persistent</a></li>
 <li class="toctree-l3"><a class="reference internal" href="../../language/pipeline/index.html">tilelang.language.pipeline</a></li>