Doc string fixes

Author: Diptorup Deb

numba_dpex/core/kernel_interface/ranges_overloads.py

@@ -33,14 +33,14 @@ def codegen(context, builder, sig, args):
             range_struct.dim1 = dim1
         else:
             range_struct.dim1 = llvmir.Constant(
-                llvmir.types.IntType(64), Range.UNDEFINED_DIMENSION
+                llvmir.types.IntType(64), Range._UNDEFINED_DIMENSION
             )
 
         if not isinstance(sig.args[2], types.NoneType):
             range_struct.dim2 = dim2
         else:
             range_struct.dim2 = llvmir.Constant(
-                llvmir.types.IntType(64), Range.UNDEFINED_DIMENSION
+                llvmir.types.IntType(64), Range._UNDEFINED_DIMENSION
             )
 
         range_struct.ndim = llvmir.Constant(llvmir.types.IntType(64), typ.ndim)

numba_dpex/experimental/launcher.py

@@ -220,18 +220,19 @@ def codegen(
 
 @dpjit
 def call_kernel(kernel_fn, index_space, *kernel_args) -> None:
-    """Calls a numba_dpex.kernel decorated function from CPython or from another
-    dpjit function. Kernel execution happens in synchronous way, so the thread
-    will be blocked till the kernel done execution.
+    """Compiles and synchronously executes a kernel function.
+
+    Kernel execution happens in synchronous way, so the main thread will be
+    blocked till the kernel done execution.
 
     Args:
         kernel_fn (numba_dpex.experimental.KernelDispatcher): A
-        numba_dpex.kernel decorated function that is compiled to a
-        KernelDispatcher by numba_dpex.
-        index_space (Range | NdRange): A numba_dpex.Range or numba_dpex.NdRange
-        type object that specifies the index space for the kernel.
+            :func:`numba_dpex.experimental.kernel` decorated function that is
+            compiled to a ``KernelDispatcher``.
+        index_space (Range | NdRange): A Range or NdRange type object that
+            specifies the index space for the kernel.
         kernel_args : List of objects that are passed to the numba_dpex.kernel
-        decorated function.
+            decorated function.
     """
     _submit_kernel_sync(  # pylint: disable=E1120
         kernel_fn,
@@ -247,25 +248,32 @@ def call_kernel_async(
     dependent_events: list[dpctl.SyclEvent],
     *kernel_args
 ) -> tuple[dpctl.SyclEvent, dpctl.SyclEvent]:
-    """Calls a numba_dpex.kernel decorated function from CPython or from another
-    dpjit function. Kernel execution happens in asynchronous way, so the thread
-    will not be blocked till the kernel done execution. That means that it is
-    user responsibility to properly use any memory used by kernel until the
-    kernel execution is completed.
+    """Compiles and asynchronously executes a kernel function.
+
+    Calls a :func:`numba_dpex.experimental.kernel` decorated function
+    asynchronously from CPython or from a :func:`numba_dpex.dpjit` function. As
+    the kernel execution happens asynchronously, so the main thread will not be
+    blocked till the kernel done execution. Instead the function returns back to
+    caller a handle for an *event* to track kernel execution. It is a user's
+    responsibility to properly track kernel execution completion and not use any
+    data that may still be used by the kernel prior to the kernel's completion.
 
     Args:
-        kernel_fn (numba_dpex.experimental.KernelDispatcher): A
-        numba_dpex.kernel decorated function that is compiled to a
-        KernelDispatcher by numba_dpex.
-        index_space (Range | NdRange): A numba_dpex.Range or numba_dpex.NdRange
-        type object that specifies the index space for the kernel.
+        kernel_fn (KernelDispatcher): A
+            :func:`numba_dpex.experimental.kernel` decorated function that is
+            compiled to a ``KernelDispatcher``.
+        index_space (Range | NdRange): A Range or NdRange type object that
+            specifies the index space for the kernel.
         kernel_args : List of objects that are passed to the numba_dpex.kernel
-        decorated function.
+            decorated function.
 
     Returns:
-        pair of host event and device event. Host event represent host task
-        that releases use of any kernel argument so it can be deallocated.
-        This task may be executed only after device task is done.
+        A pair of ``dpctl.SyclEvent`` objects. The pair of events constitute of
+        a host task and an event associated with the kernel execution. The event
+        associated with the kernel execution indicates the execution status of
+        the submitted kernel function. The host task manages the lifetime of any
+        PyObject passed in as a kernel argument and automatically decrements the
+        reference count of the object on kernel execution completion.
     """
     return _submit_kernel_async(  # pylint: disable=E1120
         kernel_fn,

numba_dpex/kernel_api/atomic_ref.py

@@ -10,9 +10,10 @@
 
 
 class AtomicRef:
-    """Analogue to the ``sycl::atomic_ref`` type. 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.
+    """Analogue to the ``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.
     """
 
     def __init__(  # pylint: disable=too-many-arguments
@@ -66,7 +67,6 @@ def fetch_add(self, val):
             val : Value to be added to the object referenced by the AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
-
         """
         old = self._ref[self._index].copy()
         self._ref[self._index] += val
@@ -79,10 +79,9 @@ def fetch_sub(self, val):
 
         Args:
             val : Value to be subtracted from the object referenced by the
-            AtomicRef.
+                AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
-
         """
         old = self._ref[self._index].copy()
         self._ref[self._index] -= val
@@ -95,10 +94,9 @@ def fetch_min(self, val):
 
         Args:
             val : Value to be compared against the object referenced by the
-            AtomicRef.
+                AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
-
         """
         old = self._ref[self._index].copy()
         self._ref[self._index] = min(old, val)
@@ -111,10 +109,9 @@ def fetch_max(self, val):
 
         Args:
             val : Value to be compared against the object referenced by the
-            AtomicRef.
+                AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
-
         """
         old = self._ref[self._index].copy()
         self._ref[self._index] = max(old, val)
@@ -127,10 +124,9 @@ def fetch_and(self, val):
 
         Args:
             val : Value to be bitwise ANDed against the object referenced by
-            the AtomicRef.
+                the AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
-
         """
         old = self._ref[self._index].copy()
         self._ref[self._index] &= val
@@ -143,10 +139,9 @@ def fetch_or(self, val):
 
         Args:
             val : Value to be bitwise ORed against the object referenced by
-            the AtomicRef.
+                the AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
-
         """
         old = self._ref[self._index].copy()
         self._ref[self._index] |= val
@@ -159,7 +154,7 @@ def fetch_xor(self, val):
 
         Args:
             val : Value to be bitwise XORed against the object referenced by
-            the AtomicRef.
+                the AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
 
@@ -172,54 +167,49 @@ def load(self):
         """Loads the value of the object referenced by the AtomicRef.
 
         Returns: The value of the object referenced by the AtomicRef.
-
         """
         return self._ref[self._index]
 
     def store(self, val):
         """Stores operand ``val`` to the object referenced by the AtomicRef.
 
         Args:
-            val : Value to be stored in the object referenced by
-            the AtomicRef.
-
+            val : Value to be stored in the object referenced by the AtomicRef.
         """
         self._ref[self._index] = val
 
     def exchange(self, val):
         """Replaces the value of the object referenced by the AtomicRef
-        with value of ``val``. Returns the original value of the referenced object.
+        with value of ``val``. Returns the original value of the referenced
+        object.
 
         Args:
             val : Value to be exchanged against the object referenced by
-            the AtomicRef.
+                the AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
-
         """
         old = self._ref[self._index].copy()
         self._ref[self._index] = val
         return old
 
     def compare_exchange(self, expected, desired, expected_idx=0):
         """Compares the value of the object referenced by the AtomicRef
-        against the value of ``expected[expected_idx]``.
-        If the values are equal, replaces the value of the
-        referenced object with the value of ``desired``.
-        Otherwise assigns the original value of the
-        referenced object to ``expected[expected_idx]``.
+        against the value of ``expected[expected_idx]``. If the values are
+        equal, replaces the value of the referenced object with the value of
+        ``desired``. Otherwise assigns the original value of the referenced
+        object to ``expected[expected_idx]``.
 
         Args:
-            expected : Array containing the expected value of the
-            object referenced by the AtomicRef.
-            desired : Value that replaces the value of the object
-            referenced by the AtomicRef.
+            expected : Array containing the expected value of the object
+                referenced by the AtomicRef.
+            desired : Value that replaces the value of the object referenced by
+                the AtomicRef.
             expected_idx: Offset in `expected` array where the expected
             value of the object referenced by the AtomicRef is present.
 
-        Returns: Returns ``True`` if the comparison operation and
-        replacement operation were successful.
-
+        Returns: ``True`` if the comparison operation and replacement operation
+            were successful.
         """
         if self._ref[self._index] == expected[expected_idx]:
             self._ref[self._index] = desired

numba_dpex/kernel_api/barrier.py

@@ -2,7 +2,7 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
-"""Python functions that simulate SYCL's barrier primitives.
+"""Python functions that simulate SYCL's group_barrier function.
 """
 
 from .index_space_ids import Group
@@ -14,13 +14,13 @@ def group_barrier(group: Group, fence_scope=MemoryScope.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
-    of the group must execute the barrier construct before any work-item
+    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
     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
-    to synchronizing at the barrier, all work-items in group g execute an
+    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.
 

numba_dpex/kernel_api/index_space_ids.py

@@ -119,8 +119,10 @@ def leader(self, work_item_id):
 
 
 class Item:
-    """Analogue to the ``sycl::item`` type. Identifies an instance of the
-    function object executing at each point in an Range.
+    """Analogue to the ``sycl::item`` class.
+
+    Identifies an instance of the function object executing at each point in an
+    :class:`.Range`.
     """
 
     def __init__(self, extent: Range, index: list):
@@ -170,7 +172,7 @@ def get_range(self, idx):
 
     @property
     def dimensions(self) -> int:
-        """Returns the rank of a Item object.
+        """Returns the number of dimensions of a Item object.
 
         Returns:
             int: Number of dimensions in the Item object
@@ -179,8 +181,10 @@ def dimensions(self) -> int:
 
 
 class NdItem:
-    """Analogue to the ``sycl::nd_item`` type. Identifies an instance of the
-    function object executing at each point in an NdRange.
+    """Analogue to the ``sycl::nd_item`` class.
+
+    Identifies an instance of the function object executing at each point in an
+    :class:`.NdRange`.
     """
 
     # TODO: define group type
@@ -201,8 +205,7 @@ def get_global_id(self, idx):
         return self._global_item.get_id(idx)
 
     def get_global_linear_id(self):
-        """Get the global linear id associated with this item for all
-        dimensions.
+        """Get the linearized global id for the item for all dimensions.
 
         Returns:
             int: The global linear id.

numba_dpex/kernel_api/memory_enums.py

@@ -11,8 +11,9 @@
 
 class MemoryOrder(FlagEnum):
     """
-    An enumeration of the supported ``sycl::memory_order`` values in dpcpp. The
-    integer values of the enums is kept consistent with the corresponding
+    An enumeration of the supported ``sycl::memory_order`` values.
+
+    The integer values of the enums is kept consistent with the corresponding
     implementation in dpcpp.
 
     =====================   ============
@@ -37,8 +38,9 @@ class MemoryOrder(FlagEnum):
 
 class MemoryScope(FlagEnum):
     """
-    An enumeration of SYCL memory scope. For more details please refer to
-    SYCL 2020 specification, section 3.8.3.2
+    An enumeration of the supported ``sycl::memory_scope`` values.
+
+    For more details please refer to SYCL 2020 specification, section 3.8.3.2
 
     ===============  ============
     Memory Scope     Enum value
@@ -59,7 +61,7 @@ class MemoryScope(FlagEnum):
 
 
 class AddressSpace(FlagEnum):
-    """The address space values supported by numba_dpex.
+    """An enumeration of the supported address space values.
 
     ==================   ============
     Address space        Value

numba_dpex/kernel_api/ranges.py

@@ -27,7 +27,7 @@ class Range(tuple):
     the behavior of `sycl::range`.
     """
 
-    UNDEFINED_DIMENSION = -1
+    _UNDEFINED_DIMENSION = -1
 
     def __new__(cls, dim0, dim1=None, dim2=None):
         """Constructs a 1, 2, or 3 dimensional range.
@@ -114,7 +114,7 @@ def dim1(self) -> int:
         try:
             return self[1]
         except IndexError:
-            return Range.UNDEFINED_DIMENSION
+            return Range._UNDEFINED_DIMENSION
 
     @property
     def dim2(self) -> int:
@@ -127,7 +127,7 @@ def dim2(self) -> int:
         try:
             return self[2]
         except IndexError:
-            return Range.UNDEFINED_DIMENSION
+            return Range._UNDEFINED_DIMENSION
 
 
 class NdRange: