Merge pull request #1331 from IntelPython/experimental/local_accessors

A sycl::local_accessor-like API for numba-dpex kernel 68b1f39

dev/.buildinfo

@@ -1,4 +1,4 @@
 # Sphinx build info version 1
 # This file hashes the configuration used when building these files. When it is not found, a full rebuild will be done.
-config: 1bf4aa731798b280d7dd0d017aa0e089
+config: b1de6db79b513c9a81d857328961eb39
 tags: 645f666f9bcd5a90fca523b33c5a78b7

dev/_sources/autoapi/numba_dpex/experimental/models/index.rst.txt

@@ -25,6 +25,10 @@ Overview
      - Data model for AtomicRefType.
    * - :py:obj:`EmptyStructModel <numba_dpex.experimental.models.EmptyStructModel>`
      - Data model that does not take space. Intended to be used with types that
+   * - :py:obj:`DpctlMDLocalAccessorModel <numba_dpex.experimental.models.DpctlMDLocalAccessorModel>`
+     - Data model to represent DpctlMDLocalAccessorType.
+   * - :py:obj:`LocalAccessorModel <numba_dpex.experimental.models.LocalAccessorModel>`
+     - Data model for the LocalAccessor type when used in a host-only function.
 
 
 
@@ -59,6 +63,29 @@ Classes
 
 
 
+.. py:class:: DpctlMDLocalAccessorModel(dmm, fe_type)
+
+   Bases: :py:obj:`numba.core.datamodel.models.StructModel`
+
+   Data model to represent DpctlMDLocalAccessorType.
+
+   Must be the same structure as
+   dpctl/syclinterface/dpctl_sycl_queue_interface.h::MDLocalAccessor.
+
+   Structure intended to be used only on host side of the kernel call.
+
+
+
+
+.. py:class:: LocalAccessorModel(dmm, fe_type)
+
+   Bases: :py:obj:`numba.core.datamodel.models.StructModel`
+
+   Data model for the LocalAccessor type when used in a host-only function.
+
+
+
+
 
 Attributes
 ----------

dev/_sources/autoapi/numba_dpex/experimental/typeof/index.rst.txt

@@ -30,6 +30,8 @@ Overview
      - Registers the type inference implementation function for a
    * - :py:obj:`typeof_nditem <numba_dpex.experimental.typeof.typeof_nditem>`\ (val, c)
      - Registers the type inference implementation function for a
+   * - :py:obj:`typeof_local_accessor <numba_dpex.experimental.typeof.typeof_local_accessor>`\ (val, c)
+     - Returns a ``numba_dpex.experimental.dpctpp_types.LocalAccessorType``
 
 
 
@@ -85,6 +87,17 @@ Functions
        instance.
 
 
+.. py:function:: typeof_local_accessor(val: numba_dpex.kernel_api.LocalAccessor, c) -> numba_dpex.core.types.kernel_api.local_accessor.LocalAccessorType
+
+   Returns a ``numba_dpex.experimental.dpctpp_types.LocalAccessorType``
+   instance for a Python LocalAccessor object.
+   :param val: Instance of the LocalAccessor type.
+   :type val: LocalAccessor
+   :param c: Numba typing context used for type inference.
+
+   Returns: LocalAccessorType object corresponding to the LocalAccessor object.
+
+
 
 
 

dev/_sources/autoapi/numba_dpex/kernel_api/index.rst.txt

@@ -29,6 +29,8 @@ Overview
      - Analogue to the ``sycl::item`` class.
    * - :py:obj:`NdItem <numba_dpex.kernel_api.NdItem>`
      - Analogue to the ``sycl::nd_item`` class.
+   * - :py:obj:`LocalAccessor <numba_dpex.kernel_api.LocalAccessor>`
+     - The ``LocalAccessor`` class is analogous to SYCL's ``local_accessor``
    * - :py:obj:`AddressSpace <numba_dpex.kernel_api.AddressSpace>`
      - An enumeration of the supported address space values.
    * - :py:obj:`MemoryOrder <numba_dpex.kernel_api.MemoryOrder>`
@@ -468,6 +470,16 @@ Classes
 
 
 
+.. py:class:: LocalAccessor(shape, dtype)
+
+   The ``LocalAccessor`` class is analogous to SYCL's ``local_accessor``
+   class. The class acts a s proxy to allocating device local memory and
+   accessing that memory from within a :func:`numba_dpex.kernel` decorated
+   function.
+
+
+
+
 .. py:class:: AddressSpace
 
    Bases: :py:obj:`numba_dpex.kernel_api.flag_enum.FlagEnum`

dev/_sources/autoapi/numba_dpex/kernel_api/local_accessor/index.rst.txt

@@ -0,0 +1,46 @@
+
+:orphan:
+
+numba_dpex.kernel_api.local_accessor
+====================================
+
+.. py:module:: numba_dpex.kernel_api.local_accessor
+
+.. autoapi-nested-parse::
+
+   Implements a Python analogue to SYCL's local_accessor class. The class is
+   intended to be used in pure Python code when prototyping a kernel function
+   and to be passed to an actual kernel function for local memory allocation.
+
+
+
+Overview
+--------
+
+.. list-table:: Classes
+   :header-rows: 0
+   :widths: auto
+   :class: summarytable
+
+   * - :py:obj:`LocalAccessor <numba_dpex.kernel_api.local_accessor.LocalAccessor>`
+     - The ``LocalAccessor`` class is analogous to SYCL's ``local_accessor``
+
+
+
+
+Classes
+-------
+
+.. py:class:: LocalAccessor(shape, dtype)
+
+   The ``LocalAccessor`` class is analogous to SYCL's ``local_accessor``
+   class. The class acts a s proxy to allocating device local memory and
+   accessing that memory from within a :func:`numba_dpex.kernel` decorated
+   function.
+
+
+
+
+
+
+

dev/autoapi/numba_dpex/experimental/models/index.html

@@ -261,6 +261,12 @@ <h2>Overview<a class="headerlink" href="#overview" title="Link to this heading">
 <tr class="row-even"><td><p><a class="reference internal" href="#numba_dpex.experimental.models.EmptyStructModel" title="numba_dpex.experimental.models.EmptyStructModel"><code class="xref py py-obj docutils literal notranslate"><span class="pre">EmptyStructModel</span></code></a></p></td>
 <td><p>Data model that does not take space. Intended to be used with types that</p></td>
 </tr>
+<tr class="row-odd"><td><p><a class="reference internal" href="#numba_dpex.experimental.models.DpctlMDLocalAccessorModel" title="numba_dpex.experimental.models.DpctlMDLocalAccessorModel"><code class="xref py py-obj docutils literal notranslate"><span class="pre">DpctlMDLocalAccessorModel</span></code></a></p></td>
+<td><p>Data model to represent DpctlMDLocalAccessorType.</p></td>
+</tr>
+<tr class="row-even"><td><p><a class="reference internal" href="#numba_dpex.experimental.models.LocalAccessorModel" title="numba_dpex.experimental.models.LocalAccessorModel"><code class="xref py py-obj docutils literal notranslate"><span class="pre">LocalAccessorModel</span></code></a></p></td>
+<td><p>Data model for the LocalAccessor type when used in a host-only function.</p></td>
+</tr>
 </tbody>
 </table>
 </div>
@@ -292,6 +298,23 @@ <h2>Classes<a class="headerlink" href="#classes" title="Link to this heading">#<
 are presented only at typing stage and not represented physically.</p>
 </dd></dl>
 
+<dl class="py class">
+<dt class="sig sig-object py" id="numba_dpex.experimental.models.DpctlMDLocalAccessorModel">
+<em class="property"><span class="pre">class</span><span class="w"> </span></em><span class="sig-name descname"><span class="pre">DpctlMDLocalAccessorModel</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">dmm</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">fe_type</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#numba_dpex.experimental.models.DpctlMDLocalAccessorModel" title="Link to this definition">#</a></dt>
+<dd><p>Bases: <code class="xref py py-obj docutils literal notranslate"><span class="pre">numba.core.datamodel.models.StructModel</span></code></p>
+<p>Data model to represent DpctlMDLocalAccessorType.</p>
+<p>Must be the same structure as
+dpctl/syclinterface/dpctl_sycl_queue_interface.h::MDLocalAccessor.</p>
+<p>Structure intended to be used only on host side of the kernel call.</p>
+</dd></dl>
+
+<dl class="py class">
+<dt class="sig sig-object py" id="numba_dpex.experimental.models.LocalAccessorModel">
+<em class="property"><span class="pre">class</span><span class="w"> </span></em><span class="sig-name descname"><span class="pre">LocalAccessorModel</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">dmm</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">fe_type</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#numba_dpex.experimental.models.LocalAccessorModel" title="Link to this definition">#</a></dt>
+<dd><p>Bases: <code class="xref py py-obj docutils literal notranslate"><span class="pre">numba.core.datamodel.models.StructModel</span></code></p>
+<p>Data model for the LocalAccessor type when used in a host-only function.</p>
+</dd></dl>
+
 </section>
 <section id="attributes">
 <h2>Attributes<a class="headerlink" href="#attributes" title="Link to this heading">#</a></h2>
@@ -348,6 +371,8 @@ <h2>Attributes<a class="headerlink" href="#attributes" title="Link to this headi
 <li><a class="reference internal" href="#classes">Classes</a><ul>
 <li><a class="reference internal" href="#numba_dpex.experimental.models.AtomicRefModel"><code class="docutils literal notranslate"><span class="pre">AtomicRefModel</span></code></a></li>
 <li><a class="reference internal" href="#numba_dpex.experimental.models.EmptyStructModel"><code class="docutils literal notranslate"><span class="pre">EmptyStructModel</span></code></a></li>
+<li><a class="reference internal" href="#numba_dpex.experimental.models.DpctlMDLocalAccessorModel"><code class="docutils literal notranslate"><span class="pre">DpctlMDLocalAccessorModel</span></code></a></li>
+<li><a class="reference internal" href="#numba_dpex.experimental.models.LocalAccessorModel"><code class="docutils literal notranslate"><span class="pre">LocalAccessorModel</span></code></a></li>
 </ul>
 </li>
 <li><a class="reference internal" href="#attributes">Attributes</a><ul>

dev/autoapi/numba_dpex/experimental/typeof/index.html

@@ -267,6 +267,9 @@ <h2>Overview<a class="headerlink" href="#overview" title="Link to this heading">
 <tr class="row-even"><td><p><a class="reference internal" href="#numba_dpex.experimental.typeof.typeof_nditem" title="numba_dpex.experimental.typeof.typeof_nditem"><code class="xref py py-obj docutils literal notranslate"><span class="pre">typeof_nditem</span></code></a>(val, c)</p></td>
 <td><p>Registers the type inference implementation function for a</p></td>
 </tr>
+<tr class="row-odd"><td><p><a class="reference internal" href="#numba_dpex.experimental.typeof.typeof_local_accessor" title="numba_dpex.experimental.typeof.typeof_local_accessor"><code class="xref py py-obj docutils literal notranslate"><span class="pre">typeof_local_accessor</span></code></a>(val, c)</p></td>
+<td><p>Returns a <code class="docutils literal notranslate"><span class="pre">numba_dpex.experimental.dpctpp_types.LocalAccessorType</span></code></p></td>
+</tr>
 </tbody>
 </table>
 </div>
@@ -346,6 +349,17 @@ <h2>Functions<a class="headerlink" href="#functions" title="Link to this heading
 </dl>
 </dd></dl>
 
+<dl class="py function">
+<dt class="sig sig-object py" id="numba_dpex.experimental.typeof.typeof_local_accessor">
+<span class="sig-name descname"><span class="pre">typeof_local_accessor</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">val</span></span><span class="p"><span class="pre">:</span></span><span class="w"> </span><span class="n"><a class="reference internal" href="../../kernel_api/index.html#numba_dpex.kernel_api.LocalAccessor" title="numba_dpex.kernel_api.LocalAccessor"><span class="pre">numba_dpex.kernel_api.LocalAccessor</span></a></span></em>, <em class="sig-param"><span class="n"><span class="pre">c</span></span></em><span class="sig-paren">)</span> <span class="sig-return"><span class="sig-return-icon">&#x2192;</span> <span class="sig-return-typehint"><span class="pre">numba_dpex.core.types.kernel_api.local_accessor.LocalAccessorType</span></span></span><a class="headerlink" href="#numba_dpex.experimental.typeof.typeof_local_accessor" title="Link to this definition">#</a></dt>
+<dd><p>Returns a <code class="docutils literal notranslate"><span class="pre">numba_dpex.experimental.dpctpp_types.LocalAccessorType</span></code>
+instance for a Python LocalAccessor object.
+:param val: Instance of the LocalAccessor type.
+:type val: LocalAccessor
+:param c: Numba typing context used for type inference.</p>
+<p>Returns: LocalAccessorType object corresponding to the LocalAccessor object.</p>
+</dd></dl>
+
 </section>
 </section>
 
@@ -396,6 +410,7 @@ <h2>Functions<a class="headerlink" href="#functions" title="Link to this heading
 <li><a class="reference internal" href="#numba_dpex.experimental.typeof.typeof_group"><code class="docutils literal notranslate"><span class="pre">typeof_group()</span></code></a></li>
 <li><a class="reference internal" href="#numba_dpex.experimental.typeof.typeof_item"><code class="docutils literal notranslate"><span class="pre">typeof_item()</span></code></a></li>
 <li><a class="reference internal" href="#numba_dpex.experimental.typeof.typeof_nditem"><code class="docutils literal notranslate"><span class="pre">typeof_nditem()</span></code></a></li>
+<li><a class="reference internal" href="#numba_dpex.experimental.typeof.typeof_local_accessor"><code class="docutils literal notranslate"><span class="pre">typeof_local_accessor()</span></code></a></li>
 </ul>
 </li>
 </ul>

dev/autoapi/numba_dpex/kernel_api/index.html

@@ -269,22 +269,25 @@ <h2>Overview<a class="headerlink" href="#overview" title="Link to this heading">
 <tr class="row-even"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.NdItem" title="numba_dpex.kernel_api.NdItem"><code class="xref py py-obj docutils literal notranslate"><span class="pre">NdItem</span></code></a></p></td>
 <td><p>Analogue to the <code class="docutils literal notranslate"><span class="pre">sycl::nd_item</span></code> class.</p></td>
 </tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.AddressSpace" title="numba_dpex.kernel_api.AddressSpace"><code class="xref py py-obj docutils literal notranslate"><span class="pre">AddressSpace</span></code></a></p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.LocalAccessor" title="numba_dpex.kernel_api.LocalAccessor"><code class="xref py py-obj docutils literal notranslate"><span class="pre">LocalAccessor</span></code></a></p></td>
+<td><p>The <code class="docutils literal notranslate"><span class="pre">LocalAccessor</span></code> class is analogous to SYCL’s <code class="docutils literal notranslate"><span class="pre">local_accessor</span></code></p></td>
+</tr>
+<tr class="row-even"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.AddressSpace" title="numba_dpex.kernel_api.AddressSpace"><code class="xref py py-obj docutils literal notranslate"><span class="pre">AddressSpace</span></code></a></p></td>
 <td><p>An enumeration of the supported address space values.</p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.MemoryOrder" title="numba_dpex.kernel_api.MemoryOrder"><code class="xref py py-obj docutils literal notranslate"><span class="pre">MemoryOrder</span></code></a></p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.MemoryOrder" title="numba_dpex.kernel_api.MemoryOrder"><code class="xref py py-obj docutils literal notranslate"><span class="pre">MemoryOrder</span></code></a></p></td>
 <td><p>An enumeration of the supported <code class="docutils literal notranslate"><span class="pre">sycl::memory_order</span></code> values.</p></td>
 </tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.MemoryScope" title="numba_dpex.kernel_api.MemoryScope"><code class="xref py py-obj docutils literal notranslate"><span class="pre">MemoryScope</span></code></a></p></td>
+<tr class="row-even"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.MemoryScope" title="numba_dpex.kernel_api.MemoryScope"><code class="xref py py-obj docutils literal notranslate"><span class="pre">MemoryScope</span></code></a></p></td>
 <td><p>An enumeration of the supported <code class="docutils literal notranslate"><span class="pre">sycl::memory_scope</span></code> values.</p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.PrivateArray" title="numba_dpex.kernel_api.PrivateArray"><code class="xref py py-obj docutils literal notranslate"><span class="pre">PrivateArray</span></code></a></p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.PrivateArray" title="numba_dpex.kernel_api.PrivateArray"><code class="xref py py-obj docutils literal notranslate"><span class="pre">PrivateArray</span></code></a></p></td>
 <td><p>The <code class="docutils literal notranslate"><span class="pre">PrivateArray</span></code> class is an simple version of array intended to be used</p></td>
 </tr>
-<tr class="row-odd"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.NdRange" title="numba_dpex.kernel_api.NdRange"><code class="xref py py-obj docutils literal notranslate"><span class="pre">NdRange</span></code></a></p></td>
+<tr class="row-even"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.NdRange" title="numba_dpex.kernel_api.NdRange"><code class="xref py py-obj docutils literal notranslate"><span class="pre">NdRange</span></code></a></p></td>
 <td><p>A class to encapsulate all kernel launch parameters.</p></td>
 </tr>
-<tr class="row-even"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.Range" title="numba_dpex.kernel_api.Range"><code class="xref py py-obj docutils literal notranslate"><span class="pre">Range</span></code></a></p></td>
+<tr class="row-odd"><td><p><a class="reference internal" href="#numba_dpex.kernel_api.Range" title="numba_dpex.kernel_api.Range"><code class="xref py py-obj docutils literal notranslate"><span class="pre">Range</span></code></a></p></td>
 <td><p>A data structure to encapsulate a single kernel launch parameter.</p></td>
 </tr>
 </tbody>
@@ -828,6 +831,15 @@ <h2>Classes<a class="headerlink" href="#classes" title="Link to this heading">#<
 
 </dd></dl>
 
+<dl class="py class">
+<dt class="sig sig-object py" id="numba_dpex.kernel_api.LocalAccessor">
+<em class="property"><span class="pre">class</span><span class="w"> </span></em><span class="sig-name descname"><span class="pre">LocalAccessor</span></span><span class="sig-paren">(</span><em class="sig-param"><span class="n"><span class="pre">shape</span></span></em>, <em class="sig-param"><span class="n"><span class="pre">dtype</span></span></em><span class="sig-paren">)</span><a class="headerlink" href="#numba_dpex.kernel_api.LocalAccessor" title="Link to this definition">#</a></dt>
+<dd><p>The <code class="docutils literal notranslate"><span class="pre">LocalAccessor</span></code> class is analogous to SYCL’s <code class="docutils literal notranslate"><span class="pre">local_accessor</span></code>
+class. The class acts a s proxy to allocating device local memory and
+accessing that memory from within a <code class="xref py py-func docutils literal notranslate"><span class="pre">numba_dpex.kernel()</span></code> decorated
+function.</p>
+</dd></dl>
+
 <dl class="py class">
 <dt class="sig sig-object py" id="numba_dpex.kernel_api.AddressSpace">
 <em class="property"><span class="pre">class</span><span class="w"> </span></em><span class="sig-name descname"><span class="pre">AddressSpace</span></span><a class="headerlink" href="#numba_dpex.kernel_api.AddressSpace" title="Link to this definition">#</a></dt>
@@ -1383,6 +1395,7 @@ <h2>Functions<a class="headerlink" href="#functions" title="Link to this heading
 <li><a class="reference internal" href="#numba_dpex.kernel_api.NdItem.get_group"><code class="docutils literal notranslate"><span class="pre">NdItem.get_group()</span></code></a></li>
 </ul>
 </li>
+<li><a class="reference internal" href="#numba_dpex.kernel_api.LocalAccessor"><code class="docutils literal notranslate"><span class="pre">LocalAccessor</span></code></a></li>
 <li><a class="reference internal" href="#numba_dpex.kernel_api.AddressSpace"><code class="docutils literal notranslate"><span class="pre">AddressSpace</span></code></a><ul>
 <li><a class="reference internal" href="#numba_dpex.kernel_api.AddressSpace.PRIVATE"><code class="docutils literal notranslate"><span class="pre">AddressSpace.PRIVATE</span></code></a></li>
 <li><a class="reference internal" href="#numba_dpex.kernel_api.AddressSpace.GLOBAL"><code class="docutils literal notranslate"><span class="pre">AddressSpace.GLOBAL</span></code></a></li>