Update docs

Author: github-actions[bot]

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

@@ -31,7 +31,7 @@ Module Contents
 
    Perform an atomic maximum on the value stored at dst with an optional memory-order.
 
-   Supports scalar/addressed extern atomic max when neither argument exposes extents, or tile-region-based atomic max for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional `memory_order` (one of "relaxed","consume","acquire","release","acq_rel","seq_cst") is used only for the direct extern `AtomicMax` path when no extents are available — otherwise the tile-region path ignores `memory_order`.
+   Supports scalar/addressed extern atomic max when neither argument exposes extents, or tile-region-based atomic max for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional `memory_order` (one of "relaxed","consume","acquire","release","acq_rel","seq_cst") is honored on CUDA targets by both the direct extern `AtomicMax` path and the tile-region path. HIP and CuteDSL codegen currently ignore it.
 
    :param dst: Destination buffer/address to apply the atomic max.
    :type dst: Buffer
@@ -75,7 +75,7 @@ Module Contents
 
    Atomically update the value at dst to the minimum of its current value and value.
 
-   Supports scalar/addressed extern atomic min when neither argument exposes extents, or tile-region-based atomic min for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional `memory_order` (one of "relaxed","consume","acquire","release","acq_rel","seq_cst") is used only for the direct extern `AtomicMin` path when no extents are available — otherwise the tile-region path ignores `memory_order`.
+   Supports scalar/addressed extern atomic min when neither argument exposes extents, or tile-region-based atomic min for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional `memory_order` (one of "relaxed","consume","acquire","release","acq_rel","seq_cst") is honored on CUDA targets by both the direct extern `AtomicMin` path and the tile-region path. HIP and CuteDSL codegen currently ignore it.
 
    :param dst: Destination buffer/address to apply the atomic min.
    :type dst: Buffer
@@ -119,7 +119,7 @@ Module Contents
 
    Atomically add `value` into `dst`, returning a handle to the operation.
 
-   Supports scalar/addressed extern atomic add when neither argument exposes extents, or tile-region-based atomic add for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional `memory_order` (one of "relaxed","consume","acquire","release","acq_rel","seq_cst") is used only for the direct extern `AtomicAdd` path when no extents are available — otherwise the tile-region path ignores `memory_order`.
+   Supports scalar/addressed extern atomic add when neither argument exposes extents, or tile-region-based atomic add for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional `memory_order` (one of "relaxed","consume","acquire","release","acq_rel","seq_cst") is honored on CUDA targets by both paths: the direct extern `AtomicAdd` path when no extents are available, and the tile-region path (including when it auto-vectorizes to `AtomicAddx2`/`AtomicAddx4`). HIP and CuteDSL codegen currently ignore it.
 
    :param dst: Destination buffer/address to apply the atomic add.
    :type dst: Buffer

autoapi/tilelang/language/atomic/index.html

@@ -780,7 +780,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <dt class="sig sig-object py" id="tilelang.language.atomic.atomic_max">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.atomic.</span></span><span class="sig-name descname"><span class="pre">atomic_max</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">dst</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">value</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">memory_order</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">None</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">return_prev</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">annotations</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">None</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.atomic.atomic_max" title="Link to this definition">¶</a></dt>
 <dd><p>Perform an atomic maximum on the value stored at dst with an optional memory-order.</p>
-<p>Supports scalar/addressed extern atomic max when neither argument exposes extents, or tile-region-based atomic max for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional <cite>memory_order</cite> (one of “relaxed”,”consume”,”acquire”,”release”,”acq_rel”,”seq_cst”) is used only for the direct extern <cite>AtomicMax</cite> path when no extents are available — otherwise the tile-region path ignores <cite>memory_order</cite>.</p>
+<p>Supports scalar/addressed extern atomic max when neither argument exposes extents, or tile-region-based atomic max for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional <cite>memory_order</cite> (one of “relaxed”,”consume”,”acquire”,”release”,”acq_rel”,”seq_cst”) is honored on CUDA targets by both the direct extern <cite>AtomicMax</cite> path and the tile-region path. HIP and CuteDSL codegen currently ignore it.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
 <dd class="field-odd"><ul class="simple">
@@ -833,7 +833,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <dt class="sig sig-object py" id="tilelang.language.atomic.atomic_min">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.atomic.</span></span><span class="sig-name descname"><span class="pre">atomic_min</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">dst</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">value</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">memory_order</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">None</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">return_prev</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">annotations</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">None</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.atomic.atomic_min" title="Link to this definition">¶</a></dt>
 <dd><p>Atomically update the value at dst to the minimum of its current value and value.</p>
-<p>Supports scalar/addressed extern atomic min when neither argument exposes extents, or tile-region-based atomic min for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional <cite>memory_order</cite> (one of “relaxed”,”consume”,”acquire”,”release”,”acq_rel”,”seq_cst”) is used only for the direct extern <cite>AtomicMin</cite> path when no extents are available — otherwise the tile-region path ignores <cite>memory_order</cite>.</p>
+<p>Supports scalar/addressed extern atomic min when neither argument exposes extents, or tile-region-based atomic min for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional <cite>memory_order</cite> (one of “relaxed”,”consume”,”acquire”,”release”,”acq_rel”,”seq_cst”) is honored on CUDA targets by both the direct extern <cite>AtomicMin</cite> path and the tile-region path. HIP and CuteDSL codegen currently ignore it.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
 <dd class="field-odd"><ul class="simple">
@@ -886,7 +886,7 @@ <h2>Module Contents<a class="headerlink" href="#module-contents" title="Link to
 <dt class="sig sig-object py" id="tilelang.language.atomic.atomic_add">
 <span class="sig-prename descclassname"><span class="pre">tilelang.language.atomic.</span></span><span class="sig-name descname"><span class="pre">atomic_add</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">dst</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">value</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">memory_order</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">None</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">return_prev</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">use_tma</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">annotations</span></span><span class="o"><span class="pre">=</span></span><span class="default_value"><span class="pre">None</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#tilelang.language.atomic.atomic_add" title="Link to this definition">¶</a></dt>
 <dd><p>Atomically add <cite>value</cite> into <cite>dst</cite>, returning a handle to the operation.</p>
-<p>Supports scalar/addressed extern atomic add when neither argument exposes extents, or tile-region-based atomic add for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional <cite>memory_order</cite> (one of “relaxed”,”consume”,”acquire”,”release”,”acq_rel”,”seq_cst”) is used only for the direct extern <cite>AtomicAdd</cite> path when no extents are available — otherwise the tile-region path ignores <cite>memory_order</cite>.</p>
+<p>Supports scalar/addressed extern atomic add when neither argument exposes extents, or tile-region-based atomic add for Buffer/BufferRegion/BufferLoad inputs. If both arguments are plain Buffers their shapes must be structurally equal. If at least one side exposes extents, extents are aligned (missing dimensions are treated as size 1); an assertion is raised if extents cannot be deduced. The optional <cite>memory_order</cite> (one of “relaxed”,”consume”,”acquire”,”release”,”acq_rel”,”seq_cst”) is honored on CUDA targets by both paths: the direct extern <cite>AtomicAdd</cite> path when no extents are available, and the tile-region path (including when it auto-vectorizes to <cite>AtomicAddx2</cite>/<cite>AtomicAddx4</cite>). HIP and CuteDSL codegen currently ignore it.</p>
 <dl class="field-list simple">
 <dt class="field-odd">Parameters<span class="colon">:</span></dt>
 <dd class="field-odd"><ul class="simple">