Merge pull request #1239 from IntelPython/experimental/memory_enum_classes

Adds memory enum classes to the experimental module

Author: ZzEeKkAa

numba_dpex/experimental/kernel_iface/__init__.py

@@ -0,0 +1,13 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""The kernel_iface provides a set of Python classes and functions that are
+analogous to the C++ SYCL API. The kernel_iface API is meant to allow
+prototyping SYCL-like kernels in pure Python before compiling them using
+numba_dpex.kernel.
+"""
+
+from .memory_enums import AddressSpace, MemoryOrder, MemoryScope
+
+__all__ = ["AddressSpace", "MemoryOrder", "MemoryScope"]

numba_dpex/experimental/kernel_iface/memory_enums.py

@@ -0,0 +1,79 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""A collection of FlagEnum classes that syntactically represents the SYCL
+memory enum classes.
+"""
+
+from numba_dpex.experimental.flag_enum import FlagEnum
+
+
+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
+    implementation in dpcpp.
+
+    =====================   ============
+    Order                   Enum value
+    =====================   ============
+    RELAXED                 0
+    ACQUIRE                 1
+    CONSUME_UNSUPPORTED     2
+    RELEASE                 3
+    ACQ_REL                 4
+    SEQ_CST                 5
+    =====================   ============
+    """
+
+    RELAXED = 0
+    ACQUIRE = 1
+    CONSUME_UNSUPPORTED = 2
+    RELEASE = 3
+    ACQ_REL = 4
+    SEQ_CST = 5
+
+
+class MemoryScope(FlagEnum):
+    """
+    An enumeration of SYCL memory scope. For more details please refer to
+    SYCL 2020 specification, section 3.8.3.2
+
+    ===============  ============
+    Memory Scope     Enum value
+    ===============  ============
+    WORK_ITEM        0
+    SUB_GROUP        1
+    WORK_GROUP       2
+    DEVICE           3
+    SYSTEM           4
+    ===============  ============
+    """
+
+    WORK_ITEM = 0
+    SUB_GROUP = 1
+    WORK_GROUP = 2
+    DEVICE = 3
+    SYSTEM = 4
+
+
+class AddressSpace(FlagEnum):
+    """The address space values supported by numba_dpex.
+
+    ==================   ============
+    Address space        Value
+    ==================   ============
+    PRIVATE              0
+    GLOBAL               1
+    CONSTANT             2
+    LOCAL                3
+    GENERIC              4
+    ==================   ============
+    """
+
+    PRIVATE = 0
+    GLOBAL = 1
+    CONSTANT = 2
+    LOCAL = 3
+    GENERIC = 4

numba_dpex/tests/experimental/kernel_iface/__init__.py

@@ -0,0 +1,3 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0

numba_dpex/tests/experimental/kernel_iface/test_memory_enum_compilation.py

@@ -0,0 +1,78 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import dpnp
+
+import numba_dpex.experimental as exp_dpex
+from numba_dpex import Range
+from numba_dpex.experimental.kernel_iface import (
+    AddressSpace,
+    MemoryOrder,
+    MemoryScope,
+)
+
+
+def test_compilation_of_memory_order():
+    """Tests if a MemoryOrder flags can be used inside a kernel function."""
+
+    @exp_dpex.kernel
+    def store_memory_order_flag(a):
+        a[0] = MemoryOrder.RELAXED
+        a[1] = MemoryOrder.CONSUME_UNSUPPORTED
+        a[2] = MemoryOrder.ACQ_REL
+        a[3] = MemoryOrder.ACQUIRE
+        a[4] = MemoryOrder.RELEASE
+        a[5] = MemoryOrder.SEQ_CST
+
+    a = dpnp.ones(10, dtype=dpnp.int64)
+    exp_dpex.call_kernel(store_memory_order_flag, Range(10), a)
+
+    assert a[0] == MemoryOrder.RELAXED
+    assert a[1] == MemoryOrder.CONSUME_UNSUPPORTED
+    assert a[2] == MemoryOrder.ACQ_REL
+    assert a[3] == MemoryOrder.ACQUIRE
+    assert a[4] == MemoryOrder.RELEASE
+    assert a[5] == MemoryOrder.SEQ_CST
+
+
+def test_compilation_of_memory_scope():
+    """Tests if a MemoryScope flags can be used inside a kernel function."""
+
+    @exp_dpex.kernel
+    def store_memory_scope_flag(a):
+        a[0] = MemoryScope.DEVICE
+        a[1] = MemoryScope.SUB_GROUP
+        a[2] = MemoryScope.WORK_GROUP
+        a[3] = MemoryScope.SYSTEM
+        a[4] = MemoryScope.WORK_ITEM
+
+    a = dpnp.ones(10, dtype=dpnp.int64)
+    exp_dpex.call_kernel(store_memory_scope_flag, Range(10), a)
+
+    assert a[0] == MemoryScope.DEVICE
+    assert a[1] == MemoryScope.SUB_GROUP
+    assert a[2] == MemoryScope.WORK_GROUP
+    assert a[3] == MemoryScope.SYSTEM
+    assert a[4] == MemoryScope.WORK_ITEM
+
+
+def test_compilation_of_address_space():
+    """Tests if a AddressSpace flags can be used inside a kernel function."""
+
+    @exp_dpex.kernel
+    def store_address_space_flag(a):
+        a[0] = AddressSpace.CONSTANT
+        a[1] = AddressSpace.GENERIC
+        a[2] = AddressSpace.GLOBAL
+        a[3] = AddressSpace.LOCAL
+        a[4] = AddressSpace.PRIVATE
+
+    a = dpnp.ones(10, dtype=dpnp.int64)
+    exp_dpex.call_kernel(store_address_space_flag, Range(10), a)
+
+    assert a[0] == AddressSpace.CONSTANT
+    assert a[1] == AddressSpace.GENERIC
+    assert a[2] == AddressSpace.GLOBAL
+    assert a[3] == AddressSpace.LOCAL
+    assert a[4] == AddressSpace.PRIVATE