From 837109e0bab57318cf857551823092cac6cdc5c4 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Sun, 10 Dec 2023 22:48:22 -0600 Subject: [PATCH 1/2] Adds a kernel_iface module in experimental. - A new experimental module kernel_iface was added to store Python implementation of the kernel API. The initial module has the memory enum classes that can be used as flags in a numba_dpex.experimental.kernel --- .../experimental/kernel_iface/__init__.py | 13 +++ .../experimental/kernel_iface/memory_enums.py | 79 +++++++++++++++++++ 2 files changed, 92 insertions(+) create mode 100644 numba_dpex/experimental/kernel_iface/__init__.py create mode 100644 numba_dpex/experimental/kernel_iface/memory_enums.py diff --git a/numba_dpex/experimental/kernel_iface/__init__.py b/numba_dpex/experimental/kernel_iface/__init__.py new file mode 100644 index 0000000000..ef1e3855dc --- /dev/null +++ b/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"] diff --git a/numba_dpex/experimental/kernel_iface/memory_enums.py b/numba_dpex/experimental/kernel_iface/memory_enums.py new file mode 100644 index 0000000000..6c741624e0 --- /dev/null +++ b/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 From d175a3b5114d3e0ebfa17f4c0d7cb010d82fad34 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 11 Dec 2023 15:12:16 -0600 Subject: [PATCH 2/2] Unit test --- .../experimental/kernel_iface/__init__.py | 3 + .../test_memory_enum_compilation.py | 78 +++++++++++++++++++ 2 files changed, 81 insertions(+) create mode 100644 numba_dpex/tests/experimental/kernel_iface/__init__.py create mode 100644 numba_dpex/tests/experimental/kernel_iface/test_memory_enum_compilation.py diff --git a/numba_dpex/tests/experimental/kernel_iface/__init__.py b/numba_dpex/tests/experimental/kernel_iface/__init__.py new file mode 100644 index 0000000000..3a217e6325 --- /dev/null +++ b/numba_dpex/tests/experimental/kernel_iface/__init__.py @@ -0,0 +1,3 @@ +# SPDX-FileCopyrightText: 2023 Intel Corporation +# +# SPDX-License-Identifier: Apache-2.0 diff --git a/numba_dpex/tests/experimental/kernel_iface/test_memory_enum_compilation.py b/numba_dpex/tests/experimental/kernel_iface/test_memory_enum_compilation.py new file mode 100644 index 0000000000..dba2c0de32 --- /dev/null +++ b/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