Repository: IntelPython/numba-dpex (language: python)

Issue:
Updated Changelog with all changes going into 0.23
- [X] Have you provided a meaningful PR description?

Updated CHANGELOG with latest changes since 4da339d




Discussion / hints:

Fix patch:
diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml
index f6dc0dcc72..f5496bae39 100644
--- a/.github/workflows/conda-package.yml
+++ b/.github/workflows/conda-package.yml
@@ -30,7 +30,7 @@ jobs:
       fail-fast: false
       matrix:
         python: ['3.9', '3.10', '3.11']
-        os: [ubuntu-latest, windows-latest]
+        os: [ubuntu-latest, windows-2019]
 
     runs-on: ${{ matrix.os }}
 
diff --git a/.github/workflows/coverage.yml b/.github/workflows/coverage.yml
index c949f6dfc6..c9b2bc9154 100644
--- a/.github/workflows/coverage.yml
+++ b/.github/workflows/coverage.yml
@@ -54,7 +54,7 @@ jobs:
       # Ignoring test due to opencl driver optimization bug
       - name: Run tests with coverage
         run: |
-          pytest  -q --cov --cov-report xml --pyargs numba_dpex \
+          pytest  -q --cov=./ --cov-report xml --pyargs numba_dpex \
             -k 'not test_1d_strided_dpnp_array_in_kernel[2]'
 
       - name: Coveralls
diff --git a/.github/workflows/cpp_style_checks.yml b/.github/workflows/cpp_style_checks.yml
index 71c431b760..92ead9fb16 100644
--- a/.github/workflows/cpp_style_checks.yml
+++ b/.github/workflows/cpp_style_checks.yml
@@ -18,7 +18,7 @@ jobs:
     steps:
     - uses: actions/checkout@v4
     - name: Run clang-format style check for C/C++ programs.
-      uses: jidicula/clang-format-action@v4.11.0
+      uses: jidicula/clang-format-action@v4.13.0
       with:
         clang-format-version: '14'
         check-path: 'numba_dpex/dpctl_iface'
diff --git a/.github/workflows/gh-pages.yml b/.github/workflows/gh-pages.yml
index c4eff984b2..bed674efea 100644
--- a/.github/workflows/gh-pages.yml
+++ b/.github/workflows/gh-pages.yml
@@ -51,7 +51,7 @@ jobs:
         run: make html
 
       - name: GitHub Pages [main]
-        uses: peaceiris/actions-gh-pages@v3.9.3
+        uses: peaceiris/actions-gh-pages@v4.0.0
         if: ${{ github.ref == 'refs/heads/main' }}
         with:
           github_token: ${{ secrets.GITHUB_TOKEN }}
@@ -64,7 +64,7 @@ jobs:
           user_email: 'github-actions[bot]@users.noreply.github.com'
 
       - name: GitHub Pages [PR]
-        uses: peaceiris/actions-gh-pages@v3.9.3
+        uses: peaceiris/actions-gh-pages@v4.0.0
         if: ${{ github.event.pull_request && github.event.action != 'closed' }}
         with:
           github_token: ${{ secrets.GITHUB_TOKEN }}
@@ -88,7 +88,7 @@ jobs:
 
       - name: Publish release
         if: startsWith(github.ref, 'refs/heads/release')
-        uses: peaceiris/actions-gh-pages@v3.9.3
+        uses: peaceiris/actions-gh-pages@v4.0.0
         with:
           github_token: ${{ secrets.GITHUB_TOKEN }}
           destination_dir : next_release
@@ -104,7 +104,7 @@ jobs:
 
       - name: Publish tag
         if: startsWith(github.ref, 'refs/tags/')
-        uses: peaceiris/actions-gh-pages@v3.9.3
+        uses: peaceiris/actions-gh-pages@v4.0.0
         with:
           github_token: ${{ secrets.GITHUB_TOKEN }}
           destination_dir : ${{ steps.capture_tag.outputs.tag_number }}
diff --git a/.github/workflows/openssf-scorecard.yml b/.github/workflows/openssf-scorecard.yml
index b5ea6ac4c6..6f953e3b8d 100644
--- a/.github/workflows/openssf-scorecard.yml
+++ b/.github/workflows/openssf-scorecard.yml
@@ -40,7 +40,7 @@ jobs:
           persist-credentials: false
 
       - name: "Run analysis"
-        uses: ossf/scorecard-action@0864cf19026789058feabb7e87baa5f140aac736 # v2.3.1
+        uses: ossf/scorecard-action@dc50aa9510b46c811795eb24b2f1ba02a914e534 # v2.3.3
         with:
           results_file: results.sarif
           results_format: sarif
@@ -71,6 +71,6 @@ jobs:
       # Upload the results to GitHub's code scanning dashboard.
       - name: "Upload to code-scanning"
         if: ${{ github.event_name == 'push' }}
-        uses: github/codeql-action/upload-sarif@1b1aada464948af03b950897e5eb522f92603cc2 # v3.24.9
+        uses: github/codeql-action/upload-sarif@9fdb3e49720b44c48891d036bb502feb25684276 # v3.25.6
         with:
           sarif_file: results.sarif
diff --git a/.gitignore b/.gitignore
index bd383c0a5b..4735becd1e 100644
--- a/.gitignore
+++ b/.gitignore
@@ -26,6 +26,9 @@ _skbuild
 
 docs/source/developer/autogen*
 
+# Ignore versioneer generated files
+numba_dpex/_version.py
+
 # Ignore generated cpp files
 numba_dpex/dpnp_iface/*.cpp
 numba_dpex/dpnp_iface/*.h
diff --git a/CHANGELOG.md b/CHANGELOG.md
index ebab23aa0c..56f5f6bee8 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -4,13 +4,21 @@ All notable changes to this project will be documented in this file.
 The format is based on [Keep a Changelog](https://keepachangelog.com/en/1.0.0/),
 and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html).
 
-## [0.23.0] - 2024-04-XX
+## [0.23.0] - 2024-05-28
 
 ### Fixed
 * Array alignment problem for stack arrays allocated for kernel arguments. (#1357)
 * Issue #892, #906 caused by incorrect code generation for indexing (#1377)
-* Fix `KernelHasReturnValueError` inside `KernelDispatcher`. (#1394)
+* Generation of `KernelHasReturnValueError` error inside `KernelDispatcher`. (#1394)
 * Issue #1390: broken support for slicing into `dpctl.tensor.usm_ndarray` in kernels (#1425)
+* Support for Wheels package on Windows (#1430)
+* Incorrect mangled name for kernel function arguments (#1443)
+* Remove artifacts from conda/wheel packages residing in root level (#1450)
+* GDB tests to work properly on Intel Max GPU (#1451)
+* Improper wheels installation on unsupported platforms (#1452)
+* Ref-counting of Python object temporaries in unboxing code (#1454)
+* Segfault caused by using `malloc` to allocate `NRT_MemInfo`. Replaced with Numba's NRT `alloc` (#1458)
+* Incorrect package name in README.md (#1463)
 
 ### Added
 * A new overloaded `dimensions` attribute for all index-space id classes (#1359)
@@ -24,18 +32,24 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 * A `sycl::local_accessor`-like API (`kernel_api.LocalAccessor`) for numba-dpex kernel (#1331)
 * Specialization support for `device_func` decorator (#1398)
 * Support for all `kernel_api` functions inside the `numba_dpex.kernel` decorator. (#1400)
+* Support for dpnp 0.15 (#1434, #1464)
+* Improvements to pyproject.toml configs to build numba-dpex from source. (#1449)
+* Load the `SPV_INTEL_variable_length_array` SPIR-V extension to supporting arrays in private address-space on Intel Max GPU. (#1451)
 
 ### Changed
 * Default inline threshold value set to `2` from `None`. (#1385)
 * Port parfor kernel templates to `kernel_api` (#1416), (#1424)
+* Use `SPIRVKernelDispatcher` for parfor kernel dispatch (#1435, #1448)
+* All examples use the latest dpctl API (#1431)
 * Minimum required dpctl version is now 0.16.1
-* Minimum required numba version is now 0.59.0
+* Minimum required numba version is now 0.59.0 (#1462)
 
 ### Removed
 * OpenCL-like kernel API functions (#1420)
 * `func` decorator (replaced by `device_func`) (#1400)
 * `numba_dpex.experimental.kernel` and `numba_dpex.experimental.device_func` (#1400)
 
+
 ## [0.22.0] - 2024-02-19
 
 ### Fixed
diff --git a/MANIFEST.in b/MANIFEST.in
deleted file mode 100644
index 8c457d17ad..0000000000
--- a/MANIFEST.in
+++ /dev/null
@@ -1,6 +0,0 @@
-include MANIFEST.in
-include README.md setup.py LICENSE
-
-include numba_dpex/_version.py
-
-recursive-include numba_dpex/examples *
diff --git a/README.md b/README.md
index dcd5f51638..84a1afde6d 100644
--- a/README.md
+++ b/README.md
@@ -12,7 +12,7 @@
 
 
 
-Data-parallel Extension for Numba* (numba-dpex) is an open-source standalone
+Data Parallel Extension for Numba* (numba-dpex) is an open-source standalone
 extension for the [Numba](http://numba.pydata.org) Python JIT compiler.
 Numba-dpex provides a [SYCL*](https://sycl.tech/)-like API for kernel
 programming Python. SYCL* is an open standard developed by the [Unified
diff --git a/conda-recipe/bld.bat b/conda-recipe/bld.bat
index f2ca6a3723..277662827a 100644
--- a/conda-recipe/bld.bat
+++ b/conda-recipe/bld.bat
@@ -35,14 +35,18 @@ for /f %%f in ('dir /b /S .\dist') do (
 
 :: wheel file was renamed
 for /f %%f in ('dir /b /S .\dist') do (
-    %PYTHON% -m pip install %%f
+    %PYTHON% -m pip install %%f ^
+      --no-build-isolation ^
+      --no-deps ^
+      --only-binary :all: ^
+      --no-index ^
+      --prefix %PREFIX% ^
+      -vv
     if %ERRORLEVEL% neq 0 exit 1
 )
 
 :: Copy wheel package
 if NOT "%WHEELS_OUTPUT_FOLDER%"=="" (
-    for /f %%f in ('dir /b /S .\dist') do (
-        copy %%f %WHEELS_OUTPUT_FOLDER%
-        if %ERRORLEVEL% neq 0 exit 1
-    )
+    copy dist\numba_dpex*.whl %WHEELS_OUTPUT_FOLDER%
+    if errorlevel 1 exit 1
 )
diff --git a/conda-recipe/build.sh b/conda-recipe/build.sh
index 68f12528f9..a6fb6607f7 100644
--- a/conda-recipe/build.sh
+++ b/conda-recipe/build.sh
@@ -9,6 +9,9 @@ echo "--gcc-toolchain=${BUILD_PREFIX} --sysroot=${BUILD_PREFIX}/${HOST}/sysroot
 ICPXCFG="$(pwd)/icpx_for_conda.cfg"
 ICXCFG="$(pwd)/icpx_for_conda.cfg"
 
+read -r GLIBC_MAJOR GLIBC_MINOR <<<"$(conda list '^sysroot_linux-64$' \
+    | tail -n 1 | awk '{print $2}' | grep -oP '\d+' | head -n 2 | tr '\n' ' ')"
+
 export ICXCFG
 export ICPXCFG
 
@@ -26,8 +29,15 @@ export PATH=$CONDA_PREFIX/bin-llvm:$PATH
 # -wnx flags mean: --wheel --no-isolation --skip-dependency-check
 ${PYTHON} -m build -w -n -x
 ${PYTHON} -m wheel tags --remove --build "$GIT_DESCRIBE_NUMBER" \
-    --platform-tag manylinux2014_x86_64 dist/numba_dpex*.whl
-${PYTHON} -m pip install dist/numba_dpex*.whl
+    --platform-tag "manylinux_${GLIBC_MAJOR}_${GLIBC_MINOR}_x86_64" \
+    dist/numba_dpex*.whl
+${PYTHON} -m pip install dist/numba_dpex*.whl \
+    --no-build-isolation \
+    --no-deps \
+    --only-binary :all: \
+    --no-index \
+    --prefix "${PREFIX}" \
+    -vv
 
 # Copy wheel package
 if [[ -v WHEELS_OUTPUT_FOLDER ]]; then
diff --git a/conda-recipe/meta.yaml b/conda-recipe/meta.yaml
index e70aa7a522..6e27f82491 100644
--- a/conda-recipe/meta.yaml
+++ b/conda-recipe/meta.yaml
@@ -3,6 +3,9 @@
 {% set excluded_compiler_version2 = "2024.0.2" %}
 {% set excluded_compiler_version3 = "2024.0.3" %}
 
+{% set pyproject = load_file_data('pyproject.toml') %}
+{% set py_build_deps = pyproject.get('build-system', {}).get('requires', []) %}
+
 package:
     name: numba-dpex
     version: {{ GIT_DESCRIBE_TAG }}
@@ -21,34 +24,38 @@ requirements:
         - {{ compiler('cxx') }}
         - {{ compiler('dpcpp') }} >={{ required_compiler_version }},!={{ excluded_compiler_version1 }},!={{ excluded_compiler_version2 }},!={{ excluded_compiler_version3 }}  # [win]
         - {{ compiler('dpcpp') }} >={{ required_compiler_version }},!={{ excluded_compiler_version1 }},!={{ excluded_compiler_version2 }}  # [linux]
-        # specific version of sysroot required by dpcpp, but 2024.0.0 package
-        # does not have it in meta data
-        - sysroot_linux-64 >=2.28  # [linux]
+        # Minimal supported version of sysroot (which is version of glibc) to
+        # have compatibility with wider range of linux distributions.
+        # 2.28 is the minimal supported version by dpcpp
+        - sysroot_linux-64 =2.28  # [linux]
     host:
+        - python
+        - pip >=24.0
         - dpcpp-cpp-rt >={{ required_compiler_version }},!={{ excluded_compiler_version1 }},!={{ excluded_compiler_version2 }},!={{ excluded_compiler_version3 }}  # [win]
         - dpcpp-cpp-rt >={{ required_compiler_version }},!={{ excluded_compiler_version1 }},!={{ excluded_compiler_version2 }}  # [linux]
-        - python
+        # ensure we are using latest version of setuptools, since we don't need
+        # editable environments for release.
         - setuptools >=69
-        - scikit-build >=0.17
-        - ninja >=1.11.1 # [not win]
-        - cmake >=3.29
-        - numba >=0.59
-        - llvmlite >=0.42.0
-        - dpctl >=0.16.1
-        - dpnp >=0.14
-        - numpy >=1.24
-        # TODO: there is no 2024 release for python 3.11
-        # - dpcpp-llvm-spirv >={{ required_compiler_version }}
-        - dpcpp-llvm-spirv >=2023.0
-        - wheel >=0.43
-        - pip >=24.0
-        - python-build >=1.1
-        - versioneer==0.29
+        {% for dep in py_build_deps %}
+        {% if dep.startswith('ninja') %}
+        - {{ dep.split(';')[0] }} # [not win]
+        {% elif dep.startswith('cmake') %}
+        - {{ dep }}
+        {% elif dep.startswith('build>=') %}
+        - {{ 'python-' ~ dep }}
+        {% else %}
+        - {{ dep|replace('_','-') }}
+        {% endif %}
+        {% endfor %}
         # versioneer dependency
         - tomli # [py<311]
+        # While we don't need it for build, but install it here, so we can
+        # pin_compatible at run section.
+        - dpcpp-llvm-spirv >={{ required_compiler_version }}
     run:
         - {{ pin_compatible('dpcpp-cpp-rt', min_pin='x.x', max_pin='x') }}
         - {{ pin_compatible('intel-cmplr-lib-rt', min_pin='x.x', max_pin='x') }}
+        # TODO: pick up min version from dep
         - {{ pin_compatible('dpcpp-llvm-spirv', min_pin='x.x', max_pin='x') }}
         - {{ pin_compatible('dpnp', min_pin='x.x.x', max_pin='x.x') }}
         - {{ pin_compatible('dpctl', min_pin='x.x.x', max_pin='x.x') }}
diff --git a/numba_dpex/__init__.py b/numba_dpex/__init__.py
index 4a24e28502..52228ac303 100644
--- a/numba_dpex/__init__.py
+++ b/numba_dpex/__init__.py
@@ -17,7 +17,7 @@
 from numba import __version__ as numba_version
 
 from .kernel_api_impl.spirv import target as spirv_kernel_target
-from .numba_patches import patch_is_ufunc
+from .numba_patches import patch_ufuncs
 from .register_kernel_api_overloads import init_kernel_api_spirv_overloads
 
 
@@ -70,7 +70,7 @@ def parse_sem_version(version_string: str) -> Tuple[int, int, int]:
 dpctl_sem_version = parse_sem_version(dpctl.__version__)
 
 # Monkey patches
-patch_is_ufunc.patch()
+patch_ufuncs.patch()
 
 from numba import prange  # noqa E402
 
diff --git a/numba_dpex/_version.py b/numba_dpex/_version.py
deleted file mode 100644
index f076bfbd66..0000000000
--- a/numba_dpex/_version.py
+++ /dev/null
@@ -1,683 +0,0 @@
-
-# This file helps to compute a version number in source trees obtained from
-# git-archive tarball (such as those provided by githubs download-from-tag
-# feature). Distribution tarballs (built by setup.py sdist) and build
-# directories (produced by setup.py build) will contain a much shorter file
-# that just contains the computed version number.
-
-# This file is released into the public domain.
-# Generated by versioneer-0.29
-# https://github.com/python-versioneer/python-versioneer
-
-"""Git implementation of _version.py."""
-
-import errno
-import os
-import re
-import subprocess
-import sys
-from typing import Any, Callable, Dict, List, Optional, Tuple
-import functools
-
-
-def get_keywords() -> Dict[str, str]:
-    """Get the keywords needed to look up the version information."""
-    # these strings will be replaced by git during git-archive.
-    # setup.py/versioneer.py will grep for the variable names, so they must
-    # each be defined on a line of their own. _version.py will just call
-    # get_keywords().
-    git_refnames = "$Format:%d$"
-    git_full = "$Format:%H$"
-    git_date = "$Format:%ci$"
-    keywords = {"refnames": git_refnames, "full": git_full, "date": git_date}
-    return keywords
-
-
-class VersioneerConfig:
-    """Container for Versioneer configuration parameters."""
-
-    VCS: str
-    style: str
-    tag_prefix: str
-    parentdir_prefix: str
-    versionfile_source: str
-    verbose: bool
-
-
-def get_config() -> VersioneerConfig:
-    """Create, populate and return the VersioneerConfig() object."""
-    # these strings are filled in when 'setup.py versioneer' creates
-    # _version.py
-    cfg = VersioneerConfig()
-    cfg.VCS = "git"
-    cfg.style = "pep440"
-    cfg.tag_prefix = ""
-    cfg.parentdir_prefix = ""
-    cfg.versionfile_source = "numba_dpex/_version.py"
-    cfg.verbose = False
-    return cfg
-
-
-class NotThisMethod(Exception):
-    """Exception raised if a method is not valid for the current scenario."""
-
-
-LONG_VERSION_PY: Dict[str, str] = {}
-HANDLERS: Dict[str, Dict[str, Callable]] = {}
-
-
-def register_vcs_handler(vcs: str, method: str) -> Callable:  # decorator
-    """Create decorator to mark a method as the handler of a VCS."""
-    def decorate(f: Callable) -> Callable:
-        """Store f in HANDLERS[vcs][method]."""
-        if vcs not in HANDLERS:
-            HANDLERS[vcs] = {}
-        HANDLERS[vcs][method] = f
-        return f
-    return decorate
-
-
-def run_command(
-    commands: List[str],
-    args: List[str],
-    cwd: Optional[str] = None,
-    verbose: bool = False,
-    hide_stderr: bool = False,
-    env: Optional[Dict[str, str]] = None,
-) -> Tuple[Optional[str], Optional[int]]:
-    """Call the given command(s)."""
-    assert isinstance(commands, list)
-    process = None
-
-    popen_kwargs: Dict[str, Any] = {}
-    if sys.platform == "win32":
-        # This hides the console window if pythonw.exe is used
-        startupinfo = subprocess.STARTUPINFO()
-        startupinfo.dwFlags |= subprocess.STARTF_USESHOWWINDOW
-        popen_kwargs["startupinfo"] = startupinfo
-
-    for command in commands:
-        try:
-            dispcmd = str([command] + args)
-            # remember shell=False, so use git.cmd on windows, not just git
-            process = subprocess.Popen([command] + args, cwd=cwd, env=env,
-                                       stdout=subprocess.PIPE,
-                                       stderr=(subprocess.PIPE if hide_stderr
-                                               else None), **popen_kwargs)
-            break
-        except OSError as e:
-            if e.errno == errno.ENOENT:
-                continue
-            if verbose:
-                print("unable to run %s" % dispcmd)
-                print(e)
-            return None, None
-    else:
-        if verbose:
-            print("unable to find command, tried %s" % (commands,))
-        return None, None
-    stdout = process.communicate()[0].strip().decode()
-    if process.returncode != 0:
-        if verbose:
-            print("unable to run %s (error)" % dispcmd)
-            print("stdout was %s" % stdout)
-        return None, process.returncode
-    return stdout, process.returncode
-
-
-def versions_from_parentdir(
-    parentdir_prefix: str,
-    root: str,
-    verbose: bool,
-) -> Dict[str, Any]:
-    """Try to determine the version from the parent directory name.
-
-    Source tarballs conventionally unpack into a directory that includes both
-    the project name and a version string. We will also support searching up
-    two directory levels for an appropriately named parent directory
-    """
-    rootdirs = []
-
-    for _ in range(3):
-        dirname = os.path.basename(root)
-        if dirname.startswith(parentdir_prefix):
-            return {"version": dirname[len(parentdir_prefix):],
-                    "full-revisionid": None,
-                    "dirty": False, "error": None, "date": None}
-        rootdirs.append(root)
-        root = os.path.dirname(root)  # up a level
-
-    if verbose:
-        print("Tried directories %s but none started with prefix %s" %
-              (str(rootdirs), parentdir_prefix))
-    raise NotThisMethod("rootdir doesn't start with parentdir_prefix")
-
-
-@register_vcs_handler("git", "get_keywords")
-def git_get_keywords(versionfile_abs: str) -> Dict[str, str]:
-    """Extract version information from the given file."""
-    # the code embedded in _version.py can just fetch the value of these
-    # keywords. When used from setup.py, we don't want to import _version.py,
-    # so we do it with a regexp instead. This function is not used from
-    # _version.py.
-    keywords: Dict[str, str] = {}
-    try:
-        with open(versionfile_abs, "r") as fobj:
-            for line in fobj:
-                if line.strip().startswith("git_refnames ="):
-                    mo = re.search(r'=\s*"(.*)"', line)
-                    if mo:
-                        keywords["refnames"] = mo.group(1)
-                if line.strip().startswith("git_full ="):
-                    mo = re.search(r'=\s*"(.*)"', line)
-                    if mo:
-                        keywords["full"] = mo.group(1)
-                if line.strip().startswith("git_date ="):
-                    mo = re.search(r'=\s*"(.*)"', line)
-                    if mo:
-                        keywords["date"] = mo.group(1)
-    except OSError:
-        pass
-    return keywords
-
-
-@register_vcs_handler("git", "keywords")
-def git_versions_from_keywords(
-    keywords: Dict[str, str],
-    tag_prefix: str,
-    verbose: bool,
-) -> Dict[str, Any]:
-    """Get version information from git keywords."""
-    if "refnames" not in keywords:
-        raise NotThisMethod("Short version file found")
-    date = keywords.get("date")
-    if date is not None:
-        # Use only the last line.  Previous lines may contain GPG signature
-        # information.
-        date = date.splitlines()[-1]
-
-        # git-2.2.0 added "%cI", which expands to an ISO-8601 -compliant
-        # datestamp. However we prefer "%ci" (which expands to an "ISO-8601
-        # -like" string, which we must then edit to make compliant), because
-        # it's been around since git-1.5.3, and it's too difficult to
-        # discover which version we're using, or to work around using an
-        # older one.
-        date = date.strip().replace(" ", "T", 1).replace(" ", "", 1)
-    refnames = keywords["refnames"].strip()
-    if refnames.startswith("$Format"):
-        if verbose:
-            print("keywords are unexpanded, not using")
-        raise NotThisMethod("unexpanded keywords, not a git-archive tarball")
-    refs = {r.strip() for r in refnames.strip("()").split(",")}
-    # starting in git-1.8.3, tags are listed as "tag: foo-1.0" instead of
-    # just "foo-1.0". If we see a "tag: " prefix, prefer those.
-    TAG = "tag: "
-    tags = {r[len(TAG):] for r in refs if r.startswith(TAG)}
-    if not tags:
-        # Either we're using git < 1.8.3, or there really are no tags. We use
-        # a heuristic: assume all version tags have a digit. The old git %d
-        # expansion behaves like git log --decorate=short and strips out the
-        # refs/heads/ and refs/tags/ prefixes that would let us distinguish
-        # between branches and tags. By ignoring refnames without digits, we
-        # filter out many common branch names like "release" and
-        # "stabilization", as well as "HEAD" and "master".
-        tags = {r for r in refs if re.search(r'\d', r)}
-        if verbose:
-            print("discarding '%s', no digits" % ",".join(refs - tags))
-    if verbose:
-        print("likely tags: %s" % ",".join(sorted(tags)))
-    for ref in sorted(tags):
-        # sorting will prefer e.g. "2.0" over "2.0rc1"
-        if ref.startswith(tag_prefix):
-            r = ref[len(tag_prefix):]
-            # Filter out refs that exactly match prefix or that don't start
-            # with a number once the prefix is stripped (mostly a concern
-            # when prefix is '')
-            if not re.match(r'\d', r):
-                continue
-            if verbose:
-                print("picking %s" % r)
-            return {"version": r,
-                    "full-revisionid": keywords["full"].strip(),
-                    "dirty": False, "error": None,
-                    "date": date}
-    # no suitable tags, so version is "0+unknown", but full hex is still there
-    if verbose:
-        print("no suitable tags, using unknown + full revision id")
-    return {"version": "0+unknown",
-            "full-revisionid": keywords["full"].strip(),
-            "dirty": False, "error": "no suitable tags", "date": None}
-
-
-@register_vcs_handler("git", "pieces_from_vcs")
-def git_pieces_from_vcs(
-    tag_prefix: str,
-    root: str,
-    verbose: bool,
-    runner: Callable = run_command
-) -> Dict[str, Any]:
-    """Get version from 'git describe' in the root of the source tree.
-
-    This only gets called if the git-archive 'subst' keywords were *not*
-    expanded, and _version.py hasn't already been rewritten with a short
-    version string, meaning we're inside a checked out source tree.
-    """
-    GITS = ["git"]
-    if sys.platform == "win32":
-        GITS = ["git.cmd", "git.exe"]
-
-    # GIT_DIR can interfere with correct operation of Versioneer.
-    # It may be intended to be passed to the Versioneer-versioned project,
-    # but that should not change where we get our version from.
-    env = os.environ.copy()
-    env.pop("GIT_DIR", None)
-    runner = functools.partial(runner, env=env)
-
-    _, rc = runner(GITS, ["rev-parse", "--git-dir"], cwd=root,
-                   hide_stderr=not verbose)
-    if rc != 0:
-        if verbose:
-            print("Directory %s not under git control" % root)
-        raise NotThisMethod("'git rev-parse --git-dir' returned error")
-
-    # if there is a tag matching tag_prefix, this yields TAG-NUM-gHEX[-dirty]
-    # if there isn't one, this yields HEX[-dirty] (no NUM)
-    describe_out, rc = runner(GITS, [
-        "describe", "--tags", "--dirty", "--always", "--long",
-        "--match", f"{tag_prefix}[[:digit:]]*"
-    ], cwd=root)
-    # --long was added in git-1.5.5
-    if describe_out is None:
-        raise NotThisMethod("'git describe' failed")
-    describe_out = describe_out.strip()
-    full_out, rc = runner(GITS, ["rev-parse", "HEAD"], cwd=root)
-    if full_out is None:
-        raise NotThisMethod("'git rev-parse' failed")
-    full_out = full_out.strip()
-
-    pieces: Dict[str, Any] = {}
-    pieces["long"] = full_out
-    pieces["short"] = full_out[:7]  # maybe improved later
-    pieces["error"] = None
-
-    branch_name, rc = runner(GITS, ["rev-parse", "--abbrev-ref", "HEAD"],
-                             cwd=root)
-    # --abbrev-ref was added in git-1.6.3
-    if rc != 0 or branch_name is None:
-        raise NotThisMethod("'git rev-parse --abbrev-ref' returned error")
-    branch_name = branch_name.strip()
-
-    if branch_name == "HEAD":
-        # If we aren't exactly on a branch, pick a branch which represents
-        # the current commit. If all else fails, we are on a branchless
-        # commit.
-        branches, rc = runner(GITS, ["branch", "--contains"], cwd=root)
-        # --contains was added in git-1.5.4
-        if rc != 0 or branches is None:
-            raise NotThisMethod("'git branch --contains' returned error")
-        branches = branches.split("\n")
-
-        # Remove the first line if we're running detached
-        if "(" in branches[0]:
-            branches.pop(0)
-
-        # Strip off the leading "* " from the list of branches.
-        branches = [branch[2:] for branch in branches]
-        if "master" in branches:
-            branch_name = "master"
-        elif not branches:
-            branch_name = None
-        else:
-            # Pick the first branch that is returned. Good or bad.
-            branch_name = branches[0]
-
-    pieces["branch"] = branch_name
-
-    # parse describe_out. It will be like TAG-NUM-gHEX[-dirty] or HEX[-dirty]
-    # TAG might have hyphens.
-    git_describe = describe_out
-
-    # look for -dirty suffix
-    dirty = git_describe.endswith("-dirty")
-    pieces["dirty"] = dirty
-    if dirty:
-        git_describe = git_describe[:git_describe.rindex("-dirty")]
-
-    # now we have TAG-NUM-gHEX or HEX
-
-    if "-" in git_describe:
-        # TAG-NUM-gHEX
-        mo = re.search(r'^(.+)-(\d+)-g([0-9a-f]+)$', git_describe)
-        if not mo:
-            # unparsable. Maybe git-describe is misbehaving?
-            pieces["error"] = ("unable to parse git-describe output: '%s'"
-                               % describe_out)
-            return pieces
-
-        # tag
-        full_tag = mo.group(1)
-        if not full_tag.startswith(tag_prefix):
-            if verbose:
-                fmt = "tag '%s' doesn't start with prefix '%s'"
-                print(fmt % (full_tag, tag_prefix))
-            pieces["error"] = ("tag '%s' doesn't start with prefix '%s'"
-                               % (full_tag, tag_prefix))
-            return pieces
-        pieces["closest-tag"] = full_tag[len(tag_prefix):]
-
-        # distance: number of commits since tag
-        pieces["distance"] = int(mo.group(2))
-
-        # commit: short hex revision ID
-        pieces["short"] = mo.group(3)
-
-    else:
-        # HEX: no tags
-        pieces["closest-tag"] = None
-        out, rc = runner(GITS, ["rev-list", "HEAD", "--left-right"], cwd=root)
-        pieces["distance"] = len(out.split())  # total number of commits
-
-    # commit date: see ISO-8601 comment in git_versions_from_keywords()
-    date = runner(GITS, ["show", "-s", "--format=%ci", "HEAD"], cwd=root)[0].strip()
-    # Use only the last line.  Previous lines may contain GPG signature
-    # information.
-    date = date.splitlines()[-1]
-    pieces["date"] = date.strip().replace(" ", "T", 1).replace(" ", "", 1)
-
-    return pieces
-
-
-def plus_or_dot(pieces: Dict[str, Any]) -> str:
-    """Return a + if we don't already have one, else return a ."""
-    if "+" in pieces.get("closest-tag", ""):
-        return "."
-    return "+"
-
-
-def render_pep440(pieces: Dict[str, Any]) -> str:
-    """Build up version string, with post-release "local version identifier".
-
-    Our goal: TAG[+DISTANCE.gHEX[.dirty]] . Note that if you
-    get a tagged build and then dirty it, you'll get TAG+0.gHEX.dirty
-
-    Exceptions:
-    1: no tags. git_describe was just HEX. 0+untagged.DISTANCE.gHEX[.dirty]
-    """
-    if pieces["closest-tag"]:
-        rendered = pieces["closest-tag"]
-        if pieces["distance"] or pieces["dirty"]:
-            rendered += plus_or_dot(pieces)
-            rendered += "%d.g%s" % (pieces["distance"], pieces["short"])
-            if pieces["dirty"]:
-                rendered += ".dirty"
-    else:
-        # exception #1
-        rendered = "0+untagged.%d.g%s" % (pieces["distance"],
-                                          pieces["short"])
-        if pieces["dirty"]:
-            rendered += ".dirty"
-    return rendered
-
-
-def render_pep440_branch(pieces: Dict[str, Any]) -> str:
-    """TAG[[.dev0]+DISTANCE.gHEX[.dirty]] .
-
-    The ".dev0" means not master branch. Note that .dev0 sorts backwards
-    (a feature branch will appear "older" than the master branch).
-
-    Exceptions:
-    1: no tags. 0[.dev0]+untagged.DISTANCE.gHEX[.dirty]
-    """
-    if pieces["closest-tag"]:
-        rendered = pieces["closest-tag"]
-        if pieces["distance"] or pieces["dirty"]:
-            if pieces["branch"] != "master":
-                rendered += ".dev0"
-            rendered += plus_or_dot(pieces)
-            rendered += "%d.g%s" % (pieces["distance"], pieces["short"])
-            if pieces["dirty"]:
-                rendered += ".dirty"
-    else:
-        # exception #1
-        rendered = "0"
-        if pieces["branch"] != "master":
-            rendered += ".dev0"
-        rendered += "+untagged.%d.g%s" % (pieces["distance"],
-                                          pieces["short"])
-        if pieces["dirty"]:
-            rendered += ".dirty"
-    return rendered
-
-
-def pep440_split_post(ver: str) -> Tuple[str, Optional[int]]:
-    """Split pep440 version string at the post-release segment.
-
-    Returns the release segments before the post-release and the
-    post-release version number (or -1 if no post-release segment is present).
-    """
-    vc = str.split(ver, ".post")
-    return vc[0], int(vc[1] or 0) if len(vc) == 2 else None
-
-
-def render_pep440_pre(pieces: Dict[str, Any]) -> str:
-    """TAG[.postN.devDISTANCE] -- No -dirty.
-
-    Exceptions:
-    1: no tags. 0.post0.devDISTANCE
-    """
-    if pieces["closest-tag"]:
-        if pieces["distance"]:
-            # update the post release segment
-            tag_version, post_version = pep440_split_post(pieces["closest-tag"])
-            rendered = tag_version
-            if post_version is not None:
-                rendered += ".post%d.dev%d" % (post_version + 1, pieces["distance"])
-            else:
-                rendered += ".post0.dev%d" % (pieces["distance"])
-        else:
-            # no commits, use the tag as the version
-            rendered = pieces["closest-tag"]
-    else:
-        # exception #1
-        rendered = "0.post0.dev%d" % pieces["distance"]
-    return rendered
-
-
-def render_pep440_post(pieces: Dict[str, Any]) -> str:
-    """TAG[.postDISTANCE[.dev0]+gHEX] .
-
-    The ".dev0" means dirty. Note that .dev0 sorts backwards
-    (a dirty tree will appear "older" than the corresponding clean one),
-    but you shouldn't be releasing software with -dirty anyways.
-
-    Exceptions:
-    1: no tags. 0.postDISTANCE[.dev0]
-    """
-    if pieces["closest-tag"]:
-        rendered = pieces["closest-tag"]
-        if pieces["distance"] or pieces["dirty"]:
-            rendered += ".post%d" % pieces["distance"]
-            if pieces["dirty"]:
-                rendered += ".dev0"
-            rendered += plus_or_dot(pieces)
-            rendered += "g%s" % pieces["short"]
-    else:
-        # exception #1
-        rendered = "0.post%d" % pieces["distance"]
-        if pieces["dirty"]:
-            rendered += ".dev0"
-        rendered += "+g%s" % pieces["short"]
-    return rendered
-
-
-def render_pep440_post_branch(pieces: Dict[str, Any]) -> str:
-    """TAG[.postDISTANCE[.dev0]+gHEX[.dirty]] .
-
-    The ".dev0" means not master branch.
-
-    Exceptions:
-    1: no tags. 0.postDISTANCE[.dev0]+gHEX[.dirty]
-    """
-    if pieces["closest-tag"]:
-        rendered = pieces["closest-tag"]
-        if pieces["distance"] or pieces["dirty"]:
-            rendered += ".post%d" % pieces["distance"]
-            if pieces["branch"] != "master":
-                rendered += ".dev0"
-            rendered += plus_or_dot(pieces)
-            rendered += "g%s" % pieces["short"]
-            if pieces["dirty"]:
-                rendered += ".dirty"
-    else:
-        # exception #1
-        rendered = "0.post%d" % pieces["distance"]
-        if pieces["branch"] != "master":
-            rendered += ".dev0"
-        rendered += "+g%s" % pieces["short"]
-        if pieces["dirty"]:
-            rendered += ".dirty"
-    return rendered
-
-
-def render_pep440_old(pieces: Dict[str, Any]) -> str:
-    """TAG[.postDISTANCE[.dev0]] .
-
-    The ".dev0" means dirty.
-
-    Exceptions:
-    1: no tags. 0.postDISTANCE[.dev0]
-    """
-    if pieces["closest-tag"]:
-        rendered = pieces["closest-tag"]
-        if pieces["distance"] or pieces["dirty"]:
-            rendered += ".post%d" % pieces["distance"]
-            if pieces["dirty"]:
-                rendered += ".dev0"
-    else:
-        # exception #1
-        rendered = "0.post%d" % pieces["distance"]
-        if pieces["dirty"]:
-            rendered += ".dev0"
-    return rendered
-
-
-def render_git_describe(pieces: Dict[str, Any]) -> str:
-    """TAG[-DISTANCE-gHEX][-dirty].
-
-    Like 'git describe --tags --dirty --always'.
-
-    Exceptions:
-    1: no tags. HEX[-dirty]  (note: no 'g' prefix)
-    """
-    if pieces["closest-tag"]:
-        rendered = pieces["closest-tag"]
-        if pieces["distance"]:
-            rendered += "-%d-g%s" % (pieces["distance"], pieces["short"])
-    else:
-        # exception #1
-        rendered = pieces["short"]
-    if pieces["dirty"]:
-        rendered += "-dirty"
-    return rendered
-
-
-def render_git_describe_long(pieces: Dict[str, Any]) -> str:
-    """TAG-DISTANCE-gHEX[-dirty].
-
-    Like 'git describe --tags --dirty --always -long'.
-    The distance/hash is unconditional.
-
-    Exceptions:
-    1: no tags. HEX[-dirty]  (note: no 'g' prefix)
-    """
-    if pieces["closest-tag"]:
-        rendered = pieces["closest-tag"]
-        rendered += "-%d-g%s" % (pieces["distance"], pieces["short"])
-    else:
-        # exception #1
-        rendered = pieces["short"]
-    if pieces["dirty"]:
-        rendered += "-dirty"
-    return rendered
-
-
-def render(pieces: Dict[str, Any], style: str) -> Dict[str, Any]:
-    """Render the given version pieces into the requested style."""
-    if pieces["error"]:
-        return {"version": "unknown",
-                "full-revisionid": pieces.get("long"),
-                "dirty": None,
-                "error": pieces["error"],
-                "date": None}
-
-    if not style or style == "default":
-        style = "pep440"  # the default
-
-    if style == "pep440":
-        rendered = render_pep440(pieces)
-    elif style == "pep440-branch":
-        rendered = render_pep440_branch(pieces)
-    elif style == "pep440-pre":
-        rendered = render_pep440_pre(pieces)
-    elif style == "pep440-post":
-        rendered = render_pep440_post(pieces)
-    elif style == "pep440-post-branch":
-        rendered = render_pep440_post_branch(pieces)
-    elif style == "pep440-old":
-        rendered = render_pep440_old(pieces)
-    elif style == "git-describe":
-        rendered = render_git_describe(pieces)
-    elif style == "git-describe-long":
-        rendered = render_git_describe_long(pieces)
-    else:
-        raise ValueError("unknown style '%s'" % style)
-
-    return {"version": rendered, "full-revisionid": pieces["long"],
-            "dirty": pieces["dirty"], "error": None,
-            "date": pieces.get("date")}
-
-
-def get_versions() -> Dict[str, Any]:
-    """Get version information or return default if unable to do so."""
-    # I am in _version.py, which lives at ROOT/VERSIONFILE_SOURCE. If we have
-    # __file__, we can work backwards from there to the root. Some
-    # py2exe/bbfreeze/non-CPython implementations don't do __file__, in which
-    # case we can only use expanded keywords.
-
-    cfg = get_config()
-    verbose = cfg.verbose
-
-    try:
-        return git_versions_from_keywords(get_keywords(), cfg.tag_prefix,
-                                          verbose)
-    except NotThisMethod:
-        pass
-
-    try:
-        root = os.path.realpath(__file__)
-        # versionfile_source is the relative path from the top of the source
-        # tree (where the .git directory might live) to this file. Invert
-        # this to find the root from __file__.
-        for _ in cfg.versionfile_source.split('/'):
-            root = os.path.dirname(root)
-    except NameError:
-        return {"version": "0+unknown", "full-revisionid": None,
-                "dirty": None,
-                "error": "unable to find root of source tree",
-                "date": None}
-
-    try:
-        pieces = git_pieces_from_vcs(cfg.tag_prefix, root, verbose)
-        return render(pieces, cfg.style)
-    except NotThisMethod:
-        pass
-
-    try:
-        if cfg.parentdir_prefix:
-            return versions_from_parentdir(cfg.parentdir_prefix, root, verbose)
-    except NotThisMethod:
-        pass
-
-    return {"version": "0+unknown", "full-revisionid": None,
-            "dirty": None,
-            "error": "unable to compute version", "date": None}
diff --git a/numba_dpex/core/debuginfo.py b/numba_dpex/core/debuginfo.py
new file mode 100644
index 0000000000..2f16f6d265
--- /dev/null
+++ b/numba_dpex/core/debuginfo.py
@@ -0,0 +1,43 @@
+# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Implements a custom debug metadata generator class for numba-dpex kernels.
+"""
+
+from numba.core import debuginfo
+
+
+class DIBuilder(debuginfo.DIBuilder):
+    """Overrides Numba's default DIBuilder with numba-dpex-specific customizations."""
+
+    # pylint: disable=too-many-arguments
+    def mark_subprogram(self, function, qualname, argnames, argtypes, line):
+        """Sets DW_AT_name and DW_AT_linkagename tags for a kernel decorated function.
+
+        Numba generates a unique name for every function it compiles, but in
+        upstream Numba the unique name is not used as the "qualified" name of
+        the function. The behavior leads to a bug discovered in Numba-dpex when
+        a compiled function uses closure variables.
+        Refer (https://github.com/IntelPython/numba-dpex/issues/898).
+        To resolve the issue numba-dpex uses the unique_name as the qualified
+        name. Refer to
+        :class:`numba_dpex.core.passes.passes.QualNameDisambiguationLowering`.
+        However, doing so breaks setting GDB breakpoints based on function
+        name as the function name is no longer what is in the source, but what
+        is the unique name generated by Numba. To fix it, numba-dpex uses a
+        modified DISubprogram metadata generator. The name (DW_AT_name) tag is
+        set to the base function name, discarding the unique qualifier and
+        linkagename is set to an empty string.
+        """
+        name = qualname[0 : qualname.find("$")]  # noqa: E203
+        argmap = dict(zip(argnames, argtypes))
+
+        di_subp = self._add_subprogram(
+            name=name,
+            linkagename="",
+            line=line,
+            function=function,
+            argmap=argmap,
+        )
+        function.set_metadata("dbg", di_subp)
diff --git a/numba_dpex/core/descriptor.py b/numba_dpex/core/descriptor.py
index 369838ed25..f8637103c1 100644
--- a/numba_dpex/core/descriptor.py
+++ b/numba_dpex/core/descriptor.py
@@ -48,6 +48,8 @@ class DpexTargetOptions(CPUTargetOptions):
     no_compile = _option_mapping("no_compile")
     inline_threshold = _option_mapping("inline_threshold")
     _compilation_mode = _option_mapping("_compilation_mode")
+    # TODO: create separate parfor kernel target
+    _parfor_body_args = _option_mapping("_parfor_body_args")
 
     def finalize(self, flags, options):
         super().finalize(flags, options)
@@ -63,6 +65,7 @@ def finalize(self, flags, options):
         _inherit_if_not_set(
             flags, options, "_compilation_mode", CompilationMode.KERNEL
         )
+        _inherit_if_not_set(flags, options, "_parfor_body_args", None)
 
 
 class DpexKernelTarget(TargetDescriptor):
diff --git a/numba_dpex/core/parfors/compiler.py b/numba_dpex/core/parfors/compiler.py
deleted file mode 100644
index 8864cc637c..0000000000
--- a/numba_dpex/core/parfors/compiler.py
+++ /dev/null
@@ -1,88 +0,0 @@
-# SPDX-FileCopyrightText: 2022 - 2024 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-from numba.core import compiler, ir
-from numba.core import types as numba_types
-from numba.core.compiler_lock import global_compiler_lock
-
-from numba_dpex.core import config
-from numba_dpex.core.exceptions import (
-    KernelHasReturnValueError,
-    UnreachableError,
-)
-from numba_dpex.core.pipelines.kernel_compiler import KernelCompiler
-
-
-@global_compiler_lock
-def compile_numba_ir_with_dpex(
-    pyfunc,
-    pyfunc_name,
-    args,
-    return_type,
-    target_context,
-    typing_context,
-    debug=False,
-    is_kernel=True,
-    extra_compile_flags=None,
-):
-    """
-    Compiles a function using class:`numba_dpex.core.pipelines.KernelCompiler`
-    and returns the compiled result.
-
-    Args:
-        args: The list of arguments passed to the kernel.
-        debug (bool): Optional flag to turn on debug mode compilation.
-        extra_compile_flags: Extra flags passed to the compiler.
-
-    Returns:
-        cres: Compiled result.
-
-    Raises:
-        KernelHasReturnValueError: If the compiled function returns a
-            non-void value.
-    """
-    # First compilation will trigger the initialization of the backend.
-    typingctx = typing_context
-    targetctx = target_context
-
-    flags = compiler.Flags()
-    # Do not compile the function to a binary, just lower to LLVM
-    flags.debuginfo = config.DEBUGINFO_DEFAULT
-    flags.no_compile = True
-    flags.no_cpython_wrapper = True
-    flags.no_cfunc_wrapper = True
-    flags.nrt = False
-
-    if debug:
-        flags.debuginfo = debug
-
-    # Run compilation pipeline
-    if isinstance(pyfunc, ir.FunctionIR):
-        cres = compiler.compile_ir(
-            typingctx=typingctx,
-            targetctx=targetctx,
-            func_ir=pyfunc,
-            args=args,
-            return_type=return_type,
-            flags=flags,
-            locals={},
-            pipeline_class=KernelCompiler,
-        )
-    else:
-        raise UnreachableError()
-
-    if (
-        is_kernel
-        and cres.signature.return_type is not None
-        and cres.signature.return_type != numba_types.void
-    ):
-        raise KernelHasReturnValueError(
-            kernel_name=pyfunc_name,
-            return_type=cres.signature.return_type,
-        )
-    # Linking depending libraries
-    library = cres.library
-    library.finalize()
-
-    return cres
diff --git a/numba_dpex/core/parfors/kernel_builder.py b/numba_dpex/core/parfors/kernel_builder.py
index 1eda1664ae..23aa46c7f9 100644
--- a/numba_dpex/core/parfors/kernel_builder.py
+++ b/numba_dpex/core/parfors/kernel_builder.py
@@ -3,114 +3,52 @@
 # SPDX-License-Identifier: Apache-2.0
 
 import copy
-import sys
 import warnings
 
-import dpctl
-import dpctl.program as dpctl_prog
 from numba.core import ir, types
 from numba.core.errors import NumbaParallelSafetyWarning
 from numba.core.ir_utils import (
-    add_offset_to_labels,
     get_name_var_table,
     get_unused_var_name,
     legalize_names,
     mk_unique_var,
-    remove_dead,
     remove_dels,
-    rename_labels,
     replace_var_names,
 )
-from numba.core.target_extension import target_override
 from numba.core.typing import signature
 from numba.parfors import parfor
 
 from numba_dpex.core import config
+from numba_dpex.core.decorators import kernel
+from numba_dpex.core.parfors.parfor_sentinel_replace_pass import (
+    ParforBodyArguments,
+)
 from numba_dpex.core.types.kernel_api.index_space_ids import ItemType
-from numba_dpex.kernel_api_impl.spirv import spirv_generator
+from numba_dpex.core.utils.call_kernel_builder import SPIRVKernelModule
+from numba_dpex.kernel_api_impl.spirv.dispatcher import (
+    SPIRVKernelDispatcher,
+    _SPIRVKernelCompileResult,
+)
 
-from ..descriptor import dpex_kernel_target
-from ..types import DpnpNdArray
-from .compiler import compile_numba_ir_with_dpex
 from .kernel_templates.range_kernel_template import RangeKernelTemplate
 
 
 class ParforKernel:
     def __init__(
         self,
-        name,
-        kernel,
         signature,
         kernel_args,
         kernel_arg_types,
-        queue: dpctl.SyclQueue,
         local_accessors=None,
         work_group_size=None,
+        kernel_module=None,
     ):
-        self.name = name
-        self.kernel = kernel
         self.signature = signature
         self.kernel_args = kernel_args
         self.kernel_arg_types = kernel_arg_types
-        self.queue = queue
         self.local_accessors = local_accessors
         self.work_group_size = work_group_size
-
-
-def _print_block(block):
-    for i, inst in enumerate(block.body):
-        print("    ", i, inst)
-
-
-def _print_body(body_dict):
-    """Pretty-print a set of IR blocks."""
-    for label, block in body_dict.items():
-        print("label: ", label)
-        _print_block(block)
-
-
-def _compile_kernel_parfor(
-    sycl_queue, kernel_name, func_ir, argtypes, debug=False
-):
-    with target_override(dpex_kernel_target.target_context.target_name):
-        cres = compile_numba_ir_with_dpex(
-            pyfunc=func_ir,
-            pyfunc_name=kernel_name,
-            args=argtypes,
-            return_type=None,
-            debug=debug,
-            is_kernel=True,
-            typing_context=dpex_kernel_target.typing_context,
-            target_context=dpex_kernel_target.target_context,
-            extra_compile_flags=None,
-        )
-    cres.library.inline_threshold = config.INLINE_THRESHOLD
-    cres.library._optimize_final_module()
-    func = cres.library.get_function(cres.fndesc.llvm_func_name)
-    kernel = dpex_kernel_target.target_context.prepare_spir_kernel(
-        func, cres.signature.args
-    )
-    spirv_module = spirv_generator.llvm_to_spirv(
-        dpex_kernel_target.target_context,
-        kernel.module.__str__(),
-        kernel.module.as_bitcode(),
-    )
-
-    dpctl_create_program_from_spirv_flags = []
-    if debug or config.DPEX_OPT == 0:
-        # if debug is ON we need to pass additional flags to igc.
-        dpctl_create_program_from_spirv_flags = ["-g", "-cl-opt-disable"]
-
-    # create a sycl::kernel_bundle
-    kernel_bundle = dpctl_prog.create_program_from_spirv(
-        sycl_queue,
-        spirv_module,
-        " ".join(dpctl_create_program_from_spirv_flags),
-    )
-    #  create a sycl::kernel
-    sycl_kernel = kernel_bundle.get_sycl_kernel(kernel.name)
-
-    return sycl_kernel
+        self.kernel_module = kernel_module
 
 
 def _legalize_names_with_typemap(names, typemap):
@@ -189,76 +127,11 @@ def _replace_var_with_array(vars, loop_body, typemap, calltypes):
         typemap[v] = types.npytypes.Array(el_typ, 1, "C")
 
 
-def _find_setitems_block(setitems, block, typemap):
-    for inst in block.body:
-        if isinstance(inst, ir.StaticSetItem) or isinstance(inst, ir.SetItem):
-            setitems.add(inst.target.name)
-        elif isinstance(inst, parfor.Parfor):
-            _find_setitems_block(setitems, inst.init_block, typemap)
-            _find_setitems_body(setitems, inst.loop_body, typemap)
-
-
-def _find_setitems_body(setitems, loop_body, typemap):
-    """
-    Find the arrays that are written into (goes into setitems)
-    """
-    for label, block in loop_body.items():
-        _find_setitems_block(setitems, block, typemap)
-
-
-def _replace_sentinel_with_parfor_body(kernel_ir, sentinel_name, loop_body):
-    # new label for splitting sentinel block
-    new_label = max(loop_body.keys()) + 1
-
-    # Search all the block in the kernel function for the sentinel assignment.
-    for label, block in kernel_ir.blocks.items():
-        for i, inst in enumerate(block.body):
-            if (
-                isinstance(inst, ir.Assign)
-                and inst.target.name == sentinel_name
-            ):
-                # We found the sentinel assignment.
-                loc = inst.loc
-                scope = block.scope
-                # split block across __sentinel__
-                # A new block is allocated for the statements prior to the
-                # sentinel but the new block maintains the current block label.
-                prev_block = ir.Block(scope, loc)
-                prev_block.body = block.body[:i]
-
-                # The current block is used for statements after the sentinel.
-                block.body = block.body[i + 1 :]  # noqa: E203
-                # But the current block gets a new label.
-                body_first_label = min(loop_body.keys())
-
-                # The previous block jumps to the minimum labelled block of the
-                # parfor body.
-                prev_block.append(ir.Jump(body_first_label, loc))
-                # Add all the parfor loop body blocks to the kernel function's
-                # IR.
-                for loop, b in loop_body.items():
-                    kernel_ir.blocks[loop] = b
-                body_last_label = max(loop_body.keys())
-                kernel_ir.blocks[new_label] = block
-                kernel_ir.blocks[label] = prev_block
-                # Add a jump from the last parfor body block to the block
-                # containing statements after the sentinel.
-                kernel_ir.blocks[body_last_label].append(
-                    ir.Jump(new_label, loc)
-                )
-                break
-        else:
-            continue
-        break
-
-
 def create_kernel_for_parfor(
     lowerer,
     parfor_node,
     typemap,
-    flags,
     loop_ranges,
-    has_aliases,
     races,
     parfor_outputs,
 ) -> ParforKernel:
@@ -367,120 +240,38 @@ def create_kernel_for_parfor(
         loop_ranges=loop_ranges,
         param_dict=param_dict,
     )
-    kernel_ir = kernel_template.kernel_ir
 
-    if config.DEBUG_ARRAY_OPT:
-        print("kernel_ir dump ", type(kernel_ir))
-        kernel_ir.dump()
-        print("loop_body dump ", type(loop_body))
-        _print_body(loop_body)
-
-    # rename all variables in kernel_ir afresh
-    var_table = get_name_var_table(kernel_ir.blocks)
-    new_var_dict = {}
-    reserved_names = (
-        [sentinel_name] + list(param_dict.values()) + legal_loop_indices
+    kernel_dispatcher: SPIRVKernelDispatcher = kernel(
+        kernel_template.py_func,
+        _parfor_body_args=ParforBodyArguments(
+            loop_body=loop_body,
+            param_dict=param_dict,
+            legal_loop_indices=legal_loop_indices,
+        ),
     )
-    for name, var in var_table.items():
-        if not (name in reserved_names):
-            new_var_dict[name] = mk_unique_var(name)
-    replace_var_names(kernel_ir.blocks, new_var_dict)
-    if config.DEBUG_ARRAY_OPT:
-        print("kernel_ir dump after renaming ")
-        kernel_ir.dump()
-
-    kernel_param_types = param_types
 
-    if config.DEBUG_ARRAY_OPT:
-        print(
-            "kernel_param_types = ",
-            type(kernel_param_types),
-            "\n",
-            kernel_param_types,
-        )
-
-    kernel_stub_last_label = max(kernel_ir.blocks.keys()) + 1
-
-    # Add kernel stub last label to each parfor.loop_body label to prevent
-    # label conflicts.
-    loop_body = add_offset_to_labels(loop_body, kernel_stub_last_label)
-
-    _replace_sentinel_with_parfor_body(kernel_ir, sentinel_name, loop_body)
-
-    if config.DEBUG_ARRAY_OPT:
-        print("kernel_ir last dump before renaming")
-        kernel_ir.dump()
-
-    kernel_ir.blocks = rename_labels(kernel_ir.blocks)
-    remove_dels(kernel_ir.blocks)
-
-    old_alias = flags.noalias
-    if not has_aliases:
-        if config.DEBUG_ARRAY_OPT:
-            print("No aliases found so adding noalias flag.")
-        flags.noalias = True
-
-    remove_dead(kernel_ir.blocks, kernel_ir.arg_names, kernel_ir, typemap)
-
-    if config.DEBUG_ARRAY_OPT:
-        print("kernel_ir after remove dead")
-        kernel_ir.dump()
-
-    # The first argument to a range kernel is a kernel_api.Item object. The
-    # ``Item`` object is used by the kernel_api.spirv backend to generate the
+    # The first argument to a range kernel is a kernel_api.NdItem object. The
+    # ``NdItem`` object is used by the kernel_api.spirv backend to generate the
     # correct SPIR-V indexing instructions. Since, the argument is not something
     # available originally in the kernel_param_types, we add it at this point to
     # make sure the kernel signature matches the actual generated code.
     ty_item = ItemType(parfor_dim)
-    kernel_param_types = (ty_item, *kernel_param_types)
+    kernel_param_types = (ty_item, *param_types)
     kernel_sig = signature(types.none, *kernel_param_types)
 
-    if config.DEBUG_ARRAY_OPT:
-        sys.stdout.flush()
-
-    if config.DEBUG_ARRAY_OPT:
-        print("after DUFunc inline".center(80, "-"))
-        kernel_ir.dump()
-
-    # The ParforLegalizeCFD pass has already ensured that the LHS and RHS
-    # arrays are on same device. We can take the queue from the first input
-    # array and use that to compile the kernel.
-
-    exec_queue: dpctl.SyclQueue = None
-
-    for arg in parfor_args:
-        obj = typemap[arg]
-        if isinstance(obj, DpnpNdArray):
-            filter_string = obj.queue.sycl_device
-            # FIXME: A better design is required so that we do not have to
-            # create a queue every time.
-            exec_queue = dpctl.get_device_cached_queue(filter_string)
-
-    if not exec_queue:
-        raise AssertionError(
-            "No execution found for parfor. No way to compile the kernel!"
-        )
-
-    sycl_kernel = _compile_kernel_parfor(
-        exec_queue,
-        kernel_name,
-        kernel_ir,
-        kernel_param_types,
-        debug=flags.debuginfo,
+    kcres: _SPIRVKernelCompileResult = kernel_dispatcher.get_compile_result(
+        types.void(*kernel_param_types)  # kernel signature
     )
-
-    flags.noalias = old_alias
+    kernel_module: SPIRVKernelModule = kcres.kernel_device_ir_module
 
     if config.DEBUG_ARRAY_OPT:
         print("kernel_sig = ", kernel_sig)
 
     return ParforKernel(
-        name=kernel_name,
-        kernel=sycl_kernel,
         signature=kernel_sig,
         kernel_args=parfor_args,
         kernel_arg_types=func_arg_types,
-        queue=exec_queue,
+        kernel_module=kernel_module,
     )
 
 
diff --git a/numba_dpex/core/parfors/kernel_templates/kernel_template_iface.py b/numba_dpex/core/parfors/kernel_templates/kernel_template_iface.py
index 453017dbbb..9bb7457218 100644
--- a/numba_dpex/core/parfors/kernel_templates/kernel_template_iface.py
+++ b/numba_dpex/core/parfors/kernel_templates/kernel_template_iface.py
@@ -30,13 +30,9 @@ def _generate_kernel_ir(self):
     def dump_kernel_string(self):
         raise NotImplementedError
 
-    @abc.abstractmethod
-    def dump_kernel_ir(self):
-        raise NotImplementedError
-
     @property
     @abc.abstractmethod
-    def kernel_ir(self):
+    def py_func(self):
         raise NotImplementedError
 
     @property
diff --git a/numba_dpex/core/parfors/kernel_templates/range_kernel_template.py b/numba_dpex/core/parfors/kernel_templates/range_kernel_template.py
index ecbfa5df4c..f5dd4304e1 100644
--- a/numba_dpex/core/parfors/kernel_templates/range_kernel_template.py
+++ b/numba_dpex/core/parfors/kernel_templates/range_kernel_template.py
@@ -5,7 +5,6 @@
 import sys
 
 import dpnp
-from numba.core import compiler
 
 import numba_dpex as dpex
 
@@ -51,7 +50,7 @@ def __init__(
         self._param_dict = param_dict
 
         self._kernel_txt = self._generate_kernel_stub_as_string()
-        self._kernel_ir = self._generate_kernel_ir()
+        self._py_func = self._generate_kernel_ir()
 
     def _generate_kernel_stub_as_string(self):
         """Generates a stub dpex kernel for the parfor as a string.
@@ -109,17 +108,15 @@ def _generate_kernel_ir(self):
         globls = {"dpnp": dpnp, "dpex": dpex}
         locls = {}
         exec(self._kernel_txt, globls, locls)
-        kernel_fn = locls[self._kernel_name]
-
-        return compiler.run_frontend(kernel_fn)
+        return locls[self._kernel_name]
 
     @property
-    def kernel_ir(self):
-        """Returns the Numba IR generated for a RangeKernelTemplate.
-
-        Returns: The Numba functionIR object for the compiled kernel_txt string.
+    def py_func(self):
+        """Returns the python function generated for a
+            TreeReduceIntermediateKernelTemplate.
+        Returns: The python function object for the compiled kernel_txt string.
         """
-        return self._kernel_ir
+        return self._py_func
 
     @property
     def kernel_string(self):
@@ -134,7 +131,3 @@ def dump_kernel_string(self):
         """Helper to print the kernel function string."""
         print(self._kernel_txt)
         sys.stdout.flush()
-
-    def dump_kernel_ir(self):
-        """Helper to dump the Numba IR for the RangeKernelTemplate."""
-        self._kernel_ir.dump()
diff --git a/numba_dpex/core/parfors/kernel_templates/reduction_template.py b/numba_dpex/core/parfors/kernel_templates/reduction_template.py
index d78f8cd449..199a3f2b41 100644
--- a/numba_dpex/core/parfors/kernel_templates/reduction_template.py
+++ b/numba_dpex/core/parfors/kernel_templates/reduction_template.py
@@ -6,7 +6,6 @@
 import sys
 
 import dpnp
-from numba.core import compiler
 
 import numba_dpex.kernel_api as kapi
 
@@ -48,7 +47,7 @@ def __init__(
         self._typemap = typemap
 
         self._kernel_txt = self._generate_kernel_stub_as_string()
-        self._kernel_ir = self._generate_kernel_ir()
+        self._py_func = self._generate_kernel_ir()
 
     def _generate_kernel_stub_as_string(self):
         """Generate reduction main kernel template"""
@@ -161,18 +160,15 @@ def _generate_kernel_ir(self):
         globls = {"dpnp": dpnp, "kapi": kapi}
         locls = {}
         exec(self._kernel_txt, globls, locls)
-        kernel_fn = locls[self._kernel_name]
-
-        return compiler.run_frontend(kernel_fn)
+        return locls[self._kernel_name]
 
     @property
-    def kernel_ir(self):
-        """Returns the Numba IR generated for a
+    def py_func(self):
+        """Returns the python function generated for a
             TreeReduceIntermediateKernelTemplate.
-
-        Returns: The Numba functionIR object for the compiled kernel_txt string.
+        Returns: The python function object for the compiled kernel_txt string.
         """
-        return self._kernel_ir
+        return self._py_func
 
     @property
     def kernel_string(self):
@@ -190,11 +186,6 @@ def dump_kernel_string(self):
         print(self._kernel_txt)
         sys.stdout.flush()
 
-    def dump_kernel_ir(self):
-        """Helper to dump the Numba IR for a
-        TreeReduceIntermediateKernelTemplate."""
-        self._kernel_ir.dump()
-
 
 class RemainderReduceIntermediateKernelTemplate(KernelTemplateInterface):
     """The class to build reduction remainder kernel_txt template and
@@ -234,7 +225,7 @@ def __init__(
         self._reductionKernelVar = reductionKernelVar
 
         self._kernel_txt = self._generate_kernel_stub_as_string()
-        self._kernel_ir = self._generate_kernel_ir()
+        self._py_func = self._generate_kernel_ir()
 
     def _generate_kernel_stub_as_string(self):
         """Generate reduction remainder kernel template"""
@@ -320,18 +311,15 @@ def _generate_kernel_ir(self):
         globls = {"dpnp": dpnp, "kapi": kapi}
         locls = {}
         exec(self._kernel_txt, globls, locls)
-        kernel_fn = locls[self._kernel_name]
-
-        return compiler.run_frontend(kernel_fn)
+        return locls[self._kernel_name]
 
     @property
-    def kernel_ir(self):
-        """Returns the Numba IR generated for a
-            RemainderReduceIntermediateKernelTemplate.
-
-        Returns: The Numba functionIR object for the compiled kernel_txt string.
+    def py_func(self):
+        """Returns the python function generated for a
+            TreeReduceIntermediateKernelTemplate.
+        Returns: The python function object for the compiled kernel_txt string.
         """
-        return self._kernel_ir
+        return self._py_func
 
     @property
     def kernel_string(self):
@@ -349,9 +337,3 @@ def dump_kernel_string(self):
 
         print(self._kernel_txt)
         sys.stdout.flush()
-
-    def dump_kernel_ir(self):
-        """Helper to dump the Numba IR for the
-        RemainderReduceIntermediateKernelTemplate."""
-
-        self._kernel_ir.dump()
diff --git a/numba_dpex/core/parfors/parfor_lowerer.py b/numba_dpex/core/parfors/parfor_lowerer.py
index 41bf86868f..4610cf9ef4 100644
--- a/numba_dpex/core/parfors/parfor_lowerer.py
+++ b/numba_dpex/core/parfors/parfor_lowerer.py
@@ -3,7 +3,6 @@
 # SPDX-License-Identifier: Apache-2.0
 
 import copy
-from collections import namedtuple
 
 from llvmlite import ir as llvmir
 from numba.core import cgutils, ir, types
@@ -31,9 +30,6 @@
     create_reduction_remainder_kernel_for_parfor,
 )
 
-# A global list of kernels to keep the objects alive indefinitely.
-keep_alive_kernels = []
-
 
 def _getvar(lowerer, x):
     """Returns the LLVM Value corresponding to a Numba IR variable.
@@ -154,20 +150,16 @@ def _submit_parfor_kernel(
         kernel_fn: ParforKernel,
         global_range,
         local_range,
+        debug=False,
     ):
         """
         Adds a call to submit a kernel function into the function body of the
         current Numba JIT compiled function.
         """
-        # Ensure that the Python arguments are kept alive for the duration of
-        # the kernel execution
-        keep_alive_kernels.append(kernel_fn.kernel)
         kl_builder = KernelLaunchIRBuilder(
             lowerer.context, lowerer.builder, kernel_dmm
         )
 
-        queue_ref = kl_builder.get_queue(exec_queue=kernel_fn.queue)
-
         kernel_args = []
         for i, arg in enumerate(kernel_fn.kernel_args):
             if (
@@ -188,24 +180,23 @@ def _submit_parfor_kernel(
             else:
                 kernel_args.append(_getvar(lowerer, arg))
 
-        kernel_ref_addr = kernel_fn.kernel.addressof_ref()
-        kernel_ref = lowerer.builder.inttoptr(
-            lowerer.context.get_constant(types.uintp, kernel_ref_addr),
-            cgutils.voidptr_t,
-        )
-
-        kl_builder.set_kernel(kernel_ref)
-        kl_builder.set_queue(queue_ref)
         kl_builder.set_range(global_range, local_range)
         kl_builder.set_arguments(
             kernel_fn.kernel_arg_types, kernel_args=kernel_args
         )
+        kl_builder.set_queue_from_arguments()
         kl_builder.set_dependent_events([])
+        kl_builder.set_kernel_from_spirv(
+            kernel_fn.kernel_module,
+            debug=debug,
+        )
+
         event_ref = kl_builder.submit()
 
         sycl.dpctl_event_wait(lowerer.builder, event_ref)
         sycl.dpctl_event_delete(lowerer.builder, event_ref)
-        sycl.dpctl_queue_delete(lowerer.builder, queue_ref)
+
+        return kl_builder.arguments.sycl_queue_ref
 
     def _reduction_codegen(
         self,
@@ -263,8 +254,6 @@ def _reduction_codegen(
             loop_ranges,
             parfor,
             typemap,
-            flags,
-            bool(alias_map),
             reductionKernelVar,
             parfor_reddict,
         )
@@ -278,13 +267,12 @@ def _reduction_codegen(
             parfor_kernel,
             global_range,
             local_range,
+            debug=flags.debuginfo,
         )
 
         parfor_kernel = create_reduction_remainder_kernel_for_parfor(
             parfor,
             typemap,
-            flags,
-            bool(alias_map),
             reductionKernelVar,
             parfor_reddict,
             reductionHelperList,
@@ -292,14 +280,16 @@ def _reduction_codegen(
 
         global_range, local_range = self._remainder_ranges(lowerer)
 
-        self._submit_parfor_kernel(
+        # TODO: find better way to pass queue
+        queue_ref = self._submit_parfor_kernel(
             lowerer,
             parfor_kernel,
             global_range,
             local_range,
+            debug=flags.debuginfo,
         )
 
-        reductionKernelVar.copy_final_sum_to_host(parfor_kernel)
+        reductionKernelVar.copy_final_sum_to_host(queue_ref)
 
     def _lower_parfor_as_kernel(self, lowerer, parfor):
         """Lowers a parfor node created by the dpjit compiler to a
@@ -400,9 +390,7 @@ def _lower_parfor_as_kernel(self, lowerer, parfor):
                     lowerer,
                     parfor,
                     typemap,
-                    flags,
                     loop_ranges,
-                    bool(alias_map),
                     parfor.races,
                     parfor_output_arrays,
                 )
@@ -418,10 +406,9 @@ def _lower_parfor_as_kernel(self, lowerer, parfor):
                 parfor_kernel,
                 global_range,
                 local_range,
+                debug=flags.debuginfo,
             )
 
-        # TODO: free the kernel at this point
-
         # Restore the original typemap of the function that was replaced
         # temporarily at the beginning of this function.
         lowerer.fndesc.typemap = orig_typemap
diff --git a/numba_dpex/core/parfors/parfor_pass.py b/numba_dpex/core/parfors/parfor_pass.py
index c2149bf0cf..3901af8e7c 100644
--- a/numba_dpex/core/parfors/parfor_pass.py
+++ b/numba_dpex/core/parfors/parfor_pass.py
@@ -17,8 +17,8 @@
 from numba.core import config, errors, ir, types, typing
 from numba.core.compiler_machinery import register_pass
 from numba.core.ir_utils import (
+    convert_size_to_var,
     dprint_func_ir,
-    mk_alloc,
     mk_unique_var,
     next_label,
 )
@@ -43,6 +43,7 @@
 )
 from numba.stencils.stencilparfor import StencilPass
 
+from numba_dpex.core.types.dpnp_ndarray_type import DpnpNdArray
 from numba_dpex.core.typing import dpnpdecl
 
 
@@ -58,6 +59,37 @@ class ConvertDPNPPass(ConvertNumpyPass):
     def __init__(self, pass_states):
         super().__init__(pass_states)
 
+    def _get_queue(self, queue_type, expr: tuple):
+        """
+        Extracts queue from the input arguments of the array operation.
+        """
+        pass_states = self.pass_states
+        typemap: map[str, any] = pass_states.typemap
+
+        var_with_queue = None
+
+        for var in expr[1]:
+            if isinstance(var, tuple):
+                res = self._get_queue(queue_type, var)
+                if res is not None:
+                    return res
+
+                continue
+
+            if not isinstance(var, ir.Var):
+                continue
+
+            _type = typemap[var.name]
+            if not isinstance(_type, DpnpNdArray):
+                continue
+            if queue_type != _type.queue:
+                continue
+
+            var_with_queue = var
+            break
+
+        return ir.Expr.getattr(var_with_queue, "sycl_queue", var_with_queue.loc)
+
     def _arrayexpr_to_parfor(self, equiv_set, lhs, arrayexpr, avail_vars):
         """generate parfor from arrayexpr node, which is essentially a
         map with recursive tree.
@@ -77,6 +109,10 @@ def _arrayexpr_to_parfor(self, equiv_set, lhs, arrayexpr, avail_vars):
             pass_states.typemap, size_vars, scope, loc
         )
 
+        # Expr is a tuple
+        ir_queue = self._get_queue(arr_typ.queue, expr)
+        assert ir_queue is not None
+
         # generate init block and body
         init_block = ir.Block(scope, loc)
         init_block.body = mk_alloc(
@@ -89,6 +125,7 @@ def _arrayexpr_to_parfor(self, equiv_set, lhs, arrayexpr, avail_vars):
             scope,
             loc,
             pass_states.typemap[lhs.name],
+            queue_ir_val=ir_queue,
         )
         body_label = next_label()
         body_block = ir.Block(scope, loc)
@@ -469,3 +506,66 @@ def _arrayexpr_tree_to_ir(
     typemap.pop(expr_out_var.name, None)
     typemap[expr_out_var.name] = el_typ
     return out_ir
+
+
+def mk_alloc(
+    typingctx,
+    typemap,
+    calltypes,
+    lhs,
+    size_var,
+    dtype,
+    scope,
+    loc,
+    lhs_typ,
+    **kws,
+):
+    """generate an array allocation with np.empty() and return list of nodes.
+    size_var can be an int variable or tuple of int variables.
+    lhs_typ is the type of the array being allocated.
+
+    Taken from numba, added kws argument to pass it to __allocate__
+    """
+    out = []
+    ndims = 1
+    size_typ = types.intp
+    if isinstance(size_var, tuple):
+        if len(size_var) == 1:
+            size_var = size_var[0]
+            size_var = convert_size_to_var(size_var, typemap, scope, loc, out)
+        else:
+            # tuple_var = build_tuple([size_var...])
+            ndims = len(size_var)
+            tuple_var = ir.Var(scope, mk_unique_var("$tuple_var"), loc)
+            if typemap:
+                typemap[tuple_var.name] = types.containers.UniTuple(
+                    types.intp, ndims
+                )
+            # constant sizes need to be assigned to vars
+            new_sizes = [
+                convert_size_to_var(s, typemap, scope, loc, out)
+                for s in size_var
+            ]
+            tuple_call = ir.Expr.build_tuple(new_sizes, loc)
+            tuple_assign = ir.Assign(tuple_call, tuple_var, loc)
+            out.append(tuple_assign)
+            size_var = tuple_var
+            size_typ = types.containers.UniTuple(types.intp, ndims)
+    if hasattr(lhs_typ, "__allocate__"):
+        return lhs_typ.__allocate__(
+            typingctx,
+            typemap,
+            calltypes,
+            lhs,
+            size_var,
+            dtype,
+            scope,
+            loc,
+            lhs_typ,
+            size_typ,
+            out,
+            **kws,
+        )
+
+    # Unused numba's code..
+    assert False
diff --git a/numba_dpex/core/parfors/parfor_sentinel_replace_pass.py b/numba_dpex/core/parfors/parfor_sentinel_replace_pass.py
new file mode 100644
index 0000000000..b9c5310b01
--- /dev/null
+++ b/numba_dpex/core/parfors/parfor_sentinel_replace_pass.py
@@ -0,0 +1,110 @@
+# SPDX-FileCopyrightText: 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+from typing import NamedTuple, Union
+
+from numba.core import config
+from numba.core.compiler_machinery import FunctionPass, register_pass
+from numba.core.ir_utils import (
+    add_offset_to_labels,
+    get_name_var_table,
+    get_unused_var_name,
+    mk_unique_var,
+    remove_dels,
+    rename_labels,
+    replace_var_names,
+)
+
+
+class ParforBodyArguments(NamedTuple):
+    """
+    Arguments containing information to inject parfor code inside kernel.
+    """
+
+    loop_body: any
+    param_dict: any
+    legal_loop_indices: any
+
+
+def _print_block(block):
+    for i, inst in enumerate(block.body):
+        print("    ", i, inst)
+
+
+def _print_body(body_dict):
+    """Pretty-print a set of IR blocks."""
+    for label, block in body_dict.items():
+        print("label: ", label)
+        _print_block(block)
+
+
+@register_pass(mutates_CFG=True, analysis_only=False)
+class ParforSentinelReplacePass(FunctionPass):
+    _name = "sentinel_inject"
+
+    def __init__(self):
+        FunctionPass.__init__(self)
+
+    def _get_parfor_body_args(self, flags) -> Union[ParforBodyArguments, None]:
+        if not hasattr(flags, "_parfor_body_args"):
+            return None
+
+        return flags._parfor_body_args
+
+    def run_pass(self, state):
+        flags = state["flags"]
+
+        args = self._get_parfor_body_args(flags)
+
+        if args is None:
+            return True
+
+        # beginning
+        kernel_ir = state["func_ir"]
+        loop_body = args.loop_body
+
+        if config.DEBUG_ARRAY_OPT:
+            print("kernel_ir dump ", type(kernel_ir))
+            kernel_ir.dump()
+            print("loop_body dump ", type(loop_body))
+            _print_body(loop_body)
+
+        # Determine the unique names of the scheduling and kernel functions.
+        loop_body_var_table = get_name_var_table(loop_body)
+        sentinel_name = get_unused_var_name("__sentinel__", loop_body_var_table)
+
+        # rename all variables in kernel_ir afresh
+        var_table = get_name_var_table(kernel_ir.blocks)
+        new_var_dict = {}
+        reserved_names = (
+            [sentinel_name]
+            + list(args.param_dict.values())
+            + args.legal_loop_indices
+        )
+        for name, _ in var_table.items():
+            if not (name in reserved_names):
+                new_var_dict[name] = mk_unique_var(name)
+
+        replace_var_names(kernel_ir.blocks, new_var_dict)
+
+        kernel_stub_last_label = max(kernel_ir.blocks.keys()) + 1
+        loop_body = add_offset_to_labels(loop_body, kernel_stub_last_label)
+
+        # new label for splitting sentinel block
+        new_label = max(loop_body.keys()) + 1
+
+        from .kernel_builder import update_sentinel  # circular
+
+        update_sentinel(kernel_ir, sentinel_name, loop_body, new_label)
+
+        # FIXME: Why rename and remove dels causes the partial_sum array update
+        # instructions to be removed.
+        kernel_ir.blocks = rename_labels(kernel_ir.blocks)
+        remove_dels(kernel_ir.blocks)
+
+        if config.DEBUG_ARRAY_OPT:
+            print("kernel_ir after remove dead")
+            kernel_ir.dump()
+
+        return True
diff --git a/numba_dpex/core/parfors/reduction_helper.py b/numba_dpex/core/parfors/reduction_helper.py
index 552575fbfe..eeeff3e486 100644
--- a/numba_dpex/core/parfors/reduction_helper.py
+++ b/numba_dpex/core/parfors/reduction_helper.py
@@ -17,10 +17,6 @@
 from numba.parfors import parfor
 from numba.parfors.parfor_lowering_utils import ParforLoweringBuilder
 
-from numba_dpex.core.datamodel.models import (
-    dpex_data_model_manager as kernel_dmm,
-)
-from numba_dpex.core.utils.call_kernel_builder import KernelLaunchIRBuilder
 from numba_dpex.core.utils.cgutils_extra import get_llvm_type
 from numba_dpex.dpctl_iface import libsyclinterface_bindings as sycl
 
@@ -275,12 +271,14 @@ def __init__(
         param_dict = _legalize_names_with_typemap(parfor_params, typemap)
         ind_dict = _legalize_names_with_typemap(loop_indices, typemap)
         self._parfor_reddict = parfor_reddict
+        # output of reduction
         self._parfor_redvars = parfor_redvars
         self._redvars_legal_dict = legalize_names(parfor_redvars)
         # Compute a new list of legal loop index names.
         legal_loop_indices = [ind_dict[v] for v in loop_indices]
 
         tmp1 = []
+        # output of reduction computed on device
         self._final_sum_names = []
         self._parfor_redvars_to_redarrs = {}
         for ele1 in reductionHelperList:
@@ -397,16 +395,8 @@ def lowerer(self):
     def work_group_size(self):
         return self._work_group_size
 
-    def copy_final_sum_to_host(self, parfor_kernel):
+    def copy_final_sum_to_host(self, queue_ref):
         lowerer = self.lowerer
-        kl_builder = KernelLaunchIRBuilder(
-            lowerer.context, lowerer.builder, kernel_dmm
-        )
-
-        # Create a local variable storing a pointer to a DPCTLSyclQueueRef
-        # pointer.
-        queue_ref = kl_builder.get_queue(exec_queue=parfor_kernel.queue)
-
         builder = lowerer.builder
         context = lowerer.context
 
@@ -448,5 +438,3 @@ def copy_final_sum_to_host(self, parfor_kernel):
             event_ref = sycl.dpctl_queue_memcpy(builder, *args)
             sycl.dpctl_event_wait(builder, event_ref)
             sycl.dpctl_event_delete(builder, event_ref)
-
-        sycl.dpctl_queue_delete(builder, queue_ref)
diff --git a/numba_dpex/core/parfors/reduction_kernel_builder.py b/numba_dpex/core/parfors/reduction_kernel_builder.py
index 6180ff5ef1..90002f6321 100644
--- a/numba_dpex/core/parfors/reduction_kernel_builder.py
+++ b/numba_dpex/core/parfors/reduction_kernel_builder.py
@@ -4,33 +4,29 @@
 
 import warnings
 
-import dpctl
 from numba.core import types
 from numba.core.errors import NumbaParallelSafetyWarning
 from numba.core.ir_utils import (
-    add_offset_to_labels,
     get_name_var_table,
     get_unused_var_name,
     legalize_names,
-    mk_unique_var,
-    remove_dels,
-    rename_labels,
-    replace_var_names,
 )
 from numba.core.typing import signature
 
+from numba_dpex.core.decorators import kernel
+from numba_dpex.core.parfors.parfor_sentinel_replace_pass import (
+    ParforBodyArguments,
+)
 from numba_dpex.core.parfors.reduction_helper import ReductionKernelVariables
-from numba_dpex.core.types import DpctlSyclQueue
 from numba_dpex.core.types.kernel_api.index_space_ids import NdItemType
 from numba_dpex.core.types.kernel_api.local_accessor import LocalAccessorType
-
-from .kernel_builder import _print_body  # saved for debug
-from .kernel_builder import (
-    ParforKernel,
-    _compile_kernel_parfor,
-    _to_scalar_from_0d,
-    update_sentinel,
+from numba_dpex.core.utils.call_kernel_builder import SPIRVKernelModule
+from numba_dpex.kernel_api_impl.spirv.dispatcher import (
+    SPIRVKernelDispatcher,
+    _SPIRVKernelCompileResult,
 )
+
+from .kernel_builder import ParforKernel, _to_scalar_from_0d
 from .kernel_templates.reduction_template import (
     RemainderReduceIntermediateKernelTemplate,
     TreeReduceIntermediateKernelTemplate,
@@ -41,8 +37,6 @@ def create_reduction_main_kernel_for_parfor(
     loop_ranges,
     parfor_node,
     typemap,
-    flags,
-    has_aliases,
     reductionKernelVar: ReductionKernelVariables,
     parfor_reddict=None,
 ):
@@ -111,7 +105,6 @@ def create_reduction_main_kernel_for_parfor(
         local_accessors_dict=local_accessors_dict,
         typemap=typemap,
     )
-    kernel_ir = kernel_template.kernel_ir
 
     for i, name in enumerate(reductionKernelVar.parfor_params):
         try:
@@ -120,40 +113,14 @@ def create_reduction_main_kernel_for_parfor(
         except KeyError:
             pass
 
-    # rename all variables in kernel_ir afresh
-    var_table = get_name_var_table(kernel_ir.blocks)
-    new_var_dict = {}
-    reserved_names = (
-        [sentinel_name]
-        + list(reductionKernelVar.param_dict.values())
-        + reductionKernelVar.legal_loop_indices
-    )
-    for name, _ in var_table.items():
-        if not (name in reserved_names):
-            new_var_dict[name] = mk_unique_var(name)
-
-    replace_var_names(kernel_ir.blocks, new_var_dict)
-    kernel_param_types = parfor_param_types
-    kernel_stub_last_label = max(kernel_ir.blocks.keys()) + 1
-    # Add kernel stub last label to each parfor.loop_body label to prevent
-    # label conflicts.
-    loop_body = add_offset_to_labels(
-        reductionKernelVar.loop_body, kernel_stub_last_label
+    kernel_dispatcher: SPIRVKernelDispatcher = kernel(
+        kernel_template.py_func,
+        _parfor_body_args=ParforBodyArguments(
+            loop_body=reductionKernelVar.loop_body,
+            param_dict=reductionKernelVar.param_dict,
+            legal_loop_indices=reductionKernelVar.legal_loop_indices,
+        ),
     )
-    # new label for splitting sentinel block
-    new_label = max(loop_body.keys()) + 1
-
-    update_sentinel(kernel_ir, sentinel_name, loop_body, new_label)
-
-    # FIXME: Why rename and remove dels causes the partial_sum array update
-    # instructions to be removed.
-    kernel_ir.blocks = rename_labels(kernel_ir.blocks)
-    remove_dels(kernel_ir.blocks)
-
-    old_alias = flags.noalias
-
-    if not has_aliases:
-        flags.noalias = True
 
     # The first argument to a range kernel is a kernel_api.NdItem object. The
     # ``NdItem`` object is used by the kernel_api.spirv backend to generate the
@@ -161,25 +128,13 @@ def create_reduction_main_kernel_for_parfor(
     # available originally in the kernel_param_types, we add it at this point to
     # make sure the kernel signature matches the actual generated code.
     ty_item = NdItemType(parfor_dim)
-    kernel_param_types = (ty_item, *kernel_param_types)
+    kernel_param_types = (ty_item, *parfor_param_types)
     kernel_sig = signature(types.none, *kernel_param_types)
 
-    # FIXME: A better design is required so that we do not have to create a
-    # queue every time.
-    ty_queue: DpctlSyclQueue = typemap[
-        reductionKernelVar.parfor_params[0]
-    ].queue
-    exec_queue = dpctl.get_device_cached_queue(ty_queue.sycl_device)
-
-    sycl_kernel = _compile_kernel_parfor(
-        exec_queue,
-        kernel_name,
-        kernel_ir,
-        kernel_param_types,
-        debug=flags.debuginfo,
+    kcres: _SPIRVKernelCompileResult = kernel_dispatcher.get_compile_result(
+        types.void(*kernel_param_types)  # kernel signature
     )
-
-    flags.noalias = old_alias
+    kernel_module: SPIRVKernelModule = kcres.kernel_device_ir_module
 
     parfor_params = (
         reductionKernelVar.parfor_params.copy()
@@ -187,22 +142,18 @@ def create_reduction_main_kernel_for_parfor(
     )
 
     return ParforKernel(
-        name=kernel_name,
-        kernel=sycl_kernel,
         signature=kernel_sig,
         kernel_args=parfor_params,
         kernel_arg_types=parfor_param_types,
-        queue=exec_queue,
         local_accessors=set(local_accessors_dict.values()),
         work_group_size=reductionKernelVar.work_group_size,
+        kernel_module=kernel_module,
     )
 
 
 def create_reduction_remainder_kernel_for_parfor(
     parfor_node,
     typemap,
-    flags,
-    has_aliases,
     reductionKernelVar,
     parfor_reddict,
     reductionHelperList,
@@ -288,19 +239,6 @@ def create_reduction_remainder_kernel_for_parfor(
         final_sum_var_name=final_sum_var_legal_name,
         reductionKernelVar=reductionKernelVar,
     )
-    kernel_ir = kernel_template.kernel_ir
-
-    var_table = get_name_var_table(kernel_ir.blocks)
-    new_var_dict = {}
-    reserved_names = (
-        [sentinel_name]
-        + list(reductionKernelVar.param_dict.values())
-        + reductionKernelVar.legal_loop_indices
-    )
-    for name, _ in var_table.items():
-        if not (name in reserved_names):
-            new_var_dict[name] = mk_unique_var(name)
-    replace_var_names(kernel_ir.blocks, new_var_dict)
 
     for i, _ in enumerate(reductionKernelVar.parfor_redvars):
         if reductionHelperList[i].global_size_var is not None:
@@ -353,48 +291,27 @@ def create_reduction_remainder_kernel_for_parfor(
                 _to_scalar_from_0d(typemap[final_sum_var_name[i]])
             )
 
-    kernel_param_types = reductionKernelVar.param_types
-
-    kernel_stub_last_label = max(kernel_ir.blocks.keys()) + 1
-
-    # Add kernel stub last label to each parfor.loop_body label to prevent
-    # label conflicts.
-    loop_body = add_offset_to_labels(
-        reductionKernelVar.loop_body, kernel_stub_last_label
+    kernel_dispatcher: SPIRVKernelDispatcher = kernel(
+        kernel_template.py_func,
+        _parfor_body_args=ParforBodyArguments(
+            loop_body=reductionKernelVar.loop_body,
+            param_dict=reductionKernelVar.param_dict,
+            legal_loop_indices=reductionKernelVar.legal_loop_indices,
+        ),
     )
-    # new label for splitting sentinel block
-    new_label = max(loop_body.keys()) + 1
-
-    update_sentinel(kernel_ir, sentinel_name, loop_body, new_label)
 
-    old_alias = flags.noalias
-    if not has_aliases:
-        flags.noalias = True
+    kernel_param_types = reductionKernelVar.param_types
 
     kernel_sig = signature(types.none, *kernel_param_types)
 
-    # FIXME: A better design is required so that we do not have to create a
-    # queue every time.
-    ty_queue: DpctlSyclQueue = typemap[
-        reductionKernelVar.parfor_params[0]
-    ].queue
-    exec_queue = dpctl.get_device_cached_queue(ty_queue.sycl_device)
-
-    sycl_kernel = _compile_kernel_parfor(
-        exec_queue,
-        kernel_name,
-        kernel_ir,
-        kernel_param_types,
-        debug=flags.debuginfo,
+    kcres: _SPIRVKernelCompileResult = kernel_dispatcher.get_compile_result(
+        types.void(*kernel_param_types)  # kernel signature
     )
-
-    flags.noalias = old_alias
+    kernel_module: SPIRVKernelModule = kcres.kernel_device_ir_module
 
     return ParforKernel(
-        name=kernel_name,
-        kernel=sycl_kernel,
         signature=kernel_sig,
         kernel_args=reductionKernelVar.parfor_params,
         kernel_arg_types=reductionKernelVar.func_arg_types,
-        queue=exec_queue,
+        kernel_module=kernel_module,
     )
diff --git a/numba_dpex/core/pipelines/kernel_compiler.py b/numba_dpex/core/pipelines/kernel_compiler.py
index 7cb21d3817..c34a85b875 100644
--- a/numba_dpex/core/pipelines/kernel_compiler.py
+++ b/numba_dpex/core/pipelines/kernel_compiler.py
@@ -8,8 +8,12 @@
     IRLegalization,
     NoPythonSupportedFeatureValidation,
 )
+from numba.core.untyped_passes import InlineClosureLikes
 
 from numba_dpex.core.exceptions import UnsupportedCompilationModeError
+from numba_dpex.core.parfors.parfor_sentinel_replace_pass import (
+    ParforSentinelReplacePass,
+)
 from numba_dpex.core.passes.passes import (
     NoPythonBackend,
     QualNameDisambiguationLowering,
@@ -57,6 +61,8 @@ def define_nopython_pipeline(state, name="dpex_kernel_nopython"):
         pm = PassManager(name)
         untyped_passes = ndpb.define_untyped_pipeline(state)
         pm.passes.extend(untyped_passes.passes)
+        # TODO: create separate parfor kernel pass
+        pm.add_pass_after(ParforSentinelReplacePass, InlineClosureLikes)
 
         typed_passes = ndpb.define_typed_pipeline(state)
         pm.passes.extend(typed_passes.passes)
diff --git a/numba_dpex/core/runtime/_dpexrt_python.c b/numba_dpex/core/runtime/_dpexrt_python.c
index 10e441ed6c..5647a09241 100644
--- a/numba_dpex/core/runtime/_dpexrt_python.c
+++ b/numba_dpex/core/runtime/_dpexrt_python.c
@@ -45,12 +45,14 @@ static NRT_MemInfo *DPEXRT_MemInfo_fill(NRT_MemInfo *mi,
                                         bool value_is_float,
                                         int64_t value,
                                         const DPCTLSyclQueueRef qref);
-static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(PyObject *ndarrobj,
+static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(NRT_api_functions *nrt,
+                                                    PyObject *ndarrobj,
                                                     void *data,
                                                     npy_intp nitems,
                                                     npy_intp itemsize,
                                                     DPCTLSyclQueueRef qref);
-static NRT_MemInfo *DPEXRT_MemInfo_alloc(npy_intp size,
+static NRT_MemInfo *DPEXRT_MemInfo_alloc(NRT_api_functions *nrt,
+                                         npy_intp size,
                                          size_t usm_type,
                                          const DPCTLSyclQueueRef qref);
 static void usmndarray_meminfo_dtor(void *ptr, size_t size, void *info);
@@ -58,7 +60,8 @@ static PyObject *box_from_arystruct_parent(usmarystruct_t *arystruct,
                                            int ndim,
                                            PyArray_Descr *descr);
 
-static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj,
+static int DPEXRT_sycl_usm_ndarray_from_python(NRT_api_functions *nrt,
+                                               PyObject *obj,
                                                usmarystruct_t *arystruct);
 static PyObject *
 DPEXRT_sycl_usm_ndarray_to_python_acqref(usmarystruct_t *arystruct,
@@ -336,6 +339,11 @@ NRT_ExternalAllocator_new_for_usm(DPCTLSyclQueueRef qref, size_t usm_type)
 static void usmndarray_meminfo_dtor(void *ptr, size_t size, void *info)
 {
     MemInfoDtorInfo *mi_dtor_info = NULL;
+    // Warning: we are destructing sycl memory. MI destructor is called
+    // separately by numba.
+    DPEXRT_DEBUG(drt_debug_print("DPEXRT-DEBUG: Call to "
+                                 "usmndarray_meminfo_dtor at %s, line %d\n",
+                                 __FILE__, __LINE__));
 
     // Sanity-check to make sure the mi_dtor_info is an actual pointer.
     if (!(mi_dtor_info = (MemInfoDtorInfo *)info)) {
@@ -416,7 +424,8 @@ static MemInfoDtorInfo *MemInfoDtorInfo_new(NRT_MemInfo *mi, PyObject *owner)
  *                          of the dpnp.ndarray was allocated.
  * @return   {return}       A new NRT_MemInfo object
  */
-static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(PyObject *ndarrobj,
+static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(NRT_api_functions *nrt,
+                                                    PyObject *ndarrobj,
                                                     void *data,
                                                     npy_intp nitems,
                                                     npy_intp itemsize,
@@ -427,8 +436,9 @@ static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(PyObject *ndarrobj,
     MemInfoDtorInfo *midtor_info = NULL;
     DPCTLSyclContextRef cref = NULL;
 
-    // Allocate a new NRT_MemInfo object
-    if (!(mi = (NRT_MemInfo *)malloc(sizeof(NRT_MemInfo)))) {
+    // Allocate a new NRT_MemInfo object. By passing 0 we are just allocating
+    // MemInfo and not the `data` that the MemInfo object manages.
+    if (!(mi = (NRT_MemInfo *)nrt->allocate(0))) {
         DPEXRT_DEBUG(drt_debug_print(
             "DPEXRT-ERROR: Could not allocate a new NRT_MemInfo "
             "object  at %s, line %d\n",
@@ -505,7 +515,8 @@ static NRT_MemInfo *NRT_MemInfo_new_from_usmndarray(PyObject *ndarrobj,
  * @return   {return}     A new NRT_MemInfo object, NULL if no NRT_MemInfo
  *                        object could be created.
  */
-static NRT_MemInfo *DPEXRT_MemInfo_alloc(npy_intp size,
+static NRT_MemInfo *DPEXRT_MemInfo_alloc(NRT_api_functions *nrt,
+                                         npy_intp size,
                                          size_t usm_type,
                                          const DPCTLSyclQueueRef qref)
 {
@@ -517,7 +528,7 @@ static NRT_MemInfo *DPEXRT_MemInfo_alloc(npy_intp size,
         "DPEXRT-DEBUG: Inside DPEXRT_MemInfo_alloc  %s, line %d\n", __FILE__,
         __LINE__));
     // Allocate a new NRT_MemInfo object
-    if (!(mi = (NRT_MemInfo *)malloc(sizeof(NRT_MemInfo)))) {
+    if (!(mi = (NRT_MemInfo *)nrt->allocate(0))) {
         DPEXRT_DEBUG(drt_debug_print(
             "DPEXRT-ERROR: Could not allocate a new NRT_MemInfo object.\n"));
         goto error;
@@ -744,14 +755,18 @@ static struct PyUSMArrayObject *PyUSMNdArray_ARRAYOBJ(PyObject *obj)
         DPEXRT_DEBUG(
             drt_debug_print("DPEXRT-DEBUG: usm array was passed directly\n"));
         arrayobj = obj;
+        Py_INCREF(arrayobj);
     }
     else if (PyObject_HasAttrString(obj, "_array_obj")) {
+        // PyObject_GetAttrString gives reference
         arrayobj = PyObject_GetAttrString(obj, "_array_obj");
 
         if (!arrayobj)
             return NULL;
-        if (!PyObject_TypeCheck(arrayobj, &PyUSMArrayType))
+        if (!PyObject_TypeCheck(arrayobj, &PyUSMArrayType)) {
+            Py_DECREF(arrayobj);
             return NULL;
+        }
     }
 
     struct PyUSMArrayObject *pyusmarrayobj =
@@ -791,7 +806,8 @@ static npy_intp product_of_shape(npy_intp *shape, npy_intp ndim)
  *                          instance of a dpnp.ndarray
  * @return   {return}       Error code representing success (0) or failure (-1).
  */
-static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj,
+static int DPEXRT_sycl_usm_ndarray_from_python(NRT_api_functions *nrt,
+                                               PyObject *obj,
                                                usmarystruct_t *arystruct)
 {
     struct PyUSMArrayObject *arrayobj = NULL;
@@ -803,17 +819,13 @@ static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj,
     PyGILState_STATE gstate;
     npy_intp itemsize = 0;
 
-    // Increment the ref count on obj to prevent CPython from garbage
-    // collecting the array.
-    // TODO: add extra description why do we need this
-    Py_IncRef(obj);
-
     DPEXRT_DEBUG(drt_debug_print(
         "DPEXRT-DEBUG: In DPEXRT_sycl_usm_ndarray_from_python at %s, line %d\n",
         __FILE__, __LINE__));
 
     // Check if the PyObject obj has an _array_obj attribute that is of
     // dpctl.tensor.usm_ndarray type.
+    // arrayobj is a new reference, reference of obj is borrowed
     if (!(arrayobj = PyUSMNdArray_ARRAYOBJ(obj))) {
         DPEXRT_DEBUG(drt_debug_print(
             "DPEXRT-ERROR: PyUSMNdArray_ARRAYOBJ check failed at %s, line %d\n",
@@ -832,6 +844,7 @@ static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj,
     data = (void *)UsmNDArray_GetData(arrayobj);
     nitems = product_of_shape(shape, ndim);
     itemsize = (npy_intp)UsmNDArray_GetElementSize(arrayobj);
+
     if (!(qref = UsmNDArray_GetQueueRef(arrayobj))) {
         DPEXRT_DEBUG(drt_debug_print(
             "DPEXRT-ERROR: UsmNDArray_GetQueueRef returned NULL at "
@@ -841,7 +854,7 @@ static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj,
     }
 
     if (!(arystruct->meminfo = NRT_MemInfo_new_from_usmndarray(
-              obj, data, nitems, itemsize, qref)))
+              nrt, obj, data, nitems, itemsize, qref)))
     {
         DPEXRT_DEBUG(drt_debug_print(
             "DPEXRT-ERROR: NRT_MemInfo_new_from_usmndarray failed "
@@ -850,6 +863,9 @@ static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj,
         goto error;
     }
 
+    Py_XDECREF(arrayobj);
+    Py_IncRef(obj);
+
     arystruct->data = data;
     arystruct->sycl_queue = qref;
     arystruct->nitems = nitems;
@@ -906,7 +922,7 @@ static int DPEXRT_sycl_usm_ndarray_from_python(PyObject *obj,
         __FILE__, __LINE__));
     gstate = PyGILState_Ensure();
     // decref the python object
-    Py_DECREF(obj);
+    Py_XDECREF((PyObject *)arrayobj);
     // release the GIL
     PyGILState_Release(gstate);
 
@@ -938,6 +954,7 @@ static PyObject *box_from_arystruct_parent(usmarystruct_t *arystruct,
         drt_debug_print("DPEXRT-DEBUG: In box_from_arystruct_parent.\n"));
 
     if (!(arrayobj = PyUSMNdArray_ARRAYOBJ(arystruct->parent))) {
+        Py_XDECREF(arrayobj);
         DPEXRT_DEBUG(
             drt_debug_print("DPEXRT-DEBUG: Arrayobj cannot be boxed from "
                             "parent as parent pointer is NULL.\n"));
@@ -945,6 +962,7 @@ static PyObject *box_from_arystruct_parent(usmarystruct_t *arystruct,
     }
 
     if ((void *)UsmNDArray_GetData(arrayobj) != arystruct->data) {
+        Py_XDECREF(arrayobj);
         DPEXRT_DEBUG(drt_debug_print(
             "DPEXRT-DEBUG: Arrayobj cannot be boxed "
             "from parent as data pointer in the arystruct is not the same as "
@@ -952,12 +970,15 @@ static PyObject *box_from_arystruct_parent(usmarystruct_t *arystruct,
         return NULL;
     }
 
-    if (UsmNDArray_GetNDim(arrayobj) != ndim)
+    if (UsmNDArray_GetNDim(arrayobj) != ndim) {
+        Py_XDECREF(arrayobj);
         return NULL;
+    }
 
     p = arystruct->shape_and_strides;
     shape = UsmNDArray_GetShape(arrayobj);
     strides = UsmNDArray_GetStrides(arrayobj);
+    Py_XDECREF(arrayobj);
 
     // Ensure the shape of the array to be boxed matches the shape of the
     // original parent.
diff --git a/numba_dpex/core/runtime/context.py b/numba_dpex/core/runtime/context.py
index 80f0253101..4fe89abd7a 100644
--- a/numba_dpex/core/runtime/context.py
+++ b/numba_dpex/core/runtime/context.py
@@ -89,12 +89,14 @@ def meminfo_alloc_unchecked(self, builder, size, usm_type, queue_ref):
         mod = builder.module
         u64 = llvmir.IntType(64)
         fnty = llvmir.FunctionType(
-            cgutils.voidptr_t, [cgutils.intp_t, u64, cgutils.voidptr_t]
+            cgutils.voidptr_t,
+            [cgutils.voidptr_t, cgutils.intp_t, u64, cgutils.voidptr_t],
         )
         fn = cgutils.get_or_insert_function(mod, fnty, "DPEXRT_MemInfo_alloc")
         fn.return_value.add_attribute("noalias")
+        nrt_api = self._context.nrt.get_nrt_api(builder)
 
-        ret = builder.call(fn, [size, usm_type, queue_ref])
+        ret = builder.call(fn, [nrt_api, size, usm_type, queue_ref])
 
         return ret
 
@@ -168,13 +170,15 @@ def arraystruct_from_python(self, pyapi, obj, ptr):
 
         """
         fnty = llvmir.FunctionType(
-            llvmir.IntType(32), [pyapi.pyobj, pyapi.voidptr]
+            llvmir.IntType(32), [pyapi.voidptr, pyapi.pyobj, pyapi.voidptr]
         )
+        nrt_api = self._context.nrt.get_nrt_api(pyapi.builder)
         fn = pyapi._get_function(fnty, "DPEXRT_sycl_usm_ndarray_from_python")
         fn.args[0].add_attribute("nocapture")
         fn.args[1].add_attribute("nocapture")
+        fn.args[2].add_attribute("nocapture")
 
-        self.error = pyapi.builder.call(fn, (obj, ptr))
+        self.error = pyapi.builder.call(fn, (nrt_api, obj, ptr))
 
         return self.error
 
diff --git a/numba_dpex/core/types/dpnp_ndarray_type.py b/numba_dpex/core/types/dpnp_ndarray_type.py
index fb09598243..6d6622202c 100644
--- a/numba_dpex/core/types/dpnp_ndarray_type.py
+++ b/numba_dpex/core/types/dpnp_ndarray_type.py
@@ -2,14 +2,28 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
+from functools import partial
+
 import dpnp
 from numba.core import ir, types
 from numba.core.ir_utils import get_np_ufunc_typ, mk_unique_var
-from numba.core.pythonapi import NativeValue, PythonAPI, box, unbox
 
 from .usm_ndarray_type import USMNdArray
 
 
+def partialclass(cls, *args, **kwds):
+    """Creates fabric class of the original class with preset initialization
+    arguments."""
+    cls0 = partial(cls, *args, **kwds)
+    new_cls = type(
+        cls.__name__ + "Partial",
+        (cls,),
+        {"__new__": lambda cls, *args, **kwds: cls0(*args, **kwds)},
+    )
+
+    return new_cls
+
+
 class DpnpNdArray(USMNdArray):
     """
     The Numba type to represent an dpnp.ndarray. The type has the same
@@ -40,15 +54,22 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs):
         Returns: The DpnpNdArray class.
         """
         if method == "__call__":
-            if not all(
-                (
-                    isinstance(inp, DpnpNdArray)
-                    or isinstance(inp, types.abstract.Number)
-                )
-                for inp in inputs
-            ):
+            dpnp_type = None
+
+            for inp in inputs:
+                if isinstance(inp, DpnpNdArray):
+                    dpnp_type = inp
+                    continue
+                if isinstance(inp, types.abstract.Number):
+                    continue
+
                 return NotImplemented
-            return DpnpNdArray
+
+            assert dpnp_type is not None
+
+            return partialclass(
+                DpnpNdArray, queue=dpnp_type.queue, usm_type=dpnp_type.usm_type
+            )
         else:
             return
 
@@ -71,6 +92,8 @@ def __allocate__(
         lhs_typ,
         size_typ,
         out,
+        # dpex specific argument:
+        queue_ir_val=None,
     ):
         """Generates the Numba typed IR representing the allocation of a new
         DpnpNdArray using the dpnp.ndarray overload.
@@ -94,6 +117,10 @@ def __allocate__(
 
         Returns: The IR Value for the allocated array
         """
+        # TODO: it looks like it is being called only for parfor allocations,
+        # so can we rely on it? We can grab information from input arguments
+        # from rhs, but doc does not set any restriction on parfor use only.
+        assert queue_ir_val is not None
         g_np_var = ir.Var(scope, mk_unique_var("$np_g_var"), loc)
         if typemap:
             typemap[g_np_var.name] = types.misc.Module(dpnp)
@@ -132,11 +159,13 @@ def __allocate__(
         usm_typ_var = ir.Var(scope, mk_unique_var("$np_usm_type_var"), loc)
         # A default device string arg added as a placeholder
         device_typ_var = ir.Var(scope, mk_unique_var("$np_device_var"), loc)
+        queue_typ_var = ir.Var(scope, mk_unique_var("$np_queue_var"), loc)
 
         if typemap:
             typemap[layout_var.name] = types.literal(lhs_typ.layout)
             typemap[usm_typ_var.name] = types.literal(lhs_typ.usm_type)
-            typemap[device_typ_var.name] = types.literal(lhs_typ.device)
+            typemap[device_typ_var.name] = types.none
+            typemap[queue_typ_var.name] = lhs_typ.queue
 
         layout_var_assign = ir.Assign(
             ir.Const(lhs_typ.layout, loc), layout_var, loc
@@ -145,16 +174,29 @@ def __allocate__(
             ir.Const(lhs_typ.usm_type, loc), usm_typ_var, loc
         )
         device_typ_var_assign = ir.Assign(
-            ir.Const(lhs_typ.device, loc), device_typ_var, loc
+            ir.Const(None, loc), device_typ_var, loc
         )
+        queue_typ_var_assign = ir.Assign(queue_ir_val, queue_typ_var, loc)
 
         out.extend(
-            [layout_var_assign, usm_typ_var_assign, device_typ_var_assign]
+            [
+                layout_var_assign,
+                usm_typ_var_assign,
+                device_typ_var_assign,
+                queue_typ_var_assign,
+            ]
         )
 
         alloc_call = ir.Expr.call(
             attr_var,
-            [size_var, typ_var, layout_var, device_typ_var, usm_typ_var],
+            [
+                size_var,
+                typ_var,
+                layout_var,
+                device_typ_var,
+                usm_typ_var,
+                queue_typ_var,
+            ],
             (),
             loc,
         )
@@ -170,6 +212,7 @@ def __allocate__(
                         layout_var,
                         device_typ_var,
                         usm_typ_var,
+                        queue_typ_var,
                     ]
                 ],
                 {},
diff --git a/numba_dpex/core/types/kernel_api/atomic_ref.py b/numba_dpex/core/types/kernel_api/atomic_ref.py
index cdc2eb1890..6524833d7b 100644
--- a/numba_dpex/core/types/kernel_api/atomic_ref.py
+++ b/numba_dpex/core/types/kernel_api/atomic_ref.py
@@ -80,3 +80,13 @@ def cast_python_value(self, args):
         AtomicRefType throws a NotImplementedError.
         """
         raise NotImplementedError
+
+    @property
+    def mangling_args(self):
+        args = [
+            self.dtype,
+            self.memory_order,
+            self.memory_scope,
+            self.address_space,
+        ]
+        return self.__class__.__name__, args
diff --git a/numba_dpex/core/types/kernel_api/index_space_ids.py b/numba_dpex/core/types/kernel_api/index_space_ids.py
index 798aceeb2e..ea57fd05a4 100644
--- a/numba_dpex/core/types/kernel_api/index_space_ids.py
+++ b/numba_dpex/core/types/kernel_api/index_space_ids.py
@@ -31,6 +31,11 @@ def key(self):
     def cast_python_value(self, args):
         raise NotImplementedError
 
+    @property
+    def mangling_args(self):
+        args = [self.ndim]
+        return self.__class__.__name__, args
+
 
 class ItemType(types.Type):
     """Numba-dpex type corresponding to :class:`numba_dpex.kernel_api.Item`"""
@@ -53,6 +58,11 @@ def key(self):
         """Numba type specific overload"""
         return self._ndim
 
+    @property
+    def mangling_args(self):
+        args = [self.ndim]
+        return self.__class__.__name__, args
+
     def cast_python_value(self, args):
         raise NotImplementedError
 
@@ -78,5 +88,10 @@ def key(self):
         """Numba type specific overload"""
         return self._ndim
 
+    @property
+    def mangling_args(self):
+        args = [self.ndim]
+        return self.__class__.__name__, args
+
     def cast_python_value(self, args):
         raise NotImplementedError
diff --git a/numba_dpex/core/types/kernel_api/ranges.py b/numba_dpex/core/types/kernel_api/ranges.py
index af2f0125a7..9dc13e48a3 100644
--- a/numba_dpex/core/types/kernel_api/ranges.py
+++ b/numba_dpex/core/types/kernel_api/ranges.py
@@ -28,6 +28,11 @@ def ndim(self):
     def key(self):
         return self._ndim
 
+    @property
+    def mangling_args(self):
+        args = [self.ndim]
+        return self.__class__.__name__, args
+
 
 class NdRangeType(types.Type):
     """Numba-dpex type corresponding to
@@ -49,3 +54,8 @@ def ndim(self):
     @property
     def key(self):
         return self._ndim
+
+    @property
+    def mangling_args(self):
+        args = [self.ndim]
+        return self.__class__.__name__, args
diff --git a/numba_dpex/core/typing/dpnpdecl.py b/numba_dpex/core/typing/dpnpdecl.py
index 6856f8ea1f..720bada7d9 100644
--- a/numba_dpex/core/typing/dpnpdecl.py
+++ b/numba_dpex/core/typing/dpnpdecl.py
@@ -3,7 +3,6 @@
 # SPDX-License-Identifier: Apache-2.0
 
 import dpnp
-import numpy as np
 from numba.core import types
 from numba.core.typing.npydecl import (
     Numpy_rules_ufunc,
@@ -23,7 +22,7 @@ def _install_operations(cls: NumpyRulesArrayOperator):
     for op, ufunc_name in cls._op_map.items():
         infer_global(op)(
             type(
-                "NumpyRulesArrayOperator_" + ufunc_name,
+                "DpnpRulesArrayOperator_" + ufunc_name,
                 (cls,),
                 dict(key=op),
             )
@@ -33,17 +32,7 @@ def _install_operations(cls: NumpyRulesArrayOperator):
 class DpnpRulesArrayOperator(NumpyRulesArrayOperator):
     @property
     def ufunc(self):
-        try:
-            op = getattr(dpnp, self._op_map[self.key])
-            npop = getattr(np, self._op_map[self.key])
-            op.nin = npop.nin
-            op.nout = npop.nout
-            op.nargs = npop.nargs
-            op.types = npop.types
-            op.is_dpnp_ufunc = True
-            return op
-        except:
-            pass
+        return getattr(dpnp, self._op_map[self.key])
 
     @classmethod
     def install_operations(cls):
diff --git a/numba_dpex/dpnp_iface/dpnp_ufunc_db.py b/numba_dpex/dpnp_iface/dpnp_ufunc_db.py
index def71adb17..f07bf47b03 100644
--- a/numba_dpex/dpnp_iface/dpnp_ufunc_db.py
+++ b/numba_dpex/dpnp_iface/dpnp_ufunc_db.py
@@ -59,11 +59,6 @@ def _fill_ufunc_db_with_dpnp_ufuncs(ufunc_db):
     for ufuncop in dpnpdecl.supported_ufuncs:
         if ufuncop == "erf":
             op = getattr(dpnp, "erf")
-            op.nin = 1
-            op.nout = 1
-            op.nargs = 2
-            op.types = ["f->f", "d->d"]
-            op.is_dpnp_ufunc = True
 
             _unary_d_d = types.float64(types.float64)
             _unary_f_f = types.float32(types.float32)
@@ -72,20 +67,16 @@ def _fill_ufunc_db_with_dpnp_ufuncs(ufunc_db):
                 "d->d": mathimpl.lower_ocl_impl[("erf", (_unary_d_d))],
             }
         else:
-            op = getattr(dpnp, ufuncop)
+            dpnpop = getattr(dpnp, ufuncop)
             npop = getattr(np, ufuncop)
-            op.nin = npop.nin
-            op.nout = npop.nout
-            op.nargs = npop.nargs
-            op.types = npop.types
-            op.is_dpnp_ufunc = True
+
             cp = copy.copy(_ufunc_db[npop])
-            ufunc_db.update({op: cp})
-            for key in list(ufunc_db[op].keys()):
+            ufunc_db.update({dpnpop: cp})
+            for key in list(ufunc_db[dpnpop].keys()):
                 if (
                     "FF->" in key
                     or "DD->" in key
                     or "F->" in key
                     or "D->" in key
                 ):
-                    ufunc_db[op].pop(key)
+                    ufunc_db[dpnpop].pop(key)
diff --git a/numba_dpex/examples/kernel/matmul.py b/numba_dpex/examples/kernel/matmul.py
index 335d319bda..34742f877e 100644
--- a/numba_dpex/examples/kernel/matmul.py
+++ b/numba_dpex/examples/kernel/matmul.py
@@ -83,7 +83,7 @@ def _arange_reshaped(shape, dtype):
 X = dpt.asarray(X)
 Y = dpt.asarray(Y)
 device = X.device.sycl_device
-result = dpt.zeros((5, 5), dtype, device=device)
+result = dpt.zeros((5, 5), dtype=dtype, device=device)
 X_slm = kapi.LocalAccessor(shape=work_group_size, dtype=dtype)
 Y_slm = kapi.LocalAccessor(shape=work_group_size, dtype=dtype)
 
diff --git a/numba_dpex/kernel_api_impl/spirv/dispatcher.py b/numba_dpex/kernel_api_impl/spirv/dispatcher.py
index 99e6de8b6c..92224d9c28 100644
--- a/numba_dpex/kernel_api_impl/spirv/dispatcher.py
+++ b/numba_dpex/kernel_api_impl/spirv/dispatcher.py
@@ -26,7 +26,7 @@
 from numba.core.types import void
 from numba.core.typing.typeof import Purpose, typeof
 
-from numba_dpex import config
+from numba_dpex.core import config
 from numba_dpex.core.descriptor import dpex_kernel_target
 from numba_dpex.core.exceptions import (
     ExecutionQueueInferenceError,
diff --git a/numba_dpex/kernel_api_impl/spirv/spirv_generator.py b/numba_dpex/kernel_api_impl/spirv/spirv_generator.py
index c731a9c75d..192311f9e0 100644
--- a/numba_dpex/kernel_api_impl/spirv/spirv_generator.py
+++ b/numba_dpex/kernel_api_impl/spirv/spirv_generator.py
@@ -124,6 +124,7 @@ def finalize(self):
             "--spirv-ext=+SPV_EXT_shader_atomic_float_add",
             "--spirv-ext=+SPV_EXT_shader_atomic_float_min_max",
             "--spirv-ext=+SPV_INTEL_arbitrary_precision_integers",
+            "--spirv-ext=+SPV_INTEL_variable_length_array",
         ]
         for key in list(self.context.extra_compile_options.keys()):
             if key == LLVM_SPIRV_ARGS:
diff --git a/numba_dpex/kernel_api_impl/spirv/target.py b/numba_dpex/kernel_api_impl/spirv/target.py
index e1e4bcbc65..379a76b92d 100644
--- a/numba_dpex/kernel_api_impl/spirv/target.py
+++ b/numba_dpex/kernel_api_impl/spirv/target.py
@@ -11,7 +11,7 @@
 import dpnp
 from llvmlite import binding as ll
 from llvmlite import ir as llvmir
-from numba.core import cgutils, funcdesc
+from numba.core import cgutils
 from numba.core import types as nb_types
 from numba.core import typing
 from numba.core.base import BaseContext
@@ -21,8 +21,10 @@
 from numba.core.typing import cmathdecl, enumdecl
 
 from numba_dpex.core.datamodel.models import _init_kernel_data_model_manager
+from numba_dpex.core.debuginfo import DIBuilder as DpexDIbuilder
 from numba_dpex.core.types import IntEnumLiteral
 from numba_dpex.core.typing import dpnpdecl
+from numba_dpex.core.utils import itanium_mangler
 from numba_dpex.kernel_api.flag_enum import FlagEnum
 from numba_dpex.kernel_api.memory_enums import AddressSpace as address_space
 from numba_dpex.kernel_api_impl.spirv import printimpl
@@ -134,6 +136,7 @@ class SPIRVTargetContext(BaseContext):
 
     implement_powi_as_math_call = True
     allow_dynamic_globals = True
+    DIBuilder = DpexDIbuilder
 
     def __init__(self, typingctx, target=SPIRV_TARGET_NAME):
         super().__init__(typingctx, target)
@@ -188,7 +191,7 @@ def _generate_spir_kernel_wrapper(self, func, argtypes):
         wrapper_module = self._internal_codegen.create_empty_spirv_module(
             "dpex.kernel.wrapper"
         )
-        wrappername = func.name.replace("dpex_fn", "dpex_kernel")
+        wrappername = func.name + ("dpex_kernel")
         argtys = list(arginfo.argument_types)
         fnty = llvmir.FunctionType(
             llvmir.IntType(32),
@@ -319,12 +322,9 @@ def target_data(self):
 
     def mangler(self, name, types, *, abi_tags=(), uid=None):
         """
-        Generates a name for a function by appending \"dpex_fn\" to the
-        name of the function before calling Numba's default function name
-        mangler."""
-        return funcdesc.default_mangler(
-            name + "dpex_fn", types, abi_tags=abi_tags, uid=uid
-        )
+        Generates a mangled function name using numba_dpex's itanium mangler.
+        """
+        return itanium_mangler.mangle(name, types, abi_tags=abi_tags, uid=uid)
 
     def prepare_spir_kernel(self, func, argtypes):
         """Generates a wrapper function with \"spir_kernel\" calling conv that
diff --git a/numba_dpex/numba_patches/patch_is_ufunc.py b/numba_dpex/numba_patches/patch_is_ufunc.py
deleted file mode 100644
index dee4a3f375..0000000000
--- a/numba_dpex/numba_patches/patch_is_ufunc.py
+++ /dev/null
@@ -1,23 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-
-def patch():
-    """Patches the numba.np.ufunc.array_exprs._is_ufunc function to make it
-    possible to support dpnp universal functions (ufuncs).
-
-    The extra condition is the check for the "is_dpnp_ufunc" attribute to
-    identify a non-NumPy ufunc.
-    """
-    import numpy
-    from numba.np.ufunc.dufunc import DUFunc
-
-    def _is_ufunc(func):
-        return isinstance(func, (numpy.ufunc, DUFunc)) or hasattr(
-            func, "is_dpnp_ufunc"
-        )
-
-    from numba.np.ufunc import array_exprs
-
-    array_exprs._is_ufunc = _is_ufunc
diff --git a/numba_dpex/numba_patches/patch_ufuncs.py b/numba_dpex/numba_patches/patch_ufuncs.py
new file mode 100644
index 0000000000..e3e316dba2
--- /dev/null
+++ b/numba_dpex/numba_patches/patch_ufuncs.py
@@ -0,0 +1,94 @@
+# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import logging
+
+import dpnp
+import numpy as np
+
+from numba_dpex.core.typing import dpnpdecl
+
+
+def patch():
+    patch_is_ufunc()
+    patch_ufuncs()
+
+
+def patch_is_ufunc():
+    """Patches the numba.np.ufunc.array_exprs._is_ufunc function to make it
+    possible to support dpnp universal functions (ufuncs).
+
+    The extra condition is the check for the "is_dpnp_ufunc" attribute to
+    identify a non-NumPy ufunc.
+    """
+    import numpy
+    from numba.np.ufunc.dufunc import DUFunc
+
+    def _is_ufunc(func):
+        return isinstance(func, (numpy.ufunc, DUFunc)) or hasattr(
+            func, "is_dpnp_ufunc"
+        )
+
+    from numba.np.ufunc import array_exprs
+
+    array_exprs._is_ufunc = _is_ufunc
+
+
+def patch_ufuncs():
+    """Patches dpnp user functions to make them compatible with numpy, so we
+    can reuse numba's implementation.
+
+    It adds "nin", "nout", "nargs" and "is_dpnp_ufunc" attributes to ufuncs.
+    """
+    failed_dpnpop_types_lst = []
+
+    op = getattr(dpnp, "erf")
+    op.nin = 1
+    op.nout = 1
+    op.nargs = 2
+    op.types = ["f->f", "d->d"]
+    op.is_dpnp_ufunc = True
+
+    for ufuncop in dpnpdecl.supported_ufuncs:
+        if ufuncop == "erf":
+            continue
+
+        dpnpop = getattr(dpnp, ufuncop)
+        npop = getattr(np, ufuncop)
+
+        if not hasattr(dpnpop, "nin"):
+            dpnpop.nin = npop.nin
+        if not hasattr(dpnpop, "nout"):
+            dpnpop.nout = npop.nout
+        if not hasattr(dpnpop, "nargs"):
+            dpnpop.nargs = dpnpop.nin + dpnpop.nout
+
+        # Check if the dpnp operation has a `types` attribute and if an
+        # AttributeError gets raised then "monkey patch" the attribute from
+        # numpy. If the attribute lookup raised a ValueError, it indicates
+        # that dpnp could not be resolve the supported types for the
+        # operation. Dpnp will fail to resolve the `types` if no SYCL
+        # devices are available on the system. For such a scenario, we log
+        # dpnp operations for which the ValueError happened and print them
+        # as a user-level warning. It is done this way so that the failure
+        # to load the dpnpdecl registry due to the ValueError does not
+        # impede a user from importing numba-dpex.
+        try:
+            dpnpop.types
+        except ValueError:
+            failed_dpnpop_types_lst.append(ufuncop)
+        except AttributeError:
+            dpnpop.types = npop.types
+
+        dpnpop.is_dpnp_ufunc = True
+
+    if len(failed_dpnpop_types_lst) > 0:
+        try:
+            getattr(dpnp, failed_dpnpop_types_lst[0]).types
+        except ValueError:
+            ops = " ".join(failed_dpnpop_types_lst)
+            logging.exception(
+                "The types attribute for the following dpnp ops could not be "
+                f"determined: {ops}"
+            )
diff --git a/pyproject.toml b/pyproject.toml
index e44e25b640..aa7cc94eb3 100644
--- a/pyproject.toml
+++ b/pyproject.toml
@@ -1,15 +1,21 @@
 [build-system]
 build-backend = "setuptools.build_meta"
 requires = [
-  # TODO: keep in sync with [project.dependencies] and /conda-recipe/meta.yaml
+  # TODO: keep in sync with [project.dependencies]
+  "wheel>=0.43",
+  "build>=1.1",
   "setuptools>=63.0.0",
   "scikit-build>=0.17.0",
   "ninja>=1.11.1; platform_system!='Windows'",
   "cmake>=3.29.0",
-  "numba>=0.59.0",
+  # We need dpctl for UsmNdArray integration for dpcpp code
   "dpctl>=0.16.1",
-  "numpy>=1.24.0",
-  "wheel",
+  # We need numba for runtime cpp headers
+  "numba>=0.59.0,<0.60.0a0",
+  "llvmlite>=0.42.0",
+  # Do we need dpnp at build time?
+  "dpnp >=0.14",
+  "numpy >=1.24",
   # WARNING: check with doc how to upgrade
   "versioneer[toml]==0.29"
 ]
@@ -32,14 +38,16 @@ classifiers = [
 dependencies = [
   # TODO: keep in sync with [build-system.requires] and /conda-recipe/meta.yaml
   # This restrictions are for dependabot, actual restrictions are set with
-  # conda.
-  # TODO: populate it during build process
+  # conda. TODO: populate it during build process
+  # TODO: do we have to set sycl runtime dependencies here
+  # "dpcpp-cpp-rt>=0.59.0",
+  # "intel-cmplr-lib-rt>=0.59.0"
   "numba>=0.59.0",
   "llvmlite>=0.42.0",
   "dpctl>=0.16.1",
   "dpnp>=0.14.0",
   "numpy>=1.24.0",
-  "dpcpp_llvm_spirv>=2024.0"
+  "dpcpp_llvm_spirv>=2024.1"
 ]
 description = "An extension for Numba to add data-parallel offload capability"
 dynamic = ["version"]
diff --git a/setup.py b/setup.py
index a5613c1884..c545d3d1a4 100644
--- a/setup.py
+++ b/setup.py
@@ -3,10 +3,11 @@
 # SPDX-License-Identifier: Apache-2.0
 
 
+import os
 import re
 
 import versioneer
-from setuptools import find_packages
+from setuptools import find_namespace_packages, find_packages
 from skbuild import setup
 
 """Top level setup.py file. Uses scikit-build.
@@ -50,6 +51,16 @@ def to_cmake_format(version: str):
 # Get the project version
 __version__ = versioneer.get_version()
 
+# Keep same style as versioneer in case ve would want to switch back
+with open(os.path.join("numba_dpex", "_version.py"), "w") as fh:
+    print(
+        "# SPDX-FileCopyrightText: 2024 Intel Corporation\n#\n"
+        "# SPDX-License-Identifier: Apache-2.0\n\n\ndef get_versions():\n"
+        '    return {"version": "' + __version__ + '"}',
+        file=fh,
+    )
+
+
 # Main setup
 setup(
     version=__version__,
@@ -58,7 +69,8 @@ def to_cmake_format(version: str):
     # Must be passed vis setup.py:
     # https://github.com/scikit-build/scikit-build/issues/864
     # TODO: switch to pyproject toml after switching to scikit-build-core
-    packages=find_packages("."),
+    packages=find_packages(".")
+    + find_namespace_packages(".", include=["numba_dpex.examples.*"]),
     # Needs for examples.
     # TODO: change to false once move examples out of package.
     include_package_data=True,


Test patch:
diff --git a/numba_dpex/tests/_helper.py b/numba_dpex/tests/_helper.py
index c790cb5b73..41b31518be 100644
--- a/numba_dpex/tests/_helper.py
+++ b/numba_dpex/tests/_helper.py
@@ -122,7 +122,7 @@ def is_windows():
 ]
 
 skip_no_gdb = pytest.mark.skipif(
-    config.TESTING_SKIP_NO_DEBUGGING and not shutil.which("gdb-oneapi"),
+    config.TESTING_SKIP_NO_DEBUGGING or not shutil.which("gdb-oneapi"),
     reason="Intel® Distribution for GDB* is not available",
 )
 
diff --git a/numba_dpex/tests/codegen/test_intenum_literal_codegen.py b/numba_dpex/tests/codegen/test_intenum_literal_codegen.py
index 5d875f37af..ef19eb2bed 100644
--- a/numba_dpex/tests/codegen/test_intenum_literal_codegen.py
+++ b/numba_dpex/tests/codegen/test_intenum_literal_codegen.py
@@ -47,7 +47,7 @@ def pass_flags_to_func(a):
 
     pattern = re.compile(
         r"call spir_func i32 @\_Z.*bitwise\_or"
-        r"\_flags.*\(i64\* nonnull %.*, i64 1, i64 2\)"
+        r"\_flags.*\(i64\*\s(\w+)?\s*%.*, i64 1, i64 2\)"
     )
 
     assert re.search(pattern, llvm_ir_mod) is not None
diff --git a/numba_dpex/tests/debugging/gdb.py b/numba_dpex/tests/debugging/gdb.py
index f47fad117b..b743507c15 100644
--- a/numba_dpex/tests/debugging/gdb.py
+++ b/numba_dpex/tests/debugging/gdb.py
@@ -33,8 +33,8 @@
     r"layout=[A-Z],\s+"
     r"address_space=[0-4],\s+"
     r"usm_type=[a-z]+,\s+"
-    r"device=[a-z:0-9]+,\s+"
-    r"sycl_queue=[A-Za-z:0-9 ]+\)"
+    r"device=[a-z\_:0-9]+,\s+"
+    r"sycl_queue=[A-Za-z:0-9\s\_:]+\)"
 )
 
 
@@ -52,7 +52,7 @@ def spawn(self):
         env["NUMBA_DEBUGINFO"] = "1"
 
         self.child = pexpect.spawn(
-            "gdb-oneapi -q python", env=env, encoding="utf-8"
+            "gdb-oneapi -q python", env=env, encoding="utf-8", timeout=60
         )
         if config.TESTING_LOG_DEBUGGING:
             self.child.logfile = sys.stdout
@@ -109,7 +109,7 @@ def expect_eol(self):
         self.child.expect(r"[^\n]*\n")
 
     def expect_hit_breakpoint(self, expected_location=None):
-        expect = r"Thread [0-9A-Za-z \"]+ hit Breakpoint [0-9\.]+"
+        expect = r"Breakpoint [0-9\.]+"
         if expected_location is not None:
             # function name + args could be long, so we have to assume that
             # the message may be splitted in multiple lines. It potentially can
diff --git a/numba_dpex/tests/debugging/test_backtraces.py b/numba_dpex/tests/debugging/test_backtraces.py
index 117eccddc0..7a4e736d9a 100644
--- a/numba_dpex/tests/debugging/test_backtraces.py
+++ b/numba_dpex/tests/debugging/test_backtraces.py
@@ -25,7 +25,5 @@ def test_backtrace(app):
 
     app.backtrace()
 
-    app.expect(
-        r"#0.*__main__::func_sum.* at simple_dpex_func.py:12", with_eol=True
-    )
-    app.expect(r"#1.*__main__::kernel_sum", with_eol=True)
+    app.expect(r"#0.*func_sum.* at simple_dpex_func.py:12", with_eol=True)
+    app.expect(r"#1.*kernel_sum", with_eol=True)
diff --git a/numba_dpex/tests/debugging/test_breakpoints.py b/numba_dpex/tests/debugging/test_breakpoints.py
index ef178e5415..f76411b100 100644
--- a/numba_dpex/tests/debugging/test_breakpoints.py
+++ b/numba_dpex/tests/debugging/test_breakpoints.py
@@ -56,14 +56,9 @@ def test_device_func_breakpoint(
     SAT-4449
     """
 
-    if api == "numba-ndpx-kernel" and breakpoint != "side-by-side.py:15":
-        pytest.skip(
-            "Breakpoint by function name not working for numba-dpex."
-        )  # TODO: https://github.com/IntelPython/numba-dpex/issues/1242
-
     app.breakpoint(breakpoint, condition=condition)
     app.run(f"side-by-side.py --api={api}")
-    app.expect_hit_breakpoint("side-by-side.py:15")
+    app.expect_hit_breakpoint(expected_location="side-by-side.py:15")
     if exp_var is not None:
         app.print(exp_var, expected=exp_val)
 
diff --git a/numba_dpex/tests/debugging/test_info.py b/numba_dpex/tests/debugging/test_info.py
index 421aba8e16..693f487fcb 100644
--- a/numba_dpex/tests/debugging/test_info.py
+++ b/numba_dpex/tests/debugging/test_info.py
@@ -78,7 +78,7 @@ def test_info_args(
 ):
     app.breakpoint(breakpoint)
     app.run(script)
-    app.expect_hit_breakpoint(breakpoint)
+    app.expect_hit_breakpoint(expected_location=breakpoint)
     app.expect(expected_line, with_eol=True)
 
     if kind == "info":
@@ -99,12 +99,12 @@ def test_info_args(
 def test_info_functions(app):
     app.breakpoint("simple_sum.py:12")
     app.run("simple_sum.py")
-    app.expect_hit_breakpoint("simple_sum.py:12")
+    app.expect_hit_breakpoint(expected_location="simple_sum.py:12")
     app.expect(r"12\s+i = item.get_id\(0\)", with_eol=True)
 
     app.info_functions("data_parallel_sum")
 
-    app.child.expect(r"11:\s+[a-z 0-9\*]+__main__::data_parallel_sum")
+    app.child.expect(r"11:\s+[a-z 0-9\*]+data_parallel_sum")
 
 
 @pytest.mark.parametrize(
@@ -119,7 +119,7 @@ def test_print_array_element(app, api):
 
     app.breakpoint("side-by-side-2.py:17", condition="param_a == 5")
     app.run(f"side-by-side-2.py --api={api}")
-    app.expect_hit_breakpoint("side-by-side-2.py:17")
+    app.expect_hit_breakpoint(expected_location="side-by-side-2.py:17")
 
     # We can access only c_array, not python style array
     app.print("b.data[5]", 5)
@@ -142,7 +142,7 @@ def test_print_array_element(app, api):
 def test_assignment_to_variable(app, api, assign):
     app.breakpoint("side-by-side-2.py:17", condition="param_a == 5")
     app.run(f"side-by-side-2.py --api={api}")
-    app.expect_hit_breakpoint("side-by-side-2.py:17")
+    app.expect_hit_breakpoint(expected_location="side-by-side-2.py:17")
 
     app.print("param_a", expected=5)
     if assign == "print":
diff --git a/numba_dpex/tests/debugging/test_stepping.py b/numba_dpex/tests/debugging/test_stepping.py
index 58464350b4..3e802a1db0 100644
--- a/numba_dpex/tests/debugging/test_stepping.py
+++ b/numba_dpex/tests/debugging/test_stepping.py
@@ -21,7 +21,7 @@
 def test_next(app: gdb):
     app.breakpoint("simple_dpex_func.py:18")
     app.run("simple_dpex_func.py")
-    app.expect_hit_breakpoint("simple_dpex_func.py:18")
+    app.expect_hit_breakpoint(expected_location="simple_dpex_func.py:18")
     app.expect(r"18\s+i = item.get_id\(0\)", with_eol=True)
     app.set_scheduler_lock()
     app.next()
@@ -31,20 +31,20 @@ def test_next(app: gdb):
     )
     # checking that we did not step in
     app.next()
-    app.expect(r"in _ZN8__main__21kernel_sum_", with_eol=True)
+    app.expect(r"in _ZN8__main__14kernel_sum_", with_eol=True)
 
 
 def test_step(app: gdb):
     app.breakpoint("simple_dpex_func.py:19")
     app.run("simple_dpex_func.py")
-    app.expect_hit_breakpoint("simple_dpex_func.py:19")
+    app.expect_hit_breakpoint(expected_location="simple_dpex_func.py:19")
     app.expect(
         r"19\s+c_in_kernel\[i\] = func_sum\(a_in_kernel\[i\], b_in_kernel\[i\]\)",
         with_eol=True,
     )
     app.set_scheduler_lock()
     app.step()
-    app.expect(r"__main__::func_sum.* at simple_dpex_func.py:12", with_eol=True)
+    app.expect(r"func_sum.* at simple_dpex_func.py:12", with_eol=True)
     app.expect(r"12\s+result = a_in_func \+ b_in_func", with_eol=True)
     app.step()
     app.expect(
@@ -57,7 +57,7 @@ def test_step(app: gdb):
 def test_stepi(app: gdb, func: str):
     app.breakpoint("simple_dpex_func.py:19")
     app.run("simple_dpex_func.py")
-    app.expect_hit_breakpoint("simple_dpex_func.py:19")
+    app.expect_hit_breakpoint(expected_location="simple_dpex_func.py:19")
     app.expect(
         r"19\s+c_in_kernel\[i\] = func_sum\(a_in_kernel\[i\], b_in_kernel\[i\]\)",
         with_eol=True,
diff --git a/numba_dpex/tests/test_debuginfo.py b/numba_dpex/tests/test_debuginfo.py
index e95e9e6d1f..bdde118d3b 100644
--- a/numba_dpex/tests/test_debuginfo.py
+++ b/numba_dpex/tests/test_debuginfo.py
@@ -126,8 +126,8 @@ def data_parallel_sum(item, a, b, c):
         c[i] = func_sum(a[i], b[i])
 
     ir_tags = [
-        r'\!DISubprogram\(name: ".*func_sum\$?\d*"',
-        r'\!DISubprogram\(name: ".*data_parallel_sum\$?\d*"',
+        r'\!DISubprogram\(name: ".*func_sum*"',
+        r'\!DISubprogram\(name: ".*data_parallel_sum*"',
     ]
 
     sig = (itemty, f32arrty, f32arrty, f32arrty)
@@ -156,8 +156,8 @@ def data_parallel_sum(item, a, b, c):
         c[i] = func_sum(a[i], b[i])
 
     ir_tags = [
-        r'\!DISubprogram\(name: ".*func_sum\$?\d*"',
-        r'\!DISubprogram\(name: ".*data_parallel_sum\$\d*"',
+        r'\!DISubprogram\(name: ".*func_sum*"',
+        r'\!DISubprogram\(name: ".*data_parallel_sum"',
     ]
 
     sig = (itemty, f32arrty, f32arrty, f32arrty)
@@ -176,6 +176,8 @@ def data_parallel_sum(item, a, b, c):
 
 
 def test_debuginfo_DISubprogram_linkageName():
+    """Tests to check that the linkagename tag is not set by numba-dpex."""
+
     def foo(item, a, b):
         i = item.get_id(0)
         b[i] = a[i]
@@ -190,7 +192,9 @@ def foo(item, a, b):
     kernel_ir = kcres.library.get_llvm_str()
 
     for tag in ir_tags:
-        assert make_check(kernel_ir, tag)
+        # Ensure that linkagename (DW_AT_linkagename) tag is not present for
+        # the DISubprogram attribute.
+        assert not make_check(kernel_ir, tag)
 
 
 def test_debuginfo_DICompileUnit_language_and_producer():