IntelPython · ZzEeKkAa · Mar 29, 2024 · Mar 28, 2024
diff --git a/docs/source/conf.py b/docs/source/conf.py
@@ -5,10 +5,14 @@
 # coding: utf-8
 # Configuration file for the Sphinx documentation builder.
 
-import numba_dpex
-
 # -- Project information -----------------------------------------------------
 
+import sys
+
+sys.path.append(".")
+
+from sycl_spec_links import sycl_ext_links  # noqa E402
+
 project = "numba-dpex"
 copyright = "2020-2024, Intel Corporation"
 author = "Intel Corporation"
@@ -45,6 +49,8 @@
 # This pattern also affects html_static_path and html_extra_path.
 exclude_patterns = []
 
+extlinks = {}
+extlinks.update(sycl_ext_links)
 
 # -- Options for HTML output -------------------------------------------------
 
@@ -72,14 +78,7 @@
 # so a file named "default.css" will overwrite the builtin "default.css".
 html_static_path = []
 
-html_sidebars = {
-    # "**": [
-    #     "globaltoc.html",
-    #     "sourcelink.html",
-    #     "searchbox.html",
-    #     "relations.html",
-    # ],
-}
+html_sidebars = {}
 
 html_show_sourcelink = False
 
@@ -88,28 +87,8 @@
 todo_link_only = True
 
 # -- InterSphinx configuration: looks for objects in external projects -----
-# Add here external classes you want to link from Intel SDC documentation
-# Each entry of the dictionary has the following format:
-#      'class name': ('link to object.inv file for that class', None)
-# intersphinx_mapping = {
-#    'pandas': ('https://pandas.pydata.org/pandas-docs/stable/', None),
-#    'python': ('http://docs.python.org/2', None),
-#    'numpy': ('http://docs.scipy.org/doc/numpy', None)
-# }
 intersphinx_mapping = {}
 
-# -- Napoleon extension configuration (Numpy and Google docstring options) -------
-# napoleon_google_docstring = True
-# napoleon_numpy_docstring = True
-# napoleon_include_init_with_doc = True
-# napoleon_include_private_with_doc = True
-# napoleon_include_special_with_doc = True
-# napoleon_use_admonition_for_examples = False
-# napoleon_use_admonition_for_notes = False
-# napoleon_use_admonition_for_references = False
-# napoleon_use_ivar = False
-# napoleon_use_param = True
-# napoleon_use_rtype = True
 
 # -- Prepend module name to an object name or not -----------------------------------
 add_module_names = False

diff --git a/docs/source/sycl_spec_links.py b/docs/source/sycl_spec_links.py
@@ -0,0 +1,56 @@
+# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Links to the SYCL 2020 specification that are used in docstring.
+
+The module provides a dictionary in the format needed by the sphinx.ext.extlinks
+extension.
+"""
+
+sycl_ext_links = {
+    "sycl_item": (
+        "https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:item.class%s",
+        None,
+    ),
+    "sycl_group": (
+        "https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#group-class%s",
+        None,
+    ),
+    "sycl_nditem": (
+        "https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:nditem.class%s",
+        None,
+    ),
+    "sycl_ndrange": (
+        "https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#subsubsec:nd-range-class%s",
+        None,
+    ),
+    "sycl_range": (
+        "https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#range-class%s",
+        None,
+    ),
+    "sycl_atomic_ref": (
+        "https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:atomic-references%s",
+        None,
+    ),
+    "sycl_local_accessor": (
+        "https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:accessor.local%s",
+        None,
+    ),
+    "sycl_private_memory": (
+        "https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_parallel_for_hierarchical_invoke%s",
+        None,
+    ),
+    "sycl_memory_scope": (
+        "https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:memory-scope%s",
+        None,
+    ),
+    "sycl_memory_order": (
+        "https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:memory-ordering%s",
+        None,
+    ),
+    "sycl_addr_space": (
+        "https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_address_space_classes%s",
+        None,
+    ),
+}
diff --git a/numba_dpex/kernel_api/atomic_fence.py b/numba_dpex/kernel_api/atomic_fence.py
@@ -4,20 +4,25 @@
 
 """Python functions that simulate SYCL's atomic_fence primitives.
 """
+from .memory_enums import MemoryOrder, MemoryScope
 
 
-def atomic_fence(memory_order, memory_scope):  # pylint: disable=unused-argument
-    """The function for performing memory fence across all work-items.
-    Modeled after ``sycl::atomic_fence`` function.
-    It provides control over re-ordering of memory load
-    and store operations. The ``atomic_fence`` function acts as a
-    fence across all work-items and devices specified by a
-    memory_scope argument.
+def atomic_fence(
+    memory_order: MemoryOrder, memory_scope: MemoryScope
+):  # pylint: disable=unused-argument
+    """Performs a memory fence operations across all work-items.
 
-    Args:
-    memory_order: The memory synchronization order.
+    The function is equivalent to the ``sycl::atomic_fence`` function and
+    controls the order of memory accesses (loads and stores) by individual
+    work-items.
+
+    .. important::
+        The function is a no-op during CPython execution and only available in
+        JIT compiled mode of execution.
 
-    memory_scope: The set of work-items and devices to which
-    the memory ordering constraints apply.
+    Args:
+        memory_order (MemoryOrder): The memory synchronization order.
+        memory_scope (MemoryScope): The set of work-items and devices to which
+            the memory ordering constraints apply.
 
     """
diff --git a/numba_dpex/kernel_api/atomic_ref.py b/numba_dpex/kernel_api/atomic_ref.py
@@ -10,7 +10,7 @@
 
 
 class AtomicRef:
-    """Analogue to the ``sycl::atomic_ref`` class.
+    """Analogue to the :sycl_atomic_ref:`sycl::atomic_ref <>` class.
 
     An atomic reference is a view into a data container that can be then updated
     atomically using any of the ``fetch_*`` member functions of the class.

diff --git a/numba_dpex/kernel_api/barrier.py b/numba_dpex/kernel_api/barrier.py
@@ -9,24 +9,34 @@
 from .memory_enums import MemoryScope
 
 
-def group_barrier(group: Group, fence_scope=MemoryScope.WORK_GROUP):
-    """Performs a barrier operation across all work-items in a work group.
+def group_barrier(
+    group: Group, fence_scope: MemoryScope = MemoryScope.WORK_GROUP
+):
+    """Performs a barrier operation across all work-items in a work-group.
 
-    The function is modeled after the ``sycl::group_barrier`` function. It
-    synchronizes work within a group of work items. All the work-items
+    The function is equivalent to the ``sycl::group_barrier`` function. It
+    synchronizes work within a group of work-items. All the work-items
     of the group must execute the barrier call before any work-item
     continues execution beyond the barrier.
 
-    The ``group_barrier`` performs mem-fence operations ensuring that memory
+    The ``group_barrier`` performs a memory fence operation ensuring that memory
     accesses issued before the barrier are not re-ordered with those issued
-    after the barrier: all work-items in group G execute a release fence prior
+    after the barrier. All work-items in group G execute a release fence prior
     to synchronizing at the barrier, all work-items in group G execute an
     acquire fence afterwards, and there is an implicit synchronization of these
     fences as if provided by an explicit atomic operation on an atomic object.
 
+    .. important::
+        The function is not implemented yet for pure CPython execution and is
+        only supported in JIT compiled mode of execution.
+
     Args:
-        fence_scope (optional): scope of any memory consistency
-        operations that are performed by the barrier.
+        group (Group): Indicates the work-group inside which the barrier is to
+            be executed.
+        fence_scope (MemoryScope) (optional): scope of any memory
+            consistency operations that are performed by the barrier.
+    Raises:
+        NotImplementedError: When the function is called directly from Python.
     """
 
     # TODO: A pure Python simulation of a group_barrier will be added later.

diff --git a/numba_dpex/kernel_api/index_space_ids.py b/numba_dpex/kernel_api/index_space_ids.py
@@ -11,7 +11,15 @@
 
 
 class Group:
-    """Analogue to the ``sycl::group`` type."""
+    # pylint: disable=line-too-long
+    """Analogue to the :sycl_group:`sycl::group <>` class.
+
+    Represents a particular work-group within a parallel execution and
+    provides API to extract various properties of the work-group. An instance
+    of the class is not user-constructible. Users should use
+    :func:`numba_dpex.kernel_api.NdItem.get_group` to access the Group to which
+    a work-item belongs.
+    """
 
     def __init__(
         self,
@@ -27,12 +35,20 @@ def __init__(
         self._leader = False
 
     def get_group_id(self, dim):
-        """Returns the index of the work-group within the global nd-range for
-        specified dimension.
+        """Returns a specific coordinate of the multi-dimensional index of a group.
 
         Since the work-items in a work-group have a defined position within the
         global nd-range, the returned group id can be used along with the local
         id to uniquely identify the work-item in the global nd-range.
+
+        Args:
+            dim (int): An integral value between (1..3) for which the group
+                index is returned.
+        Returns:
+            int: The coordinate for the ``dim`` dimension for the group's
+            multi-dimensional index within an nd-range.
+        Raises:
+            ValueError: If the ``dim`` argument is not in the (1..3) interval.
         """
         if dim > len(self._index) - 1:
             raise ValueError(
@@ -41,7 +57,12 @@ def get_group_id(self, dim):
         return self._index[dim]
 
     def get_group_linear_id(self):
-        """Returns a linearized version of the work-group index."""
+        """Returns a linearized version of the work-group index.
+
+        Returns:
+            int: The linearized index for the group's position within an
+            nd-range.
+        """
         if self.dimensions == 1:
             return self.get_group_id(0)
         if self.dimensions == 2:
@@ -59,28 +80,48 @@ def get_group_linear_id(self):
         )
 
     def get_group_range(self, dim):
-        """Returns a the extent of the range representing the number of groups
-        in the nd-range for a specified dimension.
+        """Returns the extent of the range of groups in an nd-range for given dimension.
+
+        Args:
+            dim (int): An integral value between (1..3) for which the group
+                index is returned.
+        Returns:
+            int: The extent of group range for the specified dimension.
         """
         return self._group_range[dim]
 
     def get_group_linear_range(self):
-        """Return the total number of work-groups in the nd_range."""
+        """Returns the total number of work-groups in the nd_range.
+
+        Returns:
+            int: Returns the number of groups in a parallel execution of an
+            nd-range kernel.
+        """
         num_wg = 1
         for i in range(self.dimensions):
             num_wg *= self.get_group_range(i)
 
         return num_wg
 
     def get_local_range(self, dim):
-        """Returns the extent of the SYCL range representing all dimensions
-        of the local range for a specified dimension. This local range may
-        have been provided by the programmer, or chosen by the SYCL runtime.
+        """Returns the extent of the range of work-items in a work-group for given dimension.
+
+        Args:
+            dim (int): An integral value between (1..3) for which the group
+                index is returned.
+        Returns:
+            int: The extent of the local work-item range for the specified
+            dimension.
         """
         return self._local_range[dim]
 
     def get_local_linear_range(self):
-        """Return the total number of work-items in the work-group."""
+        """Return the total number of work-items in the work-group.
+
+        Returns:
+            int: Returns the linearized size of the local range inside an
+            nd-range.
+        """
         num_wi = 1
         for i in range(self.dimensions):
             num_wi *= self.get_local_range(i)
@@ -89,24 +130,22 @@ def get_local_linear_range(self):
 
     @property
     def leader(self):
-        """Return true for exactly one work-item in the work-group, if the
-        calling work-item is the leader of the work-group, and false for all
-        other work-items in the work-group.
+        """Return true if the caller work-item is the leader of the work-group.
 
         The leader of the work-group is determined during construction of the
         work-group, and is invariant for the lifetime of the work-group. The
         leader of the work-group is guaranteed to be the work-item with a
         local id of 0.
 
-
         Returns:
             bool: If the work item is the designated leader of the
         """
         return self._leader
 
     @property
     def dimensions(self) -> int:
-        """Returns the rank of a Group object.
+        """Returns the dimensionality of the range to which the work-group belongs.
+
         Returns:
             int: Number of dimensions in the Group object
         """
@@ -119,22 +158,21 @@ def leader(self, work_item_id):
 
 
 class Item:
-    """Analogue to the ``sycl::item`` class.
+    """Analogue to the :sycl_item:`sycl::item <>` class.
 
-    Identifies an instance of the function object executing at each point in an
-    :class:`.Range`.
+    Identifies the work-item in a parallel execution of a kernel launched with
+    the :class:`.Range` index-space class.
     """
 
     def __init__(self, extent: Range, index: list):
         self._extent = extent
         self._index = index
 
     def get_linear_id(self):
-        """Get the linear id associated with this item for all dimensions.
-        Original implementation could be found at ``sycl::item_base`` class.
+        """Returns the linear id associated with this item for all dimensions.
 
         Returns:
-            int: The linear id.
+            int: The linear id of the work item in the global range.
         """
         if self.dimensions == 1:
             return self.get_id(0)
@@ -181,7 +219,7 @@ def dimensions(self) -> int:
 
 
 class NdItem:
-    """Analogue to the ``sycl::nd_item`` class.
+    """Analogue to the :sycl_nditem:`sycl::nd_item <>` class.
 
     Identifies an instance of the function object executing at each point in an
     :class:`.NdRange`.