Adds a mock class for sycl::local_accessor

Author: Diptorup Deb

numba_dpex/kernel_api/__init__.py

@@ -14,21 +14,25 @@
 from .barrier import group_barrier
 from .index_space_ids import Group, Item, NdItem
 from .launcher import call_kernel
+from .local_accessor import LocalAccessor
 from .memory_enums import AddressSpace, MemoryOrder, MemoryScope
 from .private_array import PrivateArray
 from .ranges import NdRange, Range
 
 __all__ = [
+    "call_kernel",
+    "group_barrier",
     "AddressSpace",
     "atomic_fence",
     "AtomicRef",
+    "Group",
+    "Item",
+    "LocalAccessor",
     "MemoryOrder",
     "MemoryScope",
+    "NdItem",
     "NdRange",
     "Range",
-    "Group",
-    "NdItem",
-    "Item",
     "PrivateArray",
     "group_barrier",
     "call_kernel",

numba_dpex/kernel_api/local_accessor.py

@@ -0,0 +1,101 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""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.
+"""
+import numpy
+
+
+class LocalAccessor:
+    """
+    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.
+    """
+
+    def _verify_positive_integral_list(self, ls):
+        """Checks if all members of a list are positive integers."""
+
+        ret = False
+        try:
+            ret = all(int(val) > 0 for val in ls)
+        except ValueError:
+            pass
+
+        return ret
+
+    def __init__(self, shape, dtype) -> None:
+        """Creates a new LocalAccessor instance of the given shape and dtype."""
+
+        if not isinstance(shape, (list, tuple)):
+            if hasattr(shape, "tolist"):
+                fn = getattr(shape, "tolist")
+                if callable(fn):
+                    self._shape = shape.tolist()
+            else:
+                try:
+                    self._shape = [
+                        shape,
+                    ]
+                except Exception as e:
+                    raise TypeError(
+                        "Argument shape must a non-negative integer, "
+                        "or a list/tuple of such integers."
+                    ) from e
+        else:
+            self._shape = list(shape)
+
+        # Make sure shape is made up a supported types
+        if not self._verify_positive_integral_list(self._shape):
+            raise TypeError(
+                "Argument shape must a non-negative integer, "
+                "or a list/tuple of such integers."
+            )
+
+        # Make sure shape has a rank between (1..3)
+        if len(self._shape) < 1 or len(self._shape) > 3:
+            raise TypeError("LocalAccessor can only have up to 3 dimensions.")
+
+        self._dtype = dtype
+
+        if self._dtype not in [
+            numpy.float32,
+            numpy.float64,
+            numpy.int32,
+            numpy.int64,
+        ]:
+            raise TypeError(
+                f"Argument dtype {dtype} is not supported. numpy.float32, "
+                "numpy.float64, numpy.int32, numpy.int64 are the currently "
+                "supported dtypes."
+            )
+
+        self._data = numpy.empty(self._shape, dtype=self._dtype)
+
+    def __getitem__(self, idx_obj):
+        """Returns the value stored at the position represented by idx_obj in
+        the self._data ndarray.
+        """
+
+        return self._data[idx_obj]
+
+    def __setitem__(self, idx_obj, val):
+        """Assigns a new value to the position represented by idx_obj in
+        the self._data ndarray.
+        """
+
+        self._data[idx_obj] = val
+
+    @property
+    def data(self):
+        """Returns the numpy.ndarray that is used as the internal storage for
+        a LocalAccessor object.
+
+        Returns:
+            numpy.ndarray: A numpy.ndarray object.
+        """
+        return self._data