From 285381b4347bd3b229e126d8caae5d55fac9024c Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Sat, 24 Feb 2024 19:24:08 -0600 Subject: [PATCH] Doc string fixes --- .../core/kernel_interface/ranges_overloads.py | 4 +- numba_dpex/kernel_api/atomic_ref.py | 51 ++++++++----------- numba_dpex/kernel_api/barrier.py | 8 +-- numba_dpex/kernel_api/ranges.py | 6 +-- 4 files changed, 29 insertions(+), 40 deletions(-) diff --git a/numba_dpex/core/kernel_interface/ranges_overloads.py b/numba_dpex/core/kernel_interface/ranges_overloads.py index f0ed884c8b..696b438d4a 100644 --- a/numba_dpex/core/kernel_interface/ranges_overloads.py +++ b/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) diff --git a/numba_dpex/kernel_api/atomic_ref.py b/numba_dpex/kernel_api/atomic_ref.py index 7d9e426aa1..b10ea77a57 100644 --- a/numba_dpex/kernel_api/atomic_ref.py +++ b/numba_dpex/kernel_api/atomic_ref.py @@ -66,7 +66,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 +78,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 +93,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 +108,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 +123,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 +138,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 +153,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,7 +166,6 @@ 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] @@ -180,22 +173,20 @@ 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 @@ -203,23 +194,21 @@ def exchange(self, val): 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 diff --git a/numba_dpex/kernel_api/barrier.py b/numba_dpex/kernel_api/barrier.py index 744262403d..3fd78fe80e 100644 --- a/numba_dpex/kernel_api/barrier.py +++ b/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. diff --git a/numba_dpex/kernel_api/ranges.py b/numba_dpex/kernel_api/ranges.py index 85cea812af..e4a131544a 100644 --- a/numba_dpex/kernel_api/ranges.py +++ b/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: