Skip to content

Commit

Permalink
Doc string fixes
Browse files Browse the repository at this point in the history
  • Loading branch information
Diptorup Deb committed Feb 25, 2024
1 parent 3966474 commit 4a7419a
Show file tree
Hide file tree
Showing 4 changed files with 29 additions and 40 deletions.
4 changes: 2 additions & 2 deletions numba_dpex/core/kernel_interface/ranges_overloads.py
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
51 changes: 20 additions & 31 deletions numba_dpex/kernel_api/atomic_ref.py
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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
Expand All @@ -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)
Expand All @@ -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)
Expand All @@ -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
Expand All @@ -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
Expand All @@ -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.
Expand All @@ -172,54 +166,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
Expand Down
8 changes: 4 additions & 4 deletions numba_dpex/kernel_api/barrier.py
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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.
Expand Down
6 changes: 3 additions & 3 deletions numba_dpex/kernel_api/ranges.py
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down Expand Up @@ -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:
Expand All @@ -127,7 +127,7 @@ def dim2(self) -> int:
try:
return self[2]
except IndexError:
return Range.UNDEFINED_DIMENSION
return Range._UNDEFINED_DIMENSION


class NdRange:
Expand Down

0 comments on commit 4a7419a

Please sign in to comment.