rebase

2025-11-04 18:48:15 +08:00 · 2024-10-25 20:23:45 +00:00
parent bd10c96fd8
commit 721548a799
717 changed files with 33254 additions and 34561 deletions
--- a/docs/build/html/scan_8h.html
+++ b/docs/build/html/scan_8h.html
@@ -0,0 +1,478 @@
+<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "https://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
+<html xmlns="http://www.w3.org/1999/xhtml" lang="en-US">
+<head>
+<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
+<meta http-equiv="X-UA-Compatible" content="IE=11"/>
+<meta name="generator" content="Doxygen 1.12.0"/>
+<meta name="viewport" content="width=device-width, initial-scale=1"/>
+<title>MLX: mlx/backend/metal/kernels/scan.h File Reference</title>
+<link href="tabs.css" rel="stylesheet" type="text/css"/>
+<script type="text/javascript" src="jquery.js"></script>
+<script type="text/javascript" src="dynsections.js"></script>
+<script type="text/javascript" src="clipboard.js"></script>
+<link href="navtree.css" rel="stylesheet" type="text/css"/>
+<script type="text/javascript" src="resize.js"></script>
+<script type="text/javascript" src="cookie.js"></script>
+<link href="search/search.css" rel="stylesheet" type="text/css"/>
+<script type="text/javascript" src="search/searchdata.js"></script>
+<script type="text/javascript" src="search/search.js"></script>
+<link href="doxygen.css" rel="stylesheet" type="text/css" />
+</head>
+<body>
+<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
+<div id="titlearea">
+<table cellspacing="0" cellpadding="0">
+ <tbody>
+ <tr id="projectrow">
+  <td id="projectalign">
+   <div id="projectname">MLX
+   </div>
+  </td>
+ </tr>
+ </tbody>
+</table>
+</div>
+<!-- end header part -->
+<!-- Generated by Doxygen 1.12.0 -->
+<script type="text/javascript">
+/* @license magnet:?xt=urn:btih:d3d9a9a6595521f9666a5e94cc830dab83b65699&amp;dn=expat.txt MIT */
+var searchBox = new SearchBox("searchBox", "search/",'.html');
+/* @license-end */
+</script>
+<script type="text/javascript">
+/* @license magnet:?xt=urn:btih:d3d9a9a6595521f9666a5e94cc830dab83b65699&amp;dn=expat.txt MIT */
+$(function() { codefold.init(0); });
+/* @license-end */
+</script>
+<script type="text/javascript" src="menudata.js"></script>
+<script type="text/javascript" src="menu.js"></script>
+<script type="text/javascript">
+/* @license magnet:?xt=urn:btih:d3d9a9a6595521f9666a5e94cc830dab83b65699&amp;dn=expat.txt MIT */
+$(function() {
+  initMenu('',true,false,'search.php','Search',false);
+  $(function() { init_search(); });
+});
+/* @license-end */
+</script>
+<div id="main-nav"></div>
+<script type="text/javascript">
+/* @license magnet:?xt=urn:btih:d3d9a9a6595521f9666a5e94cc830dab83b65699&amp;dn=expat.txt MIT */
+$(function(){ initResizable(false); });
+/* @license-end */
+</script>
+<!-- window showing the filter options -->
+<div id="MSearchSelectWindow"
+     onmouseover="return searchBox.OnSearchSelectShow()"
+     onmouseout="return searchBox.OnSearchSelectHide()"
+     onkeydown="return searchBox.OnSearchSelectKey(event)">
+</div>
+
+<!-- iframe showing the search results (closed by default) -->
+<div id="MSearchResultsWindow">
+<div id="MSearchResults">
+<div class="SRPage">
+<div id="SRIndex">
+<div id="SRResults"></div>
+<div class="SRStatus" id="Loading">Loading...</div>
+<div class="SRStatus" id="Searching">Searching...</div>
+<div class="SRStatus" id="NoMatches">No Matches</div>
+</div>
+</div>
+</div>
+</div>
+
+<div id="nav-path" class="navpath">
+  <ul>
+<li class="navelem"><a class="el" href="dir_938ab0ecf10b8b860ff766c820f665fd.html">mlx</a></li><li class="navelem"><a class="el" href="dir_1d446c9bd3c99228254c9484e0bc5c06.html">backend</a></li><li class="navelem"><a class="el" href="dir_d0c977ea65824390717cdb7efc36c157.html">metal</a></li><li class="navelem"><a class="el" href="dir_70a37effa88bcbd6b791977fa1e64356.html">kernels</a></li>  </ul>
+</div>
+</div><!-- top -->
+<div id="doc-content">
+<div class="header">
+  <div class="summary">
+<a href="#nested-classes">Classes</a> &#124;
+<a href="#define-members">Macros</a> &#124;
+<a href="#func-members">Functions</a>  </div>
+  <div class="headertitle"><div class="title">scan.h File Reference</div></div>
+</div><!--header-->
+<div class="contents">
+
+<p><a href="scan_8h_source.html">Go to the source code of this file.</a></p>
+<table class="memberdecls">
+<tr class="heading"><td colspan="2"><h2 class="groupheader"><a id="nested-classes" name="nested-classes"></a>
+Classes</h2></td></tr>
+<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="struct_cum_sum.html">CumSum&lt; U &gt;</a></td></tr>
+<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
+<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="struct_cum_prod.html">CumProd&lt; U &gt;</a></td></tr>
+<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
+<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="struct_cum_prod_3_01bool_01_4.html">CumProd&lt; bool &gt;</a></td></tr>
+<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
+<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="struct_cum_max.html">CumMax&lt; U &gt;</a></td></tr>
+<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
+<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="struct_cum_min.html">CumMin&lt; U &gt;</a></td></tr>
+<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
+</table><table class="memberdecls">
+<tr class="heading"><td colspan="2"><h2 class="groupheader"><a id="define-members" name="define-members"></a>
+Macros</h2></td></tr>
+<tr class="memitem:a0d8d6a9b0f3a1263629380bda8eca7bc" id="r_a0d8d6a9b0f3a1263629380bda8eca7bc"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="#a0d8d6a9b0f3a1263629380bda8eca7bc">DEFINE_SIMD_SCAN</a>()</td></tr>
+<tr class="separator:a0d8d6a9b0f3a1263629380bda8eca7bc"><td class="memSeparator" colspan="2">&#160;</td></tr>
+<tr class="memitem:a185f66aac8c5317587e6abd43f3013fc" id="r_a185f66aac8c5317587e6abd43f3013fc"><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="#a185f66aac8c5317587e6abd43f3013fc">DEFINE_SIMD_EXCLUSIVE_SCAN</a>()</td></tr>
+<tr class="separator:a185f66aac8c5317587e6abd43f3013fc"><td class="memSeparator" colspan="2">&#160;</td></tr>
+</table><table class="memberdecls">
+<tr class="heading"><td colspan="2"><h2 class="groupheader"><a id="func-members" name="func-members"></a>
+Functions</h2></td></tr>
+<tr class="memitem:a9c415d07921f3961bad0a00a34f4a9a3" id="r_a9c415d07921f3961bad0a00a34f4a9a3"><td class="memTemplParams" colspan="2">template&lt;typename T , typename U , int N_READS, bool reverse&gt; </td></tr>
+<tr class="memitem:a9c415d07921f3961bad0a00a34f4a9a3"><td class="memTemplItemLeft" align="right" valign="top">void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="#a9c415d07921f3961bad0a00a34f4a9a3">load_unsafe</a> (U values[N_READS], const device T *input)</td></tr>
+<tr class="separator:a9c415d07921f3961bad0a00a34f4a9a3"><td class="memSeparator" colspan="2">&#160;</td></tr>
+<tr class="memitem:ae8eb101e538b85f8a4bcf451489ae0ac" id="r_ae8eb101e538b85f8a4bcf451489ae0ac"><td class="memTemplParams" colspan="2">template&lt;typename T , typename U , int N_READS, bool reverse&gt; </td></tr>
+<tr class="memitem:ae8eb101e538b85f8a4bcf451489ae0ac"><td class="memTemplItemLeft" align="right" valign="top">void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="#ae8eb101e538b85f8a4bcf451489ae0ac">load_safe</a> (U values[N_READS], const device T *input, int start, int total, U init)</td></tr>
+<tr class="separator:ae8eb101e538b85f8a4bcf451489ae0ac"><td class="memSeparator" colspan="2">&#160;</td></tr>
+<tr class="memitem:a8010e7bdf7a72cbd35ce7cd7ecb08e32" id="r_a8010e7bdf7a72cbd35ce7cd7ecb08e32"><td class="memTemplParams" colspan="2">template&lt;typename U , int N_READS, bool reverse&gt; </td></tr>
+<tr class="memitem:a8010e7bdf7a72cbd35ce7cd7ecb08e32"><td class="memTemplItemLeft" align="right" valign="top">void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="#a8010e7bdf7a72cbd35ce7cd7ecb08e32">write_unsafe</a> (U values[N_READS], device U *out)</td></tr>
+<tr class="separator:a8010e7bdf7a72cbd35ce7cd7ecb08e32"><td class="memSeparator" colspan="2">&#160;</td></tr>
+<tr class="memitem:ae86aef08e5ebc8790031eb51eefa754c" id="r_ae86aef08e5ebc8790031eb51eefa754c"><td class="memTemplParams" colspan="2">template&lt;typename U , int N_READS, bool reverse&gt; </td></tr>
+<tr class="memitem:ae86aef08e5ebc8790031eb51eefa754c"><td class="memTemplItemLeft" align="right" valign="top">void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="#ae86aef08e5ebc8790031eb51eefa754c">write_safe</a> (U values[N_READS], device U *out, int start, int total)</td></tr>
+<tr class="separator:ae86aef08e5ebc8790031eb51eefa754c"><td class="memSeparator" colspan="2">&#160;</td></tr>
+<tr class="memitem:a60d279b9add7d56639bb209408f09d79" id="r_a60d279b9add7d56639bb209408f09d79"><td class="memTemplParams" colspan="2">template&lt;typename T , typename U , typename Op , int N_READS, bool inclusive, bool reverse&gt; </td></tr>
+<tr class="memitem:a60d279b9add7d56639bb209408f09d79"><td class="memTemplItemLeft" align="right" valign="top">void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="#a60d279b9add7d56639bb209408f09d79">contiguous_scan</a> (const device T *in, device U *out, const constant size_t &amp;axis_size, uint3 gid, uint3 gsize, uint3 lid, uint3 lsize, uint simd_lane_id, uint simd_group_id)</td></tr>
+<tr class="separator:a60d279b9add7d56639bb209408f09d79"><td class="memSeparator" colspan="2">&#160;</td></tr>
+<tr class="memitem:a7abb6ffb6c3b96b88c2a63cd4cc2f7ae" id="r_a7abb6ffb6c3b96b88c2a63cd4cc2f7ae"><td class="memTemplParams" colspan="2">template&lt;typename T , typename U , typename Op , int N_READS, bool inclusive, bool reverse&gt; </td></tr>
+<tr class="memitem:a7abb6ffb6c3b96b88c2a63cd4cc2f7ae"><td class="memTemplItemLeft" align="right" valign="top">void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="#a7abb6ffb6c3b96b88c2a63cd4cc2f7ae">strided_scan</a> (const device T *in, device U *out, const constant size_t &amp;axis_size, const constant size_t &amp;stride, const constant size_t &amp;stride_blocks, uint3 gid, uint3 gsize, uint3 lid, uint simd_lane_id, uint simd_group_id)</td></tr>
+<tr class="separator:a7abb6ffb6c3b96b88c2a63cd4cc2f7ae"><td class="memSeparator" colspan="2">&#160;</td></tr>
+</table>
+<h2 class="groupheader">Macro Definition Documentation</h2>
+<a id="a185f66aac8c5317587e6abd43f3013fc" name="a185f66aac8c5317587e6abd43f3013fc"></a>
+<h2 class="memtitle"><span class="permalink"><a href="#a185f66aac8c5317587e6abd43f3013fc">&#9670;&#160;</a></span>DEFINE_SIMD_EXCLUSIVE_SCAN</h2>
+
+<div class="memitem">
+<div class="memproto">
+      <table class="memname">
+        <tr>
+          <td class="memname">#define DEFINE_SIMD_EXCLUSIVE_SCAN</td>
+          <td>(</td>
+          <td class="paramname"><span class="paramname"><em></em></span></td><td>)</td>
+          <td></td>
+        </tr>
+      </table>
+</div><div class="memdoc">
+<b>Value:</b><div class="fragment"><div class="line">  <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, metal::enable_if_t&lt;sizeof(T) &lt; 8, <span class="keywordtype">bool</span>&gt; = true&gt;  \</div>
+<div class="line">  T simd_exclusive_scan(T val) {                                         \</div>
+<div class="line">    return simd_exclusive_scan_impl(val);                                \</div>
+<div class="line">  }                                                                      \</div>
+<div class="line">                                                                         \</div>
+<div class="line">  <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, metal::enable_if_t&lt;sizeof(T) == 8, <span class="keywordtype">bool</span>&gt; = true&gt; \</div>
+<div class="line">  T simd_exclusive_scan(T val) {                                         \</div>
+<div class="line">    val = simd_scan(val);                                                \</div>
+<div class="line">    <span class="keywordflow">return</span> <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a5862d5ea154c9b76cf56a630cf6385b4">simd_shuffle_and_fill_up</a>(val, init, 1);                       \</div>
+<div class="line">  }</div>
+<div class="ttc" id="abackend_2metal_2kernels_2utils_8h_html_a5862d5ea154c9b76cf56a630cf6385b4"><div class="ttname"><a href="backend_2metal_2kernels_2utils_8h.html#a5862d5ea154c9b76cf56a630cf6385b4">simd_shuffle_and_fill_up</a></div><div class="ttdeci">uint64_t simd_shuffle_and_fill_up(uint64_t data, uint64_t filling, uint16_t delta)</div><div class="ttdef"><b>Definition</b> utils.h:342</div></div>
+</div><!-- fragment -->
+</div>
+</div>
+<a id="a0d8d6a9b0f3a1263629380bda8eca7bc" name="a0d8d6a9b0f3a1263629380bda8eca7bc"></a>
+<h2 class="memtitle"><span class="permalink"><a href="#a0d8d6a9b0f3a1263629380bda8eca7bc">&#9670;&#160;</a></span>DEFINE_SIMD_SCAN</h2>
+
+<div class="memitem">
+<div class="memproto">
+      <table class="memname">
+        <tr>
+          <td class="memname">#define DEFINE_SIMD_SCAN</td>
+          <td>(</td>
+          <td class="paramname"><span class="paramname"><em></em></span></td><td>)</td>
+          <td></td>
+        </tr>
+      </table>
+</div><div class="memdoc">
+<b>Value:</b><div class="fragment"><div class="line">  <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, metal::enable_if_t&lt;sizeof(T) &lt; 8, <span class="keywordtype">bool</span>&gt; = true&gt;  \</div>
+<div class="line">  T simd_scan(T val) {                                                   \</div>
+<div class="line">    return simd_scan_impl(val);                                          \</div>
+<div class="line">  }                                                                      \</div>
+<div class="line">                                                                         \</div>
+<div class="line">  <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, metal::enable_if_t&lt;sizeof(T) == 8, <span class="keywordtype">bool</span>&gt; = true&gt; \</div>
+<div class="line">  T simd_scan(T val) {                                                   \</div>
+<div class="line">    for (<span class="keywordtype">int</span> i = 1; i &lt;= 16; i *= 2) {                                   \</div>
+<div class="line">      val = operator()(val, <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a5862d5ea154c9b76cf56a630cf6385b4">simd_shuffle_and_fill_up</a>(val, init, i));     \</div>
+<div class="line">    }                                                                    \</div>
+<div class="line">    <span class="keywordflow">return</span> val;                                                          \</div>
+<div class="line">  }</div>
+</div><!-- fragment -->
+</div>
+</div>
+<h2 class="groupheader">Function Documentation</h2>
+<a id="a60d279b9add7d56639bb209408f09d79" name="a60d279b9add7d56639bb209408f09d79"></a>
+<h2 class="memtitle"><span class="permalink"><a href="#a60d279b9add7d56639bb209408f09d79">&#9670;&#160;</a></span>contiguous_scan()</h2>
+
+<div class="memitem">
+<div class="memproto">
+<div class="memtemplate">
+template&lt;typename T , typename U , typename Op , int N_READS, bool inclusive, bool reverse&gt; </div>
+      <table class="memname">
+        <tr>
+          <td class="memname">void contiguous_scan </td>
+          <td>(</td>
+          <td class="paramtype">const device T *</td>          <td class="paramname"><span class="paramname"><em>in</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">device U *</td>          <td class="paramname"><span class="paramname"><em>out</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">const constant size_t &amp;</td>          <td class="paramname"><span class="paramname"><em>axis_size</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">uint3</td>          <td class="paramname"><span class="paramname"><em>gid</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">uint3</td>          <td class="paramname"><span class="paramname"><em>gsize</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">uint3</td>          <td class="paramname"><span class="paramname"><em>lid</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">uint3</td>          <td class="paramname"><span class="paramname"><em>lsize</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">uint</td>          <td class="paramname"><span class="paramname"><em>simd_lane_id</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">uint</td>          <td class="paramname"><span class="paramname"><em>simd_group_id</em></span>&#160;)</td>
+        </tr>
+      </table>
+</div><div class="memdoc">
+
+</div>
+</div>
+<a id="ae8eb101e538b85f8a4bcf451489ae0ac" name="ae8eb101e538b85f8a4bcf451489ae0ac"></a>
+<h2 class="memtitle"><span class="permalink"><a href="#ae8eb101e538b85f8a4bcf451489ae0ac">&#9670;&#160;</a></span>load_safe()</h2>
+
+<div class="memitem">
+<div class="memproto">
+<div class="memtemplate">
+template&lt;typename T , typename U , int N_READS, bool reverse&gt; </div>
+<table class="mlabels">
+  <tr>
+  <td class="mlabels-left">
+      <table class="memname">
+        <tr>
+          <td class="memname">void load_safe </td>
+          <td>(</td>
+          <td class="paramtype">U</td>          <td class="paramname"><span class="paramname"><em>values</em></span>[N_READS], </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">const device T *</td>          <td class="paramname"><span class="paramname"><em>input</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">int</td>          <td class="paramname"><span class="paramname"><em>start</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">int</td>          <td class="paramname"><span class="paramname"><em>total</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">U</td>          <td class="paramname"><span class="paramname"><em>init</em></span>&#160;)</td>
+        </tr>
+      </table>
+  </td>
+  <td class="mlabels-right">
+<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
+  </tr>
+</table>
+</div><div class="memdoc">
+
+</div>
+</div>
+<a id="a9c415d07921f3961bad0a00a34f4a9a3" name="a9c415d07921f3961bad0a00a34f4a9a3"></a>
+<h2 class="memtitle"><span class="permalink"><a href="#a9c415d07921f3961bad0a00a34f4a9a3">&#9670;&#160;</a></span>load_unsafe()</h2>
+
+<div class="memitem">
+<div class="memproto">
+<div class="memtemplate">
+template&lt;typename T , typename U , int N_READS, bool reverse&gt; </div>
+<table class="mlabels">
+  <tr>
+  <td class="mlabels-left">
+      <table class="memname">
+        <tr>
+          <td class="memname">void load_unsafe </td>
+          <td>(</td>
+          <td class="paramtype">U</td>          <td class="paramname"><span class="paramname"><em>values</em></span>[N_READS], </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">const device T *</td>          <td class="paramname"><span class="paramname"><em>input</em></span>&#160;)</td>
+        </tr>
+      </table>
+  </td>
+  <td class="mlabels-right">
+<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
+  </tr>
+</table>
+</div><div class="memdoc">
+
+</div>
+</div>
+<a id="a7abb6ffb6c3b96b88c2a63cd4cc2f7ae" name="a7abb6ffb6c3b96b88c2a63cd4cc2f7ae"></a>
+<h2 class="memtitle"><span class="permalink"><a href="#a7abb6ffb6c3b96b88c2a63cd4cc2f7ae">&#9670;&#160;</a></span>strided_scan()</h2>
+
+<div class="memitem">
+<div class="memproto">
+<div class="memtemplate">
+template&lt;typename T , typename U , typename Op , int N_READS, bool inclusive, bool reverse&gt; </div>
+      <table class="memname">
+        <tr>
+          <td class="memname">void strided_scan </td>
+          <td>(</td>
+          <td class="paramtype">const device T *</td>          <td class="paramname"><span class="paramname"><em>in</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">device U *</td>          <td class="paramname"><span class="paramname"><em>out</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">const constant size_t &amp;</td>          <td class="paramname"><span class="paramname"><em>axis_size</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">const constant size_t &amp;</td>          <td class="paramname"><span class="paramname"><em>stride</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">const constant size_t &amp;</td>          <td class="paramname"><span class="paramname"><em>stride_blocks</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">uint3</td>          <td class="paramname"><span class="paramname"><em>gid</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">uint3</td>          <td class="paramname"><span class="paramname"><em>gsize</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">uint3</td>          <td class="paramname"><span class="paramname"><em>lid</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">uint</td>          <td class="paramname"><span class="paramname"><em>simd_lane_id</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">uint</td>          <td class="paramname"><span class="paramname"><em>simd_group_id</em></span>&#160;)</td>
+        </tr>
+      </table>
+</div><div class="memdoc">
+
+</div>
+</div>
+<a id="ae86aef08e5ebc8790031eb51eefa754c" name="ae86aef08e5ebc8790031eb51eefa754c"></a>
+<h2 class="memtitle"><span class="permalink"><a href="#ae86aef08e5ebc8790031eb51eefa754c">&#9670;&#160;</a></span>write_safe()</h2>
+
+<div class="memitem">
+<div class="memproto">
+<div class="memtemplate">
+template&lt;typename U , int N_READS, bool reverse&gt; </div>
+<table class="mlabels">
+  <tr>
+  <td class="mlabels-left">
+      <table class="memname">
+        <tr>
+          <td class="memname">void write_safe </td>
+          <td>(</td>
+          <td class="paramtype">U</td>          <td class="paramname"><span class="paramname"><em>values</em></span>[N_READS], </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">device U *</td>          <td class="paramname"><span class="paramname"><em>out</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">int</td>          <td class="paramname"><span class="paramname"><em>start</em></span>, </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">int</td>          <td class="paramname"><span class="paramname"><em>total</em></span>&#160;)</td>
+        </tr>
+      </table>
+  </td>
+  <td class="mlabels-right">
+<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
+  </tr>
+</table>
+</div><div class="memdoc">
+
+</div>
+</div>
+<a id="a8010e7bdf7a72cbd35ce7cd7ecb08e32" name="a8010e7bdf7a72cbd35ce7cd7ecb08e32"></a>
+<h2 class="memtitle"><span class="permalink"><a href="#a8010e7bdf7a72cbd35ce7cd7ecb08e32">&#9670;&#160;</a></span>write_unsafe()</h2>
+
+<div class="memitem">
+<div class="memproto">
+<div class="memtemplate">
+template&lt;typename U , int N_READS, bool reverse&gt; </div>
+<table class="mlabels">
+  <tr>
+  <td class="mlabels-left">
+      <table class="memname">
+        <tr>
+          <td class="memname">void write_unsafe </td>
+          <td>(</td>
+          <td class="paramtype">U</td>          <td class="paramname"><span class="paramname"><em>values</em></span>[N_READS], </td>
+        </tr>
+        <tr>
+          <td class="paramkey"></td>
+          <td></td>
+          <td class="paramtype">device U *</td>          <td class="paramname"><span class="paramname"><em>out</em></span>&#160;)</td>
+        </tr>
+      </table>
+  </td>
+  <td class="mlabels-right">
+<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
+  </tr>
+</table>
+</div><div class="memdoc">
+
+</div>
+</div>
+</div><!-- contents -->
+<!-- start footer part -->
+<hr class="footer"/><address class="footer"><small>
+Generated by&#160;<a href="https://www.doxygen.org/index.html"><img class="footer" src="doxygen.svg" width="104" height="31" alt="doxygen"/></a> 1.12.0
+</small></address>
+</div><!-- doc-content -->
+</body>
+</html>