Merge pull request #1103 from IntelPython/feature/remove_atomic_emula…

…tion

Remove atomic emulation

Dockerfile

@@ -199,7 +199,7 @@ RUN \
     --mount=type=bind,target=/opt/toolkit,source=/opt/toolkit,from=toolkit-dist \
     export http_proxy=$http_proxy https_proxy=$https_proxy \
     && apt-get update && apt-get install -y \
-    spirv-tools spirv-headers \
+    spirv-headers \
     rsync \
     && rm -rf /var/lib/apt/lists/* \
     && rsync -a /opt/toolkit/bin/ /usr/local/bin/ \

conda-recipe/run_test.sh

@@ -5,15 +5,8 @@ unset ONEAPI_DEVICE_SELECTOR
 
 for selector in $(python -c "import dpctl; print(\" \".join([dev.backend.name+\":\"+dev.device_type.name for dev in dpctl.get_devices() if dev.device_type.name in [\"cpu\",\"gpu\"]]))")
 do
-    export "ONEAPI_DEVICE_SELECTOR=$selector"
-    unset NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1
-
+    ONEAPI_DEVICE_SELECTOR=$selector \
     pytest -q -ra --disable-warnings --pyargs numba_dpex -vv
-
-    export NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1
-
-    pytest -q -ra --disable-warnings -vv \
-        --pyargs numba_dpex.tests.kernel_tests.test_atomic_op::test_atomic_fp_native
 done
 
 exit 0

docs/backups/user_guides/getting_started.rst

@@ -11,7 +11,6 @@ Numba-dpex depends on following components:
 * dpnp 0.10.1 (`Intel Python DPNP`_)
 * `dpcpp-llvm-spirv`_ (SPIRV generation from LLVM IR)
 * `llvmdev`_ (LLVM IR generation)
-* `spirv-tools`_
 * `packaging`_
 * `scipy`_ (for testing)
 * `pytest`_ (for testing)
@@ -59,7 +58,7 @@ installed in conda environment:
 .. code-block:: bash
 
     export ONEAPI_ROOT=/opt/intel/oneapi
-    conda create -n numba-dpex-env -c ${ONEAPI_ROOT}/conda_channel python=3.7 dpctl dpnp numba spirv-tools dpcpp-llvm-spirv llvmdev pytest
+    conda create -n numba-dpex-env -c ${ONEAPI_ROOT}/conda_channel python=3.7 dpctl dpnp numba dpcpp-llvm-spirv llvmdev pytest
     conda activate numba-dpex-env
 
 Activate DPC++ compiler:
@@ -150,7 +149,6 @@ Refer to :ref:`Docker <docker>` section for more options.
 .. _`Intel Python dpnp`: https://github.com/IntelPython/dpnp
 .. _`dpcpp-llvm-spirv`: https://github.com/IntelPython/dpcpp-llvm-spirv
 .. _`llvmdev`: https://anaconda.org/intel/llvmdev
-.. _`spirv-tools`: https://anaconda.org/intel/spirv-tools
 .. _`packaging`: https://packaging.pypa.io/
 .. _`scipy`: https://anaconda.org/intel/scipy
 .. _`pytest`: https://docs.pytest.org

docs/backups/user_guides/kernel_programming_guide/atomic-operations.rst

@@ -21,21 +21,6 @@ Example usage of atomic operations
     The ``numba_dpex.atomic.add`` function is analogous to The
     ``numba.cuda.atomic.add`` provided by the ``numba.cuda`` backend.
 
-Generating Native FP Atomics
-----------------------------
-Numba-dpex supports generating native floating-point atomics.
-This feature is experimental. Users will need to provide
-the following environment variables to activate it.
-
-    NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1
-
-Example command:
-
-.. code-block:: bash
-
-    NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1 \
-    python program.py
-
 Full examples
 -------------
 

docs/source/getting_started.rst

@@ -19,7 +19,7 @@ to get the latest production releases.
 .. code-block:: bash
 
     conda create -n numba-dpex-env                                             \
-        numba-dpex dpnp dpctl dpcpp-llvm-spirv spirv-tools                     \
+        numba-dpex dpnp dpctl dpcpp-llvm-spirv                     \
         -c intel -c conda-forge
 
 To try out the bleeding edge, the latest packages built from tip of the main
@@ -28,7 +28,7 @@ source trunk can be installed from the ``dppy/label/dev`` conda channel.
 .. code-block:: bash
 
     conda create -n numba-dpex-env                                             \
-        numba-dpex dpnp dpctl dpcpp-llvm-spirv spirv-tools                     \
+        numba-dpex dpnp dpctl dpcpp-llvm-spirv                     \
         -c dppy/label/dev -c intel -c conda-forge
 
 
@@ -70,7 +70,7 @@ first step.
 
     # Create a conda environment that hass needed dependencies installed
     conda create -n numba-dpex-env                                             \
-        dpctl dpnp numba spirv-tools dpcpp-llvm-spirv llvmdev pytest           \
+        dpctl dpnp numba dpcpp-llvm-spirv llvmdev pytest           \
         -c intel -c conda-forge
     # Activate the environment
     conda activate numba-dpex-env

docs/source/user_guide/kernel_programming/atomic-operations.rst

@@ -21,23 +21,6 @@ Example usage of atomic operations
     The ``numba_dpex.atomic.add`` function is analogous to The
     ``numba.cuda.atomic.add`` provided by the ``numba.cuda`` backend.
 
-Generating Native FP Atomics
-----------------------------
-Numba-dpex supports generating native floating-point atomics.
-This feature is experimental. Users will need to provide
-the following environment variables to activate it.
-
-    NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1
-    NUMBA_DPEX_LLVM_SPIRV_ROOT=/path/to/dpcpp/provided/llvm_spirv
-
-Example command:
-
-.. code-block:: bash
-
-    NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1 \
-    NUMBA_DPEX_LLVM_SPIRV_ROOT=/path/to/dpcpp/provided/llvm_spirv \
-    python program.py
-
 Full examples
 -------------
 

environment/coverage.yml

@@ -13,7 +13,6 @@ dependencies:
   - numba=0.57
   - dpctl
   - dpnp
-  - spirv-tools
   - dpcpp-llvm-spirv
   - opencl_rt
   - coverage

environment/docs.yml

@@ -13,7 +13,6 @@ dependencies:
   - numba=0.57
   - dpctl
   - dpnp
-  - spirv-tools
   - dpcpp-llvm-spirv
   - opencl_rt
   - pip

numba_dpex/config.py

@@ -53,9 +53,6 @@ def __getattr__(name):
 # Dump offload diagnostics
 OFFLOAD_DIAGNOSTICS = _readenv("NUMBA_DPEX_OFFLOAD_DIAGNOSTICS", int, 0)
 
-# Activate Native floating point atomcis support for supported devices.
-# Requires llvm-spirv supporting the FP atomics extension
-NATIVE_FP_ATOMICS = _readenv("NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE", int, 0)
 # Emit debug info
 DEBUG = _readenv("NUMBA_DPEX_DEBUG", int, config.DEBUG)
 DEBUGINFO_DEFAULT = _readenv(

numba_dpex/examples/kernel/atomic_op.py

@@ -8,20 +8,28 @@
 
 
 @ndpx.kernel
-def atomic_reduction(a):
+def atomic_reduction(a, res):
+    """Summarize all the items in a and writes it into res using atomic.add.
+
+    :param a: array of values to get sum
+    :param res: result where to add all the items from a array. It must be preset to 0.
+    """
     idx = ndpx.get_global_id(0)
-    ndpx.atomic.add(a, 0, a[idx])
+    ndpx.atomic.add(res, 0, a[idx])
 
 
 def main():
     N = 10
-    a = np.arange(N)
+
+    # We are storing sum to the first element
+    a = np.arange(0, N)
+    res = np.zeros(1, dtype=a.dtype)
 
     print("Using device ...")
     print(a.device)
 
-    atomic_reduction[ndpx.Range(N)](a)
-    print("Reduction sum =", a[0])
+    atomic_reduction[ndpx.Range(N)](a, res)
+    print("Reduction sum =", res[0])
 
     print("Done...")
 

numba_dpex/ocl/__init__.py

@@ -1,5 +1,3 @@
 # SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
 #
 # SPDX-License-Identifier: Apache-2.0
-
-from .atomics import atomic_support_present

numba_dpex/ocl/atomics/__init__.py

@@ -1,32 +1,3 @@
 # SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
 #
 # SPDX-License-Identifier: Apache-2.0
-
-import os
-import os.path
-
-
-def atomic_support_present():
-    if os.path.isfile(
-        os.path.join(os.path.dirname(__file__), "atomic_ops.spir")
-    ):
-        return True
-    else:
-        return False
-
-
-def get_atomic_spirv_path():
-    if atomic_support_present():
-        return os.path.join(os.path.dirname(__file__), "atomic_ops.spir")
-    else:
-        return None
-
-
-def read_atomic_spirv_file():
-    path = get_atomic_spirv_path()
-    if path:
-        with open(path, "rb") as fin:
-            spirv = fin.read()
-        return spirv
-    else:
-        return None