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

Issue:
Update docstring to avoid mysterious tm issue
- [X] Have you provided a meaningful PR description?
Minor change to docstring


Discussion / hints:

Fix patch:
diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
index 4733d9f49c..6377801126 100644
--- a/.github/CODEOWNERS
+++ b/.github/CODEOWNERS
@@ -1,5 +1,1 @@
-*.yml  @diptorupd
-conda-recipe/*  @diptorupd
-numba_dpex/dpnp_iface @diptorupd @chudur-budur
-numba_dpex/core @diptorupd
-numba_dpex/core/passes/parfor_lowering_pass.py @diptorupd
+conda-recipe/*  @oleksandr-pavlyk
diff --git a/.github/workflows/black.yml b/.github/workflows/black.yml
index 3c8860f6e5..bc010cce60 100644
--- a/.github/workflows/black.yml
+++ b/.github/workflows/black.yml
@@ -27,4 +27,4 @@ jobs:
       - uses: psf/black@stable
         with:
           args: ". --check"
-          version: "23.1.0"
+          version: "24.1.1"
diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml
index 2f8eada12b..d692143127 100644
--- a/.github/workflows/conda-package.yml
+++ b/.github/workflows/conda-package.yml
@@ -89,7 +89,7 @@ jobs:
           path: ${{ env.WHEELS_OUTPUT_FOLDER }}${{ env.WHEELS_NAME }}-*.whl
 
   test:
-    name: Test ['${{ matrix.os }}', python='${{ matrix.python }}']
+    name: Test ['${{ matrix.os }}', '${{ matrix.scope }}', python='${{ matrix.python }}']
 
     needs: build
 
@@ -106,10 +106,10 @@ jobs:
         os: [ubuntu-20.04, ubuntu-latest, windows-latest]
         experimental: [false]
         use_mlir: [false]
-        run_gdb: [true, false]
+        scope: ['tests', 'examples', 'gdb']
         exclude:
           - os: windows-latest
-            run_gdb: true
+            scope: 'gdb'
 
     continue-on-error: ${{ matrix.experimental || matrix.use_mlir }}
 
@@ -214,25 +214,25 @@ jobs:
         run: python -c "import numba_dpex"
 
       - name: Run tests
-        if: ${{!matrix.run_gdb}}
+        if: ${{ matrix.scope == 'tests' }}
         env:
           NUMBA_DPEX_USE_MLIR: ${{ matrix.use_mlir && '1' || '0'  }}
         run: |
           pytest -q -ra --disable-warnings --pyargs ${{ env.MODULE_NAME }} -vv
 
       - name: Run examples
-        if: ${{!matrix.run_gdb}}
+        if: ${{ matrix.scope == 'examples' }}
         shell: bash -l {0}
         run: |
           cd ${{ env.EXAMPLES_PATH }}
-          for script in $(find . \( -not -name "_*" -not -name "side-by-side*" -not -name "vectorize.py" -not -name "scan.py" -and -name "*.py" \))
+          for script in $(find . \( -not -name "_*" -not -name "side-by-side*" -not -name "scan.py" -and -name "*.py" \))
           do
             echo "Executing ${script}"
             python ${script} || exit 1
           done
 
       - name: Run gdb tests
-        if: ${{matrix.run_gdb}}
+        if: ${{ matrix.scope == 'gdb' }}
         env:
           GDB_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/fcea1bcd-6a63-4849-b304-906ff71dc2c0/l_dpcpp_dbg_p_2023.2.0.49333_offline.sh
           GDB_INSTALLER: l_dpcpp_dbg_p_2023.2.0.49333_offline.sh
diff --git a/.github/workflows/openssf-scorecard.yml b/.github/workflows/openssf-scorecard.yml
new file mode 100644
index 0000000000..545b289ea4
--- /dev/null
+++ b/.github/workflows/openssf-scorecard.yml
@@ -0,0 +1,73 @@
+# This workflow uses actions that are not certified by GitHub. They are provided
+# by a third-party and are governed by separate terms of service, privacy
+# policy, and support documentation.
+
+name: Scorecard supply-chain security
+on:
+  # For Branch-Protection check. Only the default branch is supported. See
+  # https://github.com/ossf/scorecard/blob/main/docs/checks.md#branch-protection
+  branch_protection_rule:
+  # To guarantee Maintained check is occasionally updated. See
+  # https://github.com/ossf/scorecard/blob/main/docs/checks.md#maintained
+  schedule:
+    - cron: '28 2 * * 1'
+    - cron: '28 2 * * 4'
+  push:
+    branches: [ "main" ]
+
+# Declare default permissions as read only.
+permissions: read-all
+
+jobs:
+  analysis:
+    name: Scorecard analysis
+    runs-on: ubuntu-latest
+    permissions:
+      # Needed to upload the results to code-scanning dashboard.
+      security-events: write
+      # Needed to publish results and get a badge (see publish_results below).
+      id-token: write
+      # Uncomment the permissions below if installing in a private repository.
+      # contents: read
+      # actions: read
+
+    steps:
+      - name: "Checkout code"
+        uses: actions/checkout@b4ffde65f46336ab88eb53be808477a3936bae11 # v4.1.1
+        with:
+          persist-credentials: false
+
+      - name: "Run analysis"
+        uses: ossf/scorecard-action@0864cf19026789058feabb7e87baa5f140aac736 # v2.3.1
+        with:
+          results_file: results.sarif
+          results_format: sarif
+          # (Optional) "write" PAT token. Uncomment the `repo_token` line below if:
+          # - you want to enable the Branch-Protection check on a *public* repository, or
+          # - you are installing Scorecard on a *private* repository
+          # To create the PAT, follow the steps in https://github.com/ossf/scorecard-action#authentication-with-pat.
+          # repo_token: ${{ secrets.SCORECARD_TOKEN }}
+
+          # Public repositories:
+          #   - Publish results to OpenSSF REST API for easy access by consumers
+          #   - Allows the repository to include the Scorecard badge.
+          #   - See https://github.com/ossf/scorecard-action#publishing-results.
+          # For private repositories:
+          #   - `publish_results` will always be set to `false`, regardless
+          #     of the value entered here.
+          publish_results: true
+
+      # Upload the results as artifacts (optional). Commenting out will disable uploads of run results in SARIF
+      # format to the repository Actions tab.
+      - name: "Upload artifact"
+        uses: actions/upload-artifact@3cea5372237819ed00197afe530f5a7ea3e805c8 # v3.1.0
+        with:
+          name: SARIF file
+          path: results.sarif
+          retention-days: 14
+
+      # Upload the results to GitHub's code scanning dashboard.
+      - name: "Upload to code-scanning"
+        uses: github/codeql-action/upload-sarif@e8893c57a1f3a2b659b6b55564fdfdbbd2982911 # v3.24.0
+        with:
+          sarif_file: results.sarif
diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
index 2411834df1..1199f4ffa5 100644
--- a/.pre-commit-config.yaml
+++ b/.pre-commit-config.yaml
@@ -13,7 +13,7 @@ repos:
     -   id: end-of-file-fixer
     -   id: trailing-whitespace
 -   repo: https://github.com/psf/black
-    rev: 23.1.0
+    rev: 24.1.1
     hooks:
     -   id: black
         exclude: "versioneer.py|numba_dpex/_version.py"
@@ -52,7 +52,7 @@ repos:
     - id: pylint
       name: pylint
       entry: pylint
-      files: ^numba_dpex/experimental|^numba_dpex/core/utils/kernel_launcher.py
+      files: ^numba_dpex/kernel_api|^numba_dpex/experimental|^numba_dpex/core/utils/kernel_launcher.py
       language: system
       types: [python]
       require_serial: true
diff --git a/README.md b/README.md
index 6fa00b9f3c..73c195f133 100644
--- a/README.md
+++ b/README.md
@@ -3,6 +3,7 @@
 [![pre-commit](https://img.shields.io/badge/pre--commit-enabled-brightgreen?logo=pre-commit&logoColor=white)](https://github.com/pre-commit/pre-commit)
 [![Join the chat at https://matrix.to/#/#Data-Parallel-Python_community:gitter.im](https://badges.gitter.im/Join%20Chat.svg)](https://app.gitter.im/#/room/#Data-Parallel-Python_community:gitter.im)
 [![Coverity Scan Build Status](https://scan.coverity.com/projects/29068/badge.svg)](https://scan.coverity.com/projects/intelpython-numba-dpex)
+[![OpenSSF Scorecard](https://api.securityscorecards.dev/projects/github.com/IntelPython/numba-dpex/badge)](https://securityscorecards.dev/viewer/?uri=github.com/IntelPython/numba-dpex)
 <img align="left" src="https://spec.oneapi.io/oneapi-logo-white-scaled.jpg" alt="oneAPI logo" width="75"/>
 <br/>
 <br/>
diff --git a/conda-recipe/meta.yaml b/conda-recipe/meta.yaml
index a6f3dd98b9..bb566350eb 100644
--- a/conda-recipe/meta.yaml
+++ b/conda-recipe/meta.yaml
@@ -1,4 +1,4 @@
-{% set required_compiler_version = "2023.2" %}
+{% set required_compiler_version = "2024.0" %}
 {% set excluded_compiler_version1 = "2024.0.1" %}
 {% set excluded_compiler_version2 = "2024.0.2" %}
 
@@ -28,18 +28,18 @@ requirements:
         - scikit-build >=0.15*
         - ninja # [not win]
         - cmake >=3.26*
-        - numba >=0.57*
-        - dpctl >=0.14*
-        - dpnp >=0.11*
+        - numba >=0.59*
+        - dpctl >=0.16*
+        - dpnp >=0.14*
         - dpcpp-llvm-spirv
         - wheel
     run:
         - {{ pin_compatible('dpcpp-cpp-rt', min_pin='x.x', max_pin='x') }}
         - python
-        - numba >=0.57*
-        - dpctl >=0.14*
+        - numba >=0.59*
+        - dpctl >=0.16*
         - dpcpp-llvm-spirv
-        - dpnp >=0.11*
+        - dpnp >=0.14*
         - packaging
 
 test:
diff --git a/docs/.gitignore b/docs/.gitignore
index 5c490088af..c221eeaac0 100644
--- a/docs/.gitignore
+++ b/docs/.gitignore
@@ -1,2 +1,3 @@
 apidoc
 sources/_build
+source/autoapi
diff --git a/docs/Makefile b/docs/Makefile
index 1ed4090504..e912e4ca8e 100644
--- a/docs/Makefile
+++ b/docs/Makefile
@@ -7,6 +7,7 @@ SPHINXOPTS    ?=
 SPHINXBUILD   ?= sphinx-build
 SOURCEDIR     = ./source
 BUILDDIR      = _build
+AUTOAPIDIR    = source/autoapi
 
 # Put it first so that "make" without argument is like "make help".
 help:
@@ -20,4 +21,4 @@ help:
 	@$(SPHINXBUILD) -M $@ "$(SOURCEDIR)" "$(BUILDDIR)" $(SPHINXOPTS) $(O)
 
 clean:
-	rm -rf "$(BUILDDIR)"
+	rm -rf "$(BUILDDIR)" "$(AUTOAPIDIR)"
diff --git a/docs/_static/css/custom.css b/docs/_static/css/custom.css
new file mode 100644
index 0000000000..889b915047
--- /dev/null
+++ b/docs/_static/css/custom.css
@@ -0,0 +1,15 @@
+span.summarylabel {
+    background-color: var(--color-foreground-secondary);
+    color: var(--color-background-secondary);
+    font-size: 70%;
+    padding-left: 2px;
+    padding-right: 2px;
+    border-radius: 3px;
+    vertical-align: 15%;
+    padding-bottom: 2px;
+    filter: opacity(40%);
+}
+
+table.summarytable {
+    width: 100%;
+}
diff --git a/docs/_templates/autoapi/index.rst b/docs/_templates/autoapi/index.rst
new file mode 100644
index 0000000000..1af3e006b2
--- /dev/null
+++ b/docs/_templates/autoapi/index.rst
@@ -0,0 +1,17 @@
+API Reference
+=============
+
+This page contains auto-generated API reference documentation [#f1]_.
+
+.. toctree::
+   :maxdepth: 2
+
+   numba_dpex/kernel_api/index
+
+   {% for page in pages %}
+   {% if page.top_level_object and page.display %}
+   {{ page.include_path }}
+   {% endif %}
+   {% endfor %}
+
+.. [#f1] Created with `sphinx-autoapi <https://github.com/readthedocs/sphinx-autoapi>`_
diff --git a/docs/_templates/autoapi/macros.rst b/docs/_templates/autoapi/macros.rst
new file mode 100644
index 0000000000..9703b9d2db
--- /dev/null
+++ b/docs/_templates/autoapi/macros.rst
@@ -0,0 +1,41 @@
+{% macro _render_item_name(obj, sig=False) -%}
+:py:obj:`{{ obj.name }} <{{ obj.id }}>`
+     {%- if sig -%}
+       \ (
+       {%- for arg in obj.obj.args -%}
+          {%- if arg[0] %}{{ arg[0]|replace('*', '\*') }}{% endif -%}{{  arg[1] -}}
+          {%- if not loop.last  %}, {% endif -%}
+       {%- endfor -%}
+       ){%- endif -%}
+{%- endmacro %}
+
+{% macro _item(obj, sig=False, label='') %}
+   * - {{ _render_item_name(obj, sig) }}
+     - {% if label %}:summarylabel:`{{ label }}` {% endif %}{% if obj.summary %}{{ obj.summary }}{% else %}\-{% endif +%}
+{% endmacro %}
+
+{% macro auto_summary(objs, title='') -%}
+.. list-table:: {{ title }}
+   :header-rows: 0
+   :widths: auto
+   :class: summarytable
+
+  {% for obj in objs -%}
+    {%- set sig = (obj.type in ['method', 'function'] and not 'property' in obj.properties) -%}
+
+    {%- if 'property' in obj.properties -%}
+      {%- set label = 'prop' -%}
+    {%- elif 'classmethod' in obj.properties -%}
+      {%- set label = 'class' -%}
+    {%- elif 'abstractmethod' in obj.properties -%}
+      {%- set label = 'abc' -%}
+    {%- elif 'staticmethod' in obj.properties -%}
+      {%- set label = 'static' -%}
+    {%- else -%}
+      {%- set label = '' -%}
+    {%- endif -%}
+
+    {{- _item(obj, sig=sig, label=label) -}}
+  {%- endfor -%}
+
+{% endmacro %}
diff --git a/docs/_templates/autoapi/python/attribute.rst b/docs/_templates/autoapi/python/attribute.rst
new file mode 100644
index 0000000000..ebaba555ad
--- /dev/null
+++ b/docs/_templates/autoapi/python/attribute.rst
@@ -0,0 +1,1 @@
+{% extends "python/data.rst" %}
diff --git a/docs/_templates/autoapi/python/class.rst b/docs/_templates/autoapi/python/class.rst
new file mode 100644
index 0000000000..ab998efa5e
--- /dev/null
+++ b/docs/_templates/autoapi/python/class.rst
@@ -0,0 +1,69 @@
+{% import 'macros.rst' as macros %}
+
+{% if obj.display %}
+.. py:{{ obj.type }}:: {{ obj.short_name }}{% if obj.args %}({{ obj.args }}){% endif %}
+{% for (args, return_annotation) in obj.overloads %}
+   {{ " " * (obj.type | length) }}   {{ obj.short_name }}{% if args %}({{ args }}){% endif %}
+{% endfor %}
+
+
+   {% if obj.bases %}
+   {% if "show-inheritance" in autoapi_options %}
+   Bases: {% for base in obj.bases %}{{ base|link_objs }}{% if not loop.last %}, {% endif %}{% endfor %}
+   {% endif %}
+
+
+   {% if "show-inheritance-diagram" in autoapi_options and obj.bases != ["object"] %}
+   .. autoapi-inheritance-diagram:: {{ obj.obj["full_name"] }}
+      :parts: 1
+      {% if "private-members" in autoapi_options %}
+      :private-bases:
+      {% endif %}
+
+   {% endif %}
+   {% endif %}
+   {% if obj.docstring %}
+   {{ obj.docstring|indent(3) }}
+   {% endif %}
+   {% if "inherited-members" in autoapi_options %}
+   {% set visible_classes = obj.classes|selectattr("display")|list %}
+   {% else %}
+   {% set visible_classes = obj.classes|rejectattr("inherited")|selectattr("display")|list %}
+   {% endif %}
+   {% for klass in visible_classes %}
+   {{ klass.render()|indent(3) }}
+   {% endfor %}
+   {% if "inherited-members" in autoapi_options %}
+   {% set visible_attributes = obj.attributes|selectattr("display")|list %}
+   {% else %}
+   {% set visible_attributes = obj.attributes|rejectattr("inherited")|selectattr("display")|list %}
+   {% endif %}
+   {% if "inherited-members" in autoapi_options %}
+   {% set visible_methods = obj.methods|selectattr("display")|list %}
+   {% else %}
+   {% set visible_methods = obj.methods|rejectattr("inherited")|selectattr("display")|list %}
+   {% endif %}
+
+   {% if visible_methods or visible_attributes %}
+   .. rubric:: Overview
+
+   {% set summary_methods = visible_methods|rejectattr("properties", "contains", "property")|list %}
+   {% set summary_attributes = visible_attributes + visible_methods|selectattr("properties", "contains", "property")|list %}
+   {% if summary_attributes %}
+   {{ macros.auto_summary(summary_attributes, title="Attributes")|indent(3) }}
+   {% endif %}
+
+   {% if summary_methods %}
+   {{ macros.auto_summary(summary_methods, title="Methods")|indent(3) }}
+   {% endif %}
+
+   .. rubric:: Members
+
+   {% for attribute in visible_attributes %}
+   {{ attribute.render()|indent(3) }}
+   {% endfor %}
+   {% for method in visible_methods %}
+   {{ method.render()|indent(3) }}
+   {% endfor %}
+   {% endif %}
+{% endif %}
diff --git a/docs/_templates/autoapi/python/data.rst b/docs/_templates/autoapi/python/data.rst
new file mode 100644
index 0000000000..3d12b2d0c7
--- /dev/null
+++ b/docs/_templates/autoapi/python/data.rst
@@ -0,0 +1,37 @@
+{% if obj.display %}
+.. py:{{ obj.type }}:: {{ obj.name }}
+   {%- if obj.annotation is not none %}
+
+   :type: {%- if obj.annotation %} {{ obj.annotation }}{%- endif %}
+
+   {%- endif %}
+
+   {%- if obj.value is not none %}
+
+   :value: {% if obj.value is string and obj.value.splitlines()|count > 1 -%}
+                Multiline-String
+
+    .. raw:: html
+
+        <details><summary>Show Value</summary>
+
+    .. code-block:: python
+
+        """{{ obj.value|indent(width=8,blank=true) }}"""
+
+    .. raw:: html
+
+        </details>
+
+            {%- else -%}
+              {%- if obj.value is string -%}
+                {{ "%r" % obj.value|string|truncate(100) }}
+              {%- else -%}
+                {{ obj.value|string|truncate(100) }}
+              {%- endif -%}
+            {%- endif %}
+   {%- endif %}
+
+
+   {{ obj.docstring|indent(3) }}
+{% endif %}
diff --git a/docs/_templates/autoapi/python/exception.rst b/docs/_templates/autoapi/python/exception.rst
new file mode 100644
index 0000000000..92f3d38fd5
--- /dev/null
+++ b/docs/_templates/autoapi/python/exception.rst
@@ -0,0 +1,1 @@
+{% extends "python/class.rst" %}
diff --git a/docs/_templates/autoapi/python/function.rst b/docs/_templates/autoapi/python/function.rst
new file mode 100644
index 0000000000..778ba9a883
--- /dev/null
+++ b/docs/_templates/autoapi/python/function.rst
@@ -0,0 +1,15 @@
+{% if obj.display %}
+.. py:function:: {{ obj.short_name }}({{ obj.args }}){% if obj.return_annotation is not none %} -> {{ obj.return_annotation }}{% endif %}
+
+{% for (args, return_annotation) in obj.overloads %}
+                 {{ obj.short_name }}({{ args }}){% if return_annotation is not none %} -> {{ return_annotation }}{% endif %}
+
+{% endfor %}
+   {% for property in obj.properties %}
+   :{{ property }}:
+   {% endfor %}
+
+   {% if obj.docstring %}
+   {{ obj.docstring|indent(3) }}
+   {% endif %}
+{% endif %}
diff --git a/docs/_templates/autoapi/python/method.rst b/docs/_templates/autoapi/python/method.rst
new file mode 100644
index 0000000000..f50a1bd6ee
--- /dev/null
+++ b/docs/_templates/autoapi/python/method.rst
@@ -0,0 +1,19 @@
+{%- if obj.display %}
+.. py:method:: {{ obj.short_name }}({{ obj.args }}){% if obj.return_annotation is not none %} -> {{ obj.return_annotation }}{% endif %}
+
+{% for (args, return_annotation) in obj.overloads %}
+               {{ obj.short_name }}({{ args }}){% if return_annotation is not none %} -> {{ return_annotation }}{% endif %}
+
+{% endfor %}
+   {% if obj.properties %}
+   {% for property in obj.properties %}
+   :{{ property }}:
+   {% endfor %}
+
+   {% else %}
+
+   {% endif %}
+   {% if obj.docstring %}
+   {{ obj.docstring|indent(3) }}
+   {% endif %}
+{% endif %}
diff --git a/docs/_templates/autoapi/python/module.rst b/docs/_templates/autoapi/python/module.rst
new file mode 100644
index 0000000000..32143c3957
--- /dev/null
+++ b/docs/_templates/autoapi/python/module.rst
@@ -0,0 +1,112 @@
+{% import 'macros.rst' as macros %}
+
+{% if not obj.display %}
+:orphan:
+
+{% endif %}
+{{ obj.name }}
+{{ "=" * obj.name|length }}
+
+.. py:module:: {{ obj.name }}
+
+{% if obj.docstring %}
+.. autoapi-nested-parse::
+
+   {{ obj.docstring|indent(3) }}
+
+{% endif %}
+
+{% block subpackages %}
+{% set visible_subpackages = obj.subpackages|selectattr("display")|list %}
+{% if visible_subpackages %}
+Subpackages
+-----------
+.. toctree::
+   :titlesonly:
+   :maxdepth: 3
+
+{% for subpackage in visible_subpackages %}
+   {{ subpackage.short_name }}/index.rst
+{% endfor %}
+
+
+{% endif %}
+{% endblock %}
+{% block submodules %}
+{% set visible_submodules = obj.submodules|selectattr("display")|list %}
+{% if visible_submodules %}
+Submodules
+----------
+.. toctree::
+   :titlesonly:
+   :maxdepth: 1
+
+{% for submodule in visible_submodules %}
+   {{ submodule.short_name }}/index.rst
+{% endfor %}
+
+
+{% endif %}
+{% endblock %}
+{% block content %}
+{% if obj.all is not none %}
+{% set visible_children = obj.children|selectattr("display")|selectattr("short_name", "in", obj.all)|list %}
+{% elif obj.type is equalto("package") %}
+{% set visible_children = obj.children|selectattr("display")|list %}
+{% else %}
+{% set visible_children = obj.children|selectattr("display")|rejectattr("imported")|list %}
+{% endif %}
+{% if visible_children %}
+Overview
+--------
+
+{% set visible_classes = visible_children|selectattr("type", "equalto", "class")|list %}
+{% set visible_functions = visible_children|selectattr("type", "equalto", "function")|list %}
+{% set visible_attributes = visible_children|selectattr("type", "equalto", "data")|list %}
+{% if "show-module-summary" in autoapi_options and (visible_classes or visible_functions) %}
+{% block classes scoped %}
+{% if visible_classes %}
+{{ macros.auto_summary(visible_classes, title="Classes") }}
+{% endif %}
+{% endblock %}
+
+{% block functions scoped %}
+{% if visible_functions %}
+{{ macros.auto_summary(visible_functions, title="Function") }}
+{% endif %}
+{% endblock %}
+
+{% block attributes scoped %}
+{% if visible_attributes %}
+{{ macros.auto_summary(visible_attributes, title="Attributes") }}
+{% endif %}
+{% endblock %}
+{% endif %}
+
+{% if visible_classes %}
+Classes
+-------
+{% for obj_item in visible_classes %}
+{{ obj_item.render()|indent(0) }}
+{% endfor %}
+{% endif %}
+
+{% if visible_functions %}
+Functions
+---------
+{% for obj_item in visible_functions %}
+{{ obj_item.render()|indent(0) }}
+{% endfor %}
+{% endif %}
+
+{% if visible_attributes %}
+Attributes
+----------
+{% for obj_item in visible_attributes %}
+{{ obj_item.render()|indent(0) }}
+{% endfor %}
+{% endif %}
+
+
+{% endif %}
+{% endblock %}
diff --git a/docs/_templates/autoapi/python/package.rst b/docs/_templates/autoapi/python/package.rst
new file mode 100644
index 0000000000..fb9a64965e
--- /dev/null
+++ b/docs/_templates/autoapi/python/package.rst
@@ -0,0 +1,1 @@
+{% extends "python/module.rst" %}
diff --git a/docs/_templates/autoapi/python/property.rst b/docs/_templates/autoapi/python/property.rst
new file mode 100644
index 0000000000..70af24236f
--- /dev/null
+++ b/docs/_templates/autoapi/python/property.rst
@@ -0,0 +1,15 @@
+{%- if obj.display %}
+.. py:property:: {{ obj.short_name }}
+   {% if obj.annotation %}
+   :type: {{ obj.annotation }}
+   {% endif %}
+   {% if obj.properties %}
+   {% for property in obj.properties %}
+   :{{ property }}:
+   {% endfor %}
+   {% endif %}
+
+   {% if obj.docstring %}
+   {{ obj.docstring|indent(3) }}
+   {% endif %}
+{% endif %}
diff --git a/docs/backups/docker.rst b/docs/backups/docker.rst
deleted file mode 100644
index b37f4428be..0000000000
--- a/docs/backups/docker.rst
+++ /dev/null
@@ -1,146 +0,0 @@
-.. include:: ./../ext_links.txt
-
-Docker Support
-==============
-
-Numba dpex now delivers docker support. Dockerfile is capable of building
-numba-dpex as well as direct dependencies for it: dpctl and dpnp. There are
-several prebuilt images available: for trying numba_dpex and for building
-numba-dpex.
-
-Building
---------
-
-Numba dpex ships with multistage Dockerfile, which means there are different
-`targets
-<https://docs.docker.com/build/building/multi-stage/#stop-at-a-specific-build-stage>`_
-available for build. The most useful ones:
-
-- ``runtime``
-- ``runtime-gpu``
-- ``numba-dpex-builder-runtime``
-- ``numba-dpex-builder-runtime-gpu``
-
-To build docker image
-
-.. code-block:: bash
-
-    docker build --target runtime -t numba-dpex:runtime ./
-
-
-To run docker image
-
-.. code-block:: bash
-
-    docker run -it --rm numba-dpex:runtime
-
-.. note::
-
-    If you are building docker image with gpu support it will calls github api
-    to get latest versions of intel gpu drivers. You may face Github API call
-    limits. To avoid this, you can pass your github credentials to increase this
-    limit. You can do it by providing `build args
-    <https://docs.docker.com/engine/reference/commandline/build/#build-arg>`_
-    ``GITHUB_USER`` and ``GITHUB_PASSWORD``. You can use `access token
-    <https://docs.github.com/en/authentication/keeping-your-account-and-data-secure/creating-a-personal-access-token>`_
-    instead of the password:
-
-.. note::
-
-    In case you are building docker image behind firewall and your internet
-    access requires proxy, you can provide proxy `build args
-    <https://docs.docker.com/engine/reference/commandline/build/#build-arg>`_
-    ``http_proxy`` and ``https_proxy``. Please note, that these args must be
-    lowercase.
-
-Dockerfile supports different python versions. To select the one you want,
-simply set ``PYTHON_VERSION`` build arg. By default docker image is based on
-official python image based on slim debian, so the requested python version must
-be from the available python docker images. In case you want to build on images
-on custom image you have to pass ``BASE_IMAGE`` environment variable. Be aware
-that Dockerfile is based on debian so any base image should be debian based,
-like debian or ubuntu.
-
-Build arguments that could be useful:
-
-- ``PYTHON_VERSION``
-- ``CR_TAG``
-- ``IGC_TAG``
-- ``CM_TAG``
-- ``L0_TAG``
-- ``ONEAPI_VERSION``
-- ``DPCTL_GIT_BRANCH``
-- ``DPCTL_GIT_URL``
-- ``DPNP_GIT_BRANCH``
-- ``DPNP_GIT_URL``
-- ``NUMBA_DPEX_GIT_BRANCH``
-- ``NUMBA_DPEX_GIT_URL``
-- ``CMAKE_VERSION``
-- ``CMAKE_VERSION_BUILD``
-- ``INTEL_NUMPY_VERSION``
-- ``INTEL_NUMBA_VERSION``
-- ``CYTHON_VERSION``
-- ``SCIKIT_BUILD_VERSION``
-- ``http_proxy``
-- ``https_proxy``
-- ``GITHUB_USER``
-- ``GITHUB_PASSWORD``
-- ``BASE_IMAGE``
-
-Refer to Dockerfile to see all available
-
-Running prebuilt images
------------------------
-
-An easy way you can try ``numba_dpex`` is by using prebuilt images. There are
-several prebuilt images available:
-
-- ``runtime`` package that provides runtime experience
-.. code-block:: text
-
-    ghcr.io/intelpython/numba-dpex/runtime:<numba_dpex_version>-py<python_version>[-gpu]
-
-- ``builder`` package that provides building experience
-.. code-block:: text
-
-    ghcr.io/intelpython/numba-dpex/builder:<numba_dpex_version>-py<python_version>[-gpu]
-
-- you can also see ``stages`` package, but it is used mostly for building
-  stages.
-You can use them to build your own docker that is built on top of one of them.
-
-To try out numba dpex simply run:
-
-.. code-block:: bash
-
-    docker run --rm -it ghcr.io/intelpython/numba-dpex/runtime:0.20.0-py3.10 bash
-
-Within the container you can check available devices by running
-
-.. code-block:: bash
-
-    sycl-ls
-
-or
-
-.. code-block:: bash
-
-    python -c "import dpctl; dpctl.lsplatform();".
-
-.. note::
-
-    If you want to enable GPU you need to pass it within container and use
-    ``*-gpu`` tag.
-
-    For passing GPU into container on linux use arguments ``--device=/dev/dri``.
-    However if you are using WSL you need to pass ``--device=/dev/dxg -v
-    /usr/lib/wsl:/usr/lib/wsl`` instead.
-
-So, for example, if you want to run numba dpex container with GPU support on
-WSL:
-
-.. code-block:: bash
-
-    docker run --rm -it \
-    --device=/dev/dxg -v /usr/lib/wsl:/usr/lib/wsl \
-    ghcr.io/intelpython/numba-dpex/runtime:0.20.0-py3.10-gpu
diff --git a/docs/backups/sdc-sphinx-theme/globaltoc.html b/docs/backups/sdc-sphinx-theme/globaltoc.html
deleted file mode 100644
index d96aefdd9a..0000000000
--- a/docs/backups/sdc-sphinx-theme/globaltoc.html
+++ /dev/null
@@ -1,2 +0,0 @@
-<h3>Table of Contents</h3>
-{{ toctree(maxdepth=-1, titles_only=true) }}
diff --git a/docs/backups/sdc-sphinx-theme/layout.html b/docs/backups/sdc-sphinx-theme/layout.html
deleted file mode 100644
index 9a6badaf37..0000000000
--- a/docs/backups/sdc-sphinx-theme/layout.html
+++ /dev/null
@@ -1,99 +0,0 @@
-{% extends "basic/layout.html" %}
-
-{%- block scripts %}
-    {{ super() }}
-    <script type="text/javascript" src="{{ pathto('_static/sidebar.js', 1) }}"></script>
-    {% if not embedded %}<script type="text/javascript" src="{{ pathto('_static/copybutton.js', 1) }}"></script>{% endif %}
-{%- endblock %}
-
-{# Add the google webfonts needed for the logo #}
-{% block extrahead %}
-    <link href='https://fonts.googleapis.com/css?family=Source+Sans+Pro:200,600' rel='stylesheet' type='text/css'/>
-{% endblock %}
-
-{% block header %}
-<div class="topbar">
-  <a class="brand" title="{{ _('Intel Python projects') }}" href="{{ pathto('other_intel_python_projects') }}"></a>
-  <a class="brand_sdc" title="{{ _('Documentation Home') }}" href="{{ pathto(master_doc) }}"></a>
-
-  <ul>
-    <li><a class="exampleslink" title="Examples" href="{{ pathto('examples') }}"></a></li>
-    <li><a class="issueslink" title="Issues" href="https://community.intel.com/t5/Intel-Distribution-for-Python/bd-p/distribution-python"></a></li>
-    <li><a class="emaillink" title="Email" href="mailto:scripting@intel.com"></a></li>
-    <li><a class="homelink" title="GitHub" href="https://github.com/IntelPython/numba-dpex"></a></li>
-    <li>
-      {% block sidebarsearch %}
-      {% include "searchbox.html" %}
-      {% endblock %}
-    </li>
-  </ul>
-</div>
-{% endblock %}
-
-{% block relbar1 %}
-<div class="related">
-    <h3>{{ _('Navigation') }}</h3>
-    <ul>
-      {%- if next %}
-      <li class="right">
-	<a href="{{ next.link|e }}" title="{{ next.title|striptags|e }}">
-	  next {{ "&raquo;"|safe }}
-	</a>
-      </li>
-      {%- endif %}
-      {%- if prev %}
-      <li class="right">
-	<a href="{{ prev.link|e }}" title="{{ prev.title|striptags|e }}">
-	  {{ "&laquo;"|safe }} previous
-	</a>
-	{% if next %}{{ reldelim2 }}{% endif %}
-      </li>
-      {%- endif %}
-      {%- block rootrellink %}
-      <li>
-	<a href="{{ pathto(master_doc) }}">{{ shorttitle|e }}</a>
-	{{ reldelim1 }}
-      </li>
-      {%- endblock %}
-      {%- for parent in parents %}
-      <li><a href="{{ parent.link|e }}" {% if loop.last %}{{ accesskey("U") }}{% endif %}>{{ parent.title }}</a>{{ reldelim1 }}</li>
-      {%- endfor %}
-      {# Don't put the title in the relbar for the first (index) page. #}
-      {% if prev %}<li>{{ title }}</li>{% endif %}
-      {%- block relbaritems %} {% endblock %}
-    </ul>
-</div>
-{% endblock %}
-
-{# Silence the bottom relbar. #}
-{% block relbar2 %}{% endblock %}
-
-
-{%- block footer %}
-<footer class="footer">
-  <p class="pull-right">
-    {%- if edit_on_github %}
-    <a href="{{ edit_on_github }}">{{ edit_on_github_page_message }}</a> &nbsp;
-    {%- endif %}
-    {%- if show_source and has_source and sourcename %}
-    <a href="{{ pathto('_sources/' + sourcename, true)|e }}"
-       rel="nofollow">{{ _('Page Source') }}</a>
-    {%- endif %} &nbsp;
-    <a href="#">Back to Top</a></p>
-  <p>
-    {%- if show_copyright %}
-    {%- if hasdoc('copyright') %}
-    {% trans path=pathto('copyright'), copyright=copyright|e %}&copy; <a href="{{ path }}">Copyright</a> {{ copyright }}.{% endtrans %}<br/>
-    {%- else %}
-    {% trans copyright=copyright|e %}&copy; Copyright {{ copyright }}.{% endtrans %}<br/>
-    {%- endif %}
-    {%- endif %}
-    {%- if show_sphinx %}
-    {% trans sphinx_version=sphinx_version|e %}Created using <a href="http://www.sphinx-doc.org/en/stable/">Sphinx</a> {{ sphinx_version }}.{% endtrans %} &nbsp;
-    {%- endif %}
-    {%- if last_updated %}
-    {% trans last_updated=last_updated|e %}Last built {{ last_updated }}.{% endtrans %} <br/>
-    {%- endif %}
-  </p>
-</footer>
-{%- endblock %}
diff --git a/docs/backups/sdc-sphinx-theme/localtoc.html b/docs/backups/sdc-sphinx-theme/localtoc.html
deleted file mode 100644
index 00752f4e12..0000000000
--- a/docs/backups/sdc-sphinx-theme/localtoc.html
+++ /dev/null
@@ -1,2 +0,0 @@
-<h3>Page Contents</h3>
-{{ toc }}
diff --git a/docs/backups/sdc-sphinx-theme/searchbox.html b/docs/backups/sdc-sphinx-theme/searchbox.html
deleted file mode 100644
index 0d17cc1491..0000000000
--- a/docs/backups/sdc-sphinx-theme/searchbox.html
+++ /dev/null
@@ -1,7 +0,0 @@
-{%- if pagename != "search" %}
-<form action="{{ pathto('search') }}" method="get">
-  <input type="text" name="q" placeholder="Search" />
-  <input type="hidden" name="check_keywords" value="yes" />
-  <input type="hidden" name="area" value="default" />
-</form>
-{%- endif %}
diff --git a/docs/backups/sdc-sphinx-theme/static/copybutton.js b/docs/backups/sdc-sphinx-theme/static/copybutton.js
deleted file mode 100644
index 8f12e23d6c..0000000000
--- a/docs/backups/sdc-sphinx-theme/static/copybutton.js
+++ /dev/null
@@ -1,62 +0,0 @@
-$(document).ready(function() {
-    /* Add a [>>>] button on the top-right corner of code samples to hide
-     * the >>> and ... prompts and the output and thus make the code
-     * copyable. */
-    var div = $('.highlight-python .highlight,' +
-                '.highlight-python3 .highlight,' +
-                '.highlight-default .highlight')
-    var pre = div.find('pre');
-
-    // get the styles from the current theme
-    pre.parent().parent().css('position', 'relative');
-    var hide_text = 'Hide the prompts and output';
-    var show_text = 'Show the prompts and output';
-    var border_width = pre.css('border-top-width');
-    var border_style = pre.css('border-top-style');
-    var border_color = pre.css('border-top-color');
-    var button_styles = {
-        'cursor':'pointer', 'position': 'absolute', 'top': '0', 'right': '0',
-        'border-color': border_color, 'border-style': border_style,
-        'border-width': border_width, 'color': border_color, 'text-size': '75%',
-        'font-family': 'monospace', 'padding-left': '0.2em', 'padding-right': '0.2em',
-        'border-radius': '0 3px 0 0'
-    }
-
-    // create and add the button to all the code blocks that contain >>>
-    div.each(function(index) {
-        var jthis = $(this);
-        if (jthis.find('.gp').length > 0) {
-            var button = $('<span class="copybutton">&gt;&gt;&gt;</span>');
-            button.css(button_styles)
-            button.attr('title', hide_text);
-            button.data('hidden', 'false');
-            jthis.prepend(button);
-        }
-        // tracebacks (.gt) contain bare text elements that need to be
-        // wrapped in a span to work with .nextUntil() (see later)
-        jthis.find('pre:has(.gt)').contents().filter(function() {
-            return ((this.nodeType == 3) && (this.data.trim().length > 0));
-        }).wrap('<span>');
-    });
-
-    // define the behavior of the button when it's clicked
-    $('.copybutton').click(function(e){
-        e.preventDefault();
-        var button = $(this);
-        if (button.data('hidden') === 'false') {
-            // hide the code output
-            button.parent().find('.go, .gp, .gt').hide();
-            button.next('pre').find('.gt').nextUntil('.gp, .go').css('visibility', 'hidden');
-            button.css('text-decoration', 'line-through');
-            button.attr('title', show_text);
-            button.data('hidden', 'true');
-        } else {
-            // show the code output
-            button.parent().find('.go, .gp, .gt').show();
-            button.next('pre').find('.gt').nextUntil('.gp, .go').css('visibility', 'visible');
-            button.css('text-decoration', 'none');
-            button.attr('title', hide_text);
-            button.data('hidden', 'false');
-        }
-    });
-});
diff --git a/docs/backups/sdc-sphinx-theme/static/numba-dpex-logo.png b/docs/backups/sdc-sphinx-theme/static/numba-dpex-logo.png
deleted file mode 100644
index a0f224e883..0000000000
Binary files a/docs/backups/sdc-sphinx-theme/static/numba-dpex-logo.png and /dev/null differ
diff --git a/docs/backups/sdc-sphinx-theme/static/sdc.css b/docs/backups/sdc-sphinx-theme/static/sdc.css
deleted file mode 100644
index b9e6e37205..0000000000
--- a/docs/backups/sdc-sphinx-theme/static/sdc.css
+++ /dev/null
@@ -1,694 +0,0 @@
-/*!
- * Bootstrap v1.4.0
- *
- * Copyright 2011 Twitter, Inc
- * Licensed under the Apache License v2.0
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Heavily modified by Kyle Barbary for the AstroPy Project for use with Sphinx.
- * Modified by Intel Corporation for Intel(R) Scalable Dataframe Compiler
- */
-
-@import url("basic.css");
-
-/* Caption text --------------------------------------------------------------*/
-span-caption-text {
-  font-weight: bold;
-}
-
-/* Body text -----------------------------------------------------------------*/
-body {
-  background-color: #ffffff;
-  margin: 0;
-  font-family: "Helvetica Neue", Helvetica, Arial, sans-serif;
-  font-size: 13px;
-  font-weight: normal;
-  line-height: 18px;
-  color: #001d33;
-}
-
-/* Hyperlinks ----------------------------------------------------------------*/
-
-a {
-  color: #0069d6;
-  text-decoration: none;
-  line-height: inherit;
-  font-weight: inherit;
-}
-
-a:hover {
-  color: #00438a;
-  text-decoration: underline;
-}
-
-/* Typography ----------------------------------------------------------------*/
-
-h1,h2,h3,h4,h5,h6 {
-  color: #003a66;
-  margin: 0.7em 0 0 0;
-  line-height: 1.5em;
-}
-h1 {
-  font-size: 26px;
-  margin: 0;
-}
-h2 {
-  font-size: 20px;
-  line-height: 1.2em;
-  margin: 1em 0 0.5em 0;
-  border-bottom: 1px solid #0070c5;
-}
-h3 {
-  font-size: 18px;
-}
-h4 {
-  font-size: 16px;
-}
-h5 {
-  font-size: 14px;
-}
-h6 {
-  font-size: 13px;
-  text-transform: uppercase;
-}
-
-p {
-  font-size: 13px;
-  font-weight: normal;
-  line-height: 18px;
-  margin-top: 0px;
-  margin-bottom: 9px;
-}
-
-ul, ol {
-  margin-left: 0;
-  padding: 0 0 0 25px;
-}
-ul ul, ul ol, ol ol, ol ul {
-  margin-bottom: 0;
-}
-ul {
-  list-style: disc;
-}
-ol {
-  list-style: decimal;
-}
-li {
-  line-height: 18px;
-  color: #000010;
-}
-ul.unstyled {
-  list-style: none;
-  margin-left: 0;
-}
-dl {
-  margin-bottom: 18px;
-}
-dl dt, dl dd {
-  line-height: 18px;
-}
-dl dd {
-  margin-left: 9px;
-}
-hr {
-  margin: 20px 0 19px;
-  border: 0;
-  border-bottom: 1px solid #eee;
-}
-strong {
-  font-style: inherit;
-  font-weight: bold;
-}
-em {
-  font-style: italic;
-  font-weight: inherit;
-  line-height: inherit;
-}
-.muted {
-  color: #bfbfbf;
-}
-
-address {
-  display: block;
-  line-height: 18px;
-  margin-bottom: 18px;
-}
-code, pre {
-  padding: 0 3px 2px;
-  font-family: monospace;
-  -webkit-border-radius: 3px;
-  -moz-border-radius: 3px;
-  border-radius: 3px;
-}
-tt {
-  font-family: monospace;
-}
-code {
-  padding: 1px 3px;
-}
-pre {
-  display: block;
-  padding: 8.5px;
-  margin: 0 0 18px;
-  line-height: 18px;
-  border: 1px solid #ddd;
-  border: 1px solid rgba(0, 0, 0, 0.12);
-  -webkit-border-radius: 3px;
-  -moz-border-radius: 3px;
-  border-radius: 3px;
-  white-space: pre;
-  word-wrap: break-word;
-}
-
-img {
-  margin: 9px 0;
-}
-
-/* format inline code with a rounded box */
-tt, code {
-    margin: 0 2px;
-    padding: 0 5px;
-    border: 1px solid #ddd;
-    border: 1px solid rgba(0, 0, 0, 0.12);
-    border-radius: 3px;
-}
-
-code.xref, a code {
-    margin: 0;
-    padding: 0 1px 0 1px;
-    background-color: transparent;
-    border: none;
-}
-
-/* all code has same box background color, even in headers */
-h1 tt, h2 tt, h3 tt, h4 tt, h5 tt, h6 tt,
-h1 code, h2 code, h3 code, h4 code, h5 code, h6 code,
-pre, code, tt {
-  background-color: #f8f8f8;
-}
-
-/* override box for links & other sphinx-specifc stuff */
-tt.xref, a tt, tt.descname, tt.descclassname {
-  padding: 0 1px 0 1px;
-  border: none;
-}
-
-/* override box for related bar at the top of the page */
-.related tt {
-  border: none;
-  padding: 0 1px 0 1px;
-  background-color: transparent;
-  font-weight: bold;
-}
-
-th {
-    background-color: #dddddd;
-}
-
-.viewcode-back {
-    font-family: sans-serif;
-}
-
-div.viewcode-block:target {
-    background-color: #f4debf;
-    border-top: 1px solid #ac9;
-    border-bottom: 1px solid #ac9;
-}
-
-table.docutils {
-    border-spacing: 5px;
-    border-collapse: separate;
-}
-
-/* Topbar --------------------------------------------------------------------*/
-
-/* Horizontal bar */
-div.topbar {
-  height: 70px;
-  position: absolute;
-  top: 0;
-  left: 0;
-  right: 0;
-  z-index: 10000;
-  padding: 0px 10px;
-  background-color: #0070c5;
-  background-repeat: repeat-x;
-  background-image: -khtml-gradient(linear, left top, left bottom, from(#0070c5), to(#0091ff));
-  background-image: -moz-linear-gradient(top, #0070c5, #0091ff);
-  background-image: -ms-linear-gradient(top, #0070c5, #0091ff);
-  background-image: -webkit-gradient(linear, left top, left bottom, color-stop(0%, #0070c5), color-stop(100%, #0091ff));
-  background-image: -webkit-linear-gradient(top, #0070c5, #0091ff);
-  background-image: -o-linear-gradient(top, #0070c5, #0091ff);
-  background-image: linear-gradient(to top, #0070c5, #0091ff);
-  filter: progid:DXImageTransform.Microsoft.gradient(startColorstr='#0070c5', endColorstr='#0091ff', GradientType=0);
-  overflow-x: hidden;
-  overflow-y: hidden;
-  width: 100%;
-}
-
-/* Scalable Dataframe Compiler png */
-div.topbar a.brand_sdc {
-  width: 200px;
-  float: left;
-  display: block;
-  height: 60px;
-  padding: 8px 12px 0px 45px;
-  margin-left: -10px;
-  background: transparent url("numba-dpex-logo.png") no-repeat 10px 4px;
-  background-size: 200px 50px;
-}
-
-
-div.topbar .brand:hover, div.topbar .brand_sdc:hover, div.topbar ul li a.homelink:hover {
-  background-color: #333;
-  background-color: rgba(255, 255, 255, 0.05);
-}
-
-div.topbar ul {
-    font-size: 120%;
-    list-style-type: none;
-    margin: 20px;
-    padding: 0 0 0 10px;
-    float: right;
-    color: #eeeeee;
-    text-align: center;
-    text-decoration: none;
-    height: 100%;
-    vertical-align: middle;
-}
-div.topbar ul li {
-    float: left;
-    display: inline-block;
-    height: 60px;
-    margin: 5px;
-    padding: 0px;
-    vertical-align: middle;
-}
-
-div.topbar ul li a {
-  color: #eeeeee;
-  text-decoration: none;
-  padding: 5px;
-  display: block;
-  height: auto;
-  text-align: center;
-  vertical-align: middle;
-  border-radius: 4px;
-}
-
-div.topbar ul li a:hover {
-  color: #ffffff;
-  text-decoration: none;
-}
-
-
-div.dropdown {
-    position: relative; /* Fixed this to relative */
-    display: inline-block;
-    z-index: 999999;
-}
-
-div.dropdown-content {
-    display: none; /* Fix this at none */
-    background-color: DimGray;
-    color: White;
-    width: 235px;
-    height: 155px;
-    box-shadow: 0px 8px 16px 0px rgba(0,0,0,0.2);
-    padding-top: 3px;
-    padding-right: 15px;
-    position: fixed;
-}
-
-div.dropdown:hover .dropdown-content {
-    display: block;
-}
-
-div.dropdown-content a {
-    display: block;
-}
-
-div.dropdown-content a:hover {
-    color: white;
-}
-
-div.dropdown a:hover {
-    color: #FF5000;
-}
-
-div.dropdown z:hover {
-    color: #FF5000;
-    width: 10px;
-    display: none;
-}
-
-div.dropdown z {
-    color: white;
-}
-
-div.dropdown:after {
-    content: "";
-    position: absolute;
-    right: -13px;
-    top: 9px;
-    width: 0;
-    height: 0;
-    border-left: 5px solid transparent;
-    border-right: 5px solid transparent;
-    border-top: 5px solid white;
-    z-index: 999999;
-}
-
-div.dropdown:hover {
-  cursor: pointer;
-}
-
-/* Intel(R) SDC Github page */
-div.topbar ul li a.homelink {
-  margin: -20px;
-  width: 90px;
-  display: block;
-  height: 60px;
-  padding: 5px 0px;
-  background: transparent url("sdc_github.png") no-repeat 5px 0px;
-  background-size: 60px 60px;
-}
-
-/* Intel(R) SDC Examples & Tutorials */
-div.topbar ul li a.exampleslink {
-  margin: -20px;
-  width: 90px;
-  display: block;
-  height: 60px;
-  padding: 5px 0px;
-  background: transparent url("sdc_examples.png") no-repeat 5px 0px;
-  background-size: 60px 60px;
-}
-
-/* Intel(R) SDC Email */
-div.topbar ul li a.emaillink {
-  margin: -20px;
-  width: 90px;
-  display: block;
-  height: 60px;
-  padding: 5px 0px;
-  background: transparent url("sdc_email.png") no-repeat 5px 0px;
-  background-size: 60px 60px;
-}
-
-/* Intel(R) SDC Issues */
-div.topbar ul li a.issueslink {
-  margin: -20px;
-  width: 90px;
-  display: block;
-  height: 60px;
-  padding: 5px 0px;
-  background: transparent url("sdc_issues.png") no-repeat 5px 0px;
-  background-size: 60px 60px;
-}
-
-
-div.topbar form {
-  text-align: left;
-  margin: 0 0 0 5px;
-  position: relative;
-  filter: alpha(opacity=100);
-  -khtml-opacity: 1;
-  -moz-opacity: 1;
-  opacity: 1;
-}
-
-div.topbar input {
-  background-color: #444;
-  background-color: rgba(255, 255, 255, 0.3);
-  font-family: "Helvetica Neue", Helvetica, Arial, sans-serif;
-  font-size: medium;
-  font-weight: 200;
-  line-height: 1;
-  padding: 4px 9px;
-  color: #ffffff;
-  color: rgba(255, 255, 255, 0.75);
-  border: 1px solid #111;
-  -webkit-border-radius: 4px;
-  -moz-border-radius: 4px;
-  border-radius: 4px;
-  -webkit-box-shadow: inset 0 1px 2px rgba(0, 0, 0, 0.1), 0 1px 0px rgba(255, 255, 255, 0.25);
-  -moz-box-shadow: inset 0 1px 2px rgba(0, 0, 0, 0.1), 0 1px 0px rgba(255, 255, 255, 0.25);
-  box-shadow: inset 0 1px 2px rgba(0, 0, 0, 0.1), 0 1px 0px rgba(255, 255, 255, 0.25);
-  -webkit-transition: none;
-  -moz-transition: none;
-  -ms-transition: none;
-  -o-transition: none;
-  transition: none;
-}
-div.topbar input:-moz-placeholder {
-  color: #e6e6e6;
-}
-div.topbar input::-webkit-input-placeholder {
-  color: #e6e6e6;
-}
-div.topbar input:hover {
-  background-color: #bfbfbf;
-  background-color: rgba(255, 255, 255, 0.5);
-  color: #ffffff;
-}
-div.topbar input:focus, div.topbar input.focused {
-  outline: 0;
-  background-color: #ffffff;
-  color: #0070c5;
-  text-shadow: 0 1px 0 #ffffff;
-  border: 0;
-  padding: 5px 10px;
-  -webkit-box-shadow: 0 0 3px rgba(0, 0, 0, 0.15);
-  -moz-box-shadow: 0 0 3px rgba(0, 0, 0, 0.15);
-  box-shadow: 0 0 3px rgba(0, 0, 0, 0.15);
-}
-
-
-/* Relation bar (breadcrumbs, prev, next) ------------------------------------*/
-
-div.related {
-    height: 21px;
-    width: auto;
-    margin: 0 10px;
-    position: relative;
-    top: 42px;
-    clear: both;
-    left: 0;
-    right: 0;
-    z-index: 999;
-    font-size: 100%;
-    vertical-align: middle;
-    background-color: #fff;
-    border-bottom: 1px solid #bbb;
-}
-div.related ul {
-    padding: 0;
-    margin: 0;
-}
-
-
-/* Footer --------------------------------------------------------------------*/
-
-footer {
-  display: block;
-  margin: 10px 10px 0px;
-  padding: 10px 0 0 0;
-  border-top: 1px solid #bbb;
-}
-.pull-right {
-  float: right;
-  width: 30em;
-  text-align: right;
-}
-
-
-/* Sphinx sidebar ------------------------------------------------------------*/
-div.sphinxsidebar {
-    font-size: inherit;
-    border-radius: 3px;
-    background-color: #eee;
-    border: 1px solid #bbb;
-    word-wrap: break-word;
-    /* overflow-wrap is the canonical name for word-wrap in the CSS3 text draft.
-    We include it here mainly for future-proofing. */
-    overflow-wrap: break-word;
-}
-
-div.sphinxsidebarwrapper {
-    padding: 0px 0px 0px 5px;
-}
-
-div.sphinxsidebar h3 {
-    font-family: 'Trebuchet MS', sans-serif;
-    font-size: 1.4em;
-    font-weight: normal;
-    margin: 5px 0px 0px 5px;
-    padding: 0;
-    line-height: 1.6em;
-}
-div.sphinxsidebar h4 {
-    font-family: 'Trebuchet MS', sans-serif;
-    font-size: 1.3em;
-    font-weight: normal;
-    margin: 5px 0 0 0;
-    padding: 0;
-}
-div.sphinxsidebar p {
-}
-div.sphinxsidebar p.topless {
-    margin: 5px 10px 10px 10px;
-}
-div.sphinxsidebar ul {
-    margin: 0px 0px 0px 5px;
-    padding: 0;
-}
-
-div.sphinxsidebar ul ul {
-    margin-left: 15px;
-    list-style-type: disc;
-}
-
-/* If showing the global TOC (toctree),
-   color the current page differently */
-div.sphinxsidebar a.current {
-    color: #0070c5;
-}
-div.sphinxsidebar a.current:hover {
-    color: #0070c5;
-}
-
-
-/* document, documentwrapper, body, bodywrapper ----------------------------- */
-
-div.document {
-    margin-top: 72px;
-    margin-left: 10px;
-    margin-right: 10px;
-}
-
-div.documentwrapper {
-    float: left;
-    width: 100%;
-}
-
-div.body {
-    background-color: #ffffff;
-    padding: 0 0 0px 20px;
-}
-
-div.bodywrapper {
-    margin: 0 0 0 230px;
-    max-width: 55em;
-}
-
-
-/* Header links ------------------------------------------------------------- */
-
-a.headerlink {
-    font-size: 0.8em;
-    padding: 0 4px 0 4px;
-    text-decoration: none;
-}
-
-a.headerlink:hover {
-    background-color: #0069d6;
-    color: white;
-    text-decoration: none;
-}
-
-
-/* Admonitions and warnings ------------------------------------------------- */
-
-/* Shared by admonitions and warnings */
-div.admonition,
-div.warning {
-    padding: 0px;
-    border-radius: 3px;
-    -moz-border-radius: 3px;
-    -webkit-border-radius: 3px;
-}
-div.admonition p,
-div.warning p {
-    margin: 0.5em 1em 0.5em 1em;
-    padding: 0;
-}
-div.admonition pre,
-div.warning pre {
-    margin: 0.4em 1em 0.4em 1em;
-}
-div.admonition p.admonition-title,
-div.warning p.admonition-title {
-    margin: 0;
-    padding: 0.1em 0 0.1em 0.5em;
-    color: white;
-    font-weight: bold;
-    font-size: 1.1em;
-}
-div.admonition ul, div.admonition ol,
-div.warning ul, div.warning ol {
-    margin: 0.1em 0.5em 0.5em 3em;
-    padding: 0;
-}
-
-/* Admonitions only */
-div.admonition {
-    border: 1px solid  #0066cc;
-    background-color: #b3d9ff;
-}
-div.admonition p.admonition-title {
-    background-color: #1a8cff;
-}
-
-/* Warnings only */
-div.warning {
-    border: 1px solid #805500;
-    background-color: #ffe6b3;
-}
-div.warning p.admonition-title {
-    background-color: #ffc34d;
-}
-
-
-/* Figures ------------------------------------------------------------------ */
-
-.figure.align-center {
-    clear: none;
-}
-
-/* This is a div for containing multiple figures side-by-side, for use with
- * .. container:: figures */
-div.figures {
-    border: 1px solid #CCCCCC;
-    background-color: #F8F8F8;
-    margin: 1em;
-    text-align: center;
-}
-
-div.figures .figure {
-    clear: none;
-    float: none;
-    display: inline-block;
-    border: none;
-    margin-left: 0.5em;
-    margin-right: 0.5em;
-}
-
-.field-list th {
-    white-space: nowrap;
-}
-
-table.field-list {
-    border-spacing: 0px;
-    margin-left: 1px;
-    border-left: 5px solid rgb(238, 238, 238) !important;
-}
-
-table.field-list th.field-name {
-    display: inline-block;
-    padding: 1px 8px 1px 5px;
-    white-space: nowrap;
-    background-color: rgb(238, 238, 238);
-    border-radius: 0 3px 3px 0;
-    -webkit-border-radius: 0 3px 3px 0;
-}
diff --git a/docs/backups/sdc-sphinx-theme/static/sdc_email.png b/docs/backups/sdc-sphinx-theme/static/sdc_email.png
deleted file mode 100644
index 9fde61751d..0000000000
Binary files a/docs/backups/sdc-sphinx-theme/static/sdc_email.png and /dev/null differ
diff --git a/docs/backups/sdc-sphinx-theme/static/sdc_examples.png b/docs/backups/sdc-sphinx-theme/static/sdc_examples.png
deleted file mode 100644
index dd7095b80e..0000000000
Binary files a/docs/backups/sdc-sphinx-theme/static/sdc_examples.png and /dev/null differ
diff --git a/docs/backups/sdc-sphinx-theme/static/sdc_github.png b/docs/backups/sdc-sphinx-theme/static/sdc_github.png
deleted file mode 100644
index 89467c51c4..0000000000
Binary files a/docs/backups/sdc-sphinx-theme/static/sdc_github.png and /dev/null differ
diff --git a/docs/backups/sdc-sphinx-theme/static/sdc_issues.png b/docs/backups/sdc-sphinx-theme/static/sdc_issues.png
deleted file mode 100644
index 13d597429e..0000000000
Binary files a/docs/backups/sdc-sphinx-theme/static/sdc_issues.png and /dev/null differ
diff --git a/docs/backups/sdc-sphinx-theme/static/sidebar.js b/docs/backups/sdc-sphinx-theme/static/sidebar.js
deleted file mode 100644
index 8922cda0ca..0000000000
--- a/docs/backups/sdc-sphinx-theme/static/sidebar.js
+++ /dev/null
@@ -1,160 +0,0 @@
-﻿/*
- * sidebar.js
- * ~~~~~~~~~~
- *
- * This script makes the Sphinx sidebar collapsible.
- *
- * .sphinxsidebar contains .sphinxsidebarwrapper.  This script adds
- * in .sphixsidebar, after .sphinxsidebarwrapper, the #sidebarbutton
- * used to collapse and expand the sidebar.
- *
- * When the sidebar is collapsed the .sphinxsidebarwrapper is hidden
- * and the width of the sidebar and the margin-left of the document
- * are decreased. When the sidebar is expanded the opposite happens.
- * This script saves a per-browser/per-session cookie used to
- * remember the position of the sidebar among the pages.
- * Once the browser is closed the cookie is deleted and the position
- * reset to the default (expanded).
- *
- * :copyright: Copyright 2007-2011 by the Sphinx team, see AUTHORS.
- * :license: BSD, see LICENSE for details.
- *
- */
-
-$(function() {
-  // global elements used by the functions.
-  // the 'sidebarbutton' element is defined as global after its
-  // creation, in the add_sidebar_button function
-  var bodywrapper = $('.bodywrapper');
-  var sidebar = $('.sphinxsidebar');
-  var sidebarwrapper = $('.sphinxsidebarwrapper');
-
-  // for some reason, the document has no sidebar; do not run into errors
-  if (!sidebar.length) return;
-
-  // original margin-left of the bodywrapper and width of the sidebar
-  // with the sidebar expanded
-  var bw_margin_expanded = bodywrapper.css('margin-left');
-  var ssb_width_expanded = sidebar.width();
-
-  // margin-left of the bodywrapper and width of the sidebar
-  // with the sidebar collapsed
-  var bw_margin_collapsed = 12;
-  var ssb_width_collapsed = 12;
-
-  // custom colors
-  var dark_color = '#003a66';
-  var light_color = '#0070c5';
-
-  function sidebar_is_collapsed() {
-    return sidebarwrapper.is(':not(:visible)');
-  }
-
-  function toggle_sidebar() {
-    if (sidebar_is_collapsed())
-      expand_sidebar();
-    else
-      collapse_sidebar();
-  }
-
-  function collapse_sidebar() {
-    sidebarwrapper.hide();
-    sidebar.css('width', ssb_width_collapsed);
-    bodywrapper.css('margin-left', bw_margin_collapsed);
-    sidebarbutton.css({
-        'margin-left': '-1px',
-        'height': bodywrapper.height(),
-	'border-radius': '3px'
-    });
-    sidebarbutton.find('span').text('»');
-    sidebarbutton.attr('title', _('Expand sidebar'));
-    document.cookie = 'sidebar=collapsed';
-  }
-
-  function expand_sidebar() {
-    bodywrapper.css('margin-left', bw_margin_expanded);
-    sidebar.css('width', ssb_width_expanded);
-    sidebarwrapper.show();
-    sidebarbutton.css({
-        'margin-left': ssb_width_expanded - 12,
-        'height': bodywrapper.height(),
-	'border-radius': '0px 3px 3px 0px'
-    });
-    sidebarbutton.find('span').text('«');
-    sidebarbutton.attr('title', _('Collapse sidebar'));
-    document.cookie = 'sidebar=expanded';
-  }
-
-  function add_sidebar_button() {
-    sidebarwrapper.css({
-        'float': 'left',
-        'margin-right': '0',
-        'width': ssb_width_expanded - 18
-    });
-    // create the button
-    sidebar.append('<div id="sidebarbutton"><span>&laquo;</span></div>');
-    var sidebarbutton = $('#sidebarbutton');
-
-    // find the height of the viewport to center the '<<' in the page
-    var viewport_height;
-    if (window.innerHeight)
- 	  viewport_height = window.innerHeight;
-    else
-	  viewport_height = $(window).height();
-    var sidebar_offset = sidebar.offset().top;
-    var sidebar_height = Math.max(bodywrapper.height(), sidebar.height());
-    sidebarbutton.find('span').css({
-        'font-family': '"Lucida Grande",Arial,sans-serif',
-	'display': 'block',
-	'top': Math.min(viewport_height/2, sidebar_height/2 + sidebar_offset) - 10,
-	'width': 12,
-	'position': 'fixed',
-	'text-align': 'center'
-    });
-
-    sidebarbutton.click(toggle_sidebar);
-    sidebarbutton.attr('title', _('Collapse sidebar'));
-    sidebarbutton.css({
-        'color': '#FFFFFF',
-	'background-color': light_color,
-	'border': '1px solid ' + light_color,
-        'border-radius': '0px 3px 3px 0px',
-        'font-size': '1.2em',
-        'cursor': 'pointer',
-        'height': sidebar_height,
-        'padding-top': '1px',
-	'margin': '-1px',
-        'margin-left': ssb_width_expanded - 12
-    });
-
-    sidebarbutton.hover(
-      function () {
-          $(this).css('background-color', dark_color);
-      },
-      function () {
-          $(this).css('background-color', light_color);
-      }
-    );
-  }
-
-  function set_position_from_cookie() {
-    if (!document.cookie)
-      return;
-    var items = document.cookie.split(';');
-    for(var k=0; k<items.length; k++) {
-      var key_val = items[k].split('=');
-      var key = key_val[0];
-      if (key == 'sidebar') {
-        var value = key_val[1];
-        if ((value == 'collapsed') && (!sidebar_is_collapsed()))
-          collapse_sidebar();
-        else if ((value == 'expanded') && (sidebar_is_collapsed()))
-          expand_sidebar();
-      }
-    }
-  }
-
-  add_sidebar_button();
-  var sidebarbutton = $('#sidebarbutton');
-  set_position_from_cookie();
-});
diff --git a/docs/backups/sdc-sphinx-theme/theme.conf b/docs/backups/sdc-sphinx-theme/theme.conf
deleted file mode 100644
index 4184857de0..0000000000
--- a/docs/backups/sdc-sphinx-theme/theme.conf
+++ /dev/null
@@ -1,11 +0,0 @@
-# Intel(R) SDC theme based on Bootstrap-AstroPy CSS
-[theme]
-inherit = basic
-stylesheet = sdc.css
-pygments_style = sphinx
-
-[options]
-logotext1 = astro
-logotext2 = py
-logotext3 = :docs
-astropy_project_menubar = False
diff --git a/docs/backups/user_guides/getting_started.rst b/docs/backups/user_guides/getting_started.rst
deleted file mode 100644
index 50b4b69534..0000000000
--- a/docs/backups/user_guides/getting_started.rst
+++ /dev/null
@@ -1,157 +0,0 @@
-Getting Started
-===============
-
-Installation
-------------
-
-Numba-dpex depends on following components:
-
-* numba 0.54.* or 0.55.* (`Numba`_)
-* dpctl 0.13.* (`Intel Python dpctl`_)
-* dpnp 0.10.1 (`Intel Python DPNP`_)
-* `dpcpp-llvm-spirv`_ (SPIRV generation from LLVM IR)
-* `llvmdev`_ (LLVM IR generation)
-* `packaging`_
-* `scipy`_ (for testing)
-* `pytest`_ (for testing)
-
-It is recommended to use conda packages from `Intel Distribution for Python`_
-channel or `anaconda.org/intel`_ channel.
-`Intel Distribution for Python`_ is available from `Intel oneAPI`_.
-
-Create conda environment:
-
-.. code-block:: bash
-
-    export ONEAPI_ROOT=/opt/intel/oneapi
-    conda create -n numba-dpex-env numba-dpex dpnp -c ${ONEAPI_ROOT}/conda_channel
-
-Build and Install Conda Package
--------------------------------
-
-Create and activate conda build environment:
-
-.. code-block:: bash
-
-    conda create -n build-env conda-build
-    conda activate build-env
-
-Set environment variable ``ONEAPI_ROOT`` and build conda package:
-
-.. code-block:: bash
-
-    export ONEAPI_ROOT=/opt/intel/oneapi
-    conda build conda-recipe -c ${ONEAPI_ROOT}/conda_channel
-
-Install conda package:
-
-.. code-block:: bash
-
-    conda install numba-dpex
-
-Build and Install with setuptools
----------------------------------
-
-``setup.py`` requires environment variable ``ONEAPI_ROOT`` and following packages
-installed in conda environment:
-
-.. code-block:: bash
-
-    export ONEAPI_ROOT=/opt/intel/oneapi
-    conda create -n numba-dpex-env -c ${ONEAPI_ROOT}/conda_channel python=3.7 dpctl dpnp numba dpcpp-llvm-spirv llvmdev pytest
-    conda activate numba-dpex-env
-
-Activate DPC++ compiler:
-
-.. code-block:: bash
-
-    source ${ONEAPI_ROOT}/compiler/latest/env/vars.sh
-
-For installing:
-
-.. code-block:: bash
-
-    python setup.py install
-
-For development:
-
-.. code-block:: bash
-
-    python setup.py develop
-
-
-Build and Install with docker
----------------------------------
-
-.. code-block:: bash
-
-    docker run --rm -it \
-    -v /path/to/numba-dpex/source:/build \
-    -v /path/to/dist:/dist
-    ghcr.io/intelpython/numba-dpex/builder:0.20.0-py3.10
-
-    python setup.py develop
-    python setup.py bdist_wheel
-    cp dist/numba_dpex*.whl /dist/
-
-Now you can install numba-dpex wheel in whatever compatible environment with ``pip``.
-You will find ``numba_dpex*.whl`` file in the ``/path/to/dist`` location in
-your host system.
-
-You can check what dpctl and dpnp is shipped with builder by running ``pip list``.
-In case you need another version, consider building ``builder`` target with necessary
-build args. Refer to :ref:`Docker <docker>` section for more details.
-
-
-Testing
--------
-
-See folder ``numba_dpex/tests``.
-
-To run the tests:
-
-.. code-block:: bash
-
-    python -m pytest --pyargs numba_dpex.tests
-
-Examples
---------
-
-See folder ``numba_dpex/examples``.
-
-To run the examples:
-
-.. code-block:: bash
-
-    python numba_dpex/examples/sum.py
-
-Docker
-------
-
-An easy way you can try `numba_dpex` is by using docker.
-To try out numba dpex simply run:
-
-.. code-block:: bash
-
-    docker run --rm -it ghcr.io/intelpython/numba-dpex/runtime:0.20.0-py3.10
-
-.. code-block:: python
-
-    import dpctl
-
-    dpctl.lsplatform()
-
-Refer to :ref:`Docker <docker>` section for more options.
-
-.. _`Numba`: https://github.com/numba/numba
-.. _`Intel Python Numba`: https://github.com/IntelPython/numba
-.. _`Intel Python dpctl`: https://github.com/IntelPython/dpctl
-.. _`Intel Python dpnp`: https://github.com/IntelPython/dpnp
-.. _`dpcpp-llvm-spirv`: https://github.com/IntelPython/dpcpp-llvm-spirv
-.. _`llvmdev`: https://anaconda.org/intel/llvmdev
-.. _`packaging`: https://packaging.pypa.io/
-.. _`scipy`: https://anaconda.org/intel/scipy
-.. _`pytest`: https://docs.pytest.org
-.. _`Intel Distribution for Python`: https://www.intel.com/content/www/us/en/developer/tools/oneapi/distribution-for-python.html
-.. _`anaconda.org/intel`: https://anaconda.org/intel
-.. _`Intel oneAPI`: https://www.intel.com/content/www/us/en/developer/tools/oneapi/overview.html
diff --git a/docs/backups/user_guides/migrating_from_numba_cuda.rst b/docs/backups/user_guides/migrating_from_numba_cuda.rst
deleted file mode 100644
index 9683a1dcb7..0000000000
--- a/docs/backups/user_guides/migrating_from_numba_cuda.rst
+++ /dev/null
@@ -1,18 +0,0 @@
-numba-dpex for ``numba.cuda`` Developers
-========================================
-
-.. todo::
-
-   Coming soon.
-
-.. Transition from Numba CUDA
-.. --------------------------
-
-.. Replace ``@cuda.jit(device=True)`` with ``@numba_dpex.func``.
-
-.. See also
-.. --------
-
-.. Examples:
-
-.. - ``numba_dpex/examples/dpex_func.py``
diff --git a/docs/backups/user_guides/performance_tips.rst b/docs/backups/user_guides/performance_tips.rst
deleted file mode 100644
index 2e3132ee33..0000000000
--- a/docs/backups/user_guides/performance_tips.rst
+++ /dev/null
@@ -1,5 +0,0 @@
-.. _performance_tips:
-.. include:: ./../ext_links.txt
-
-Performance Tips
-================
diff --git a/docs/backups/user_guides/troubleshooting.rst b/docs/backups/user_guides/troubleshooting.rst
deleted file mode 100644
index 4b813015e4..0000000000
--- a/docs/backups/user_guides/troubleshooting.rst
+++ /dev/null
@@ -1,5 +0,0 @@
-.. _troubleshooting:
-.. include:: ./../ext_links.txt
-
-Troubleshooting
-===============
diff --git a/docs/source/api_reference/index.rst b/docs/source/api_reference/index.rst
deleted file mode 100644
index 728caaabcf..0000000000
--- a/docs/source/api_reference/index.rst
+++ /dev/null
@@ -1,8 +0,0 @@
-.. _index:
-.. include:: ./../ext_links.txt
-
-
-API Reference
-=============
-
-Coming soon
diff --git a/docs/source/conf.py b/docs/source/conf.py
index 60d42b9b98..72a445165c 100644
--- a/docs/source/conf.py
+++ b/docs/source/conf.py
@@ -1,28 +1,6 @@
-# *****************************************************************************
-# Copyright (c) 2022, Intel Corporation All rights reserved.
+# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
 #
-# Redistribution and use in source and binary forms, with or without
-# modification, are permitted provided that the following conditions are met:
-#
-#     Redistributions of source code must retain the above copyright notice,
-#     this list of conditions and the following disclaimer.
-#
-#     Redistributions in binary form must reproduce the above copyright notice,
-#     this list of conditions and the following disclaimer in the documentation
-#     and/or other materials provided with the distribution.
-#
-# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
-# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
-# THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
-# PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR
-# CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
-# EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
-# PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
-# OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
-# WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR
-# OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE,
-# EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-# *****************************************************************************
+# SPDX-License-Identifier: Apache-2.0
 
 # coding: utf-8
 # Configuration file for the Sphinx documentation builder.
@@ -32,7 +10,7 @@
 # -- Project information -----------------------------------------------------
 
 project = "numba-dpex"
-copyright = "2020-2023, Intel Corporation"
+copyright = "2020-2024, Intel Corporation"
 author = "Intel Corporation"
 
 # The full version, including alpha/beta/rc tags
@@ -53,6 +31,7 @@
     "sphinxcontrib.programoutput",
     "sphinxcontrib.googleanalytics",
     "myst_parser",
+    "autoapi.extension",
 ]
 
 # Add any paths that contain templates here, relative to this directory.
@@ -70,22 +49,17 @@
 # The theme to use for HTML and HTML Help pages.  See the documentation for
 # a list of builtin themes.
 #
-html_theme = "pydata_sphinx_theme"
+# html_theme = "pydata_sphinx_theme"
+html_theme = "furo"
 
 html_theme_options = {
-    "icon_links": [
+    "footer_icons": [
         {
             "name": "GitHub",
             "url": "https://github.com/IntelPython/numba-dpex",
             "icon": "fab fa-github-square",
         },
-        {
-            "name": "Gitter",
-            "url": "https://app.gitter.im/#/room/#Data-Parallel-Python_community:gitter.im",
-            "icon": "fab fa-brands fa-gitter",
-        },
     ],
-    # "logo_only": True,
 }
 
 googleanalytics_id = "G-LGGL0NJK6P"
@@ -137,3 +111,52 @@
 
 # -- Prepend module name to an object name or not -----------------------------------
 add_module_names = False
+
+# -- Auto API configurations ---------------------------------------------------
+
+
+# def skip_util_classes(app, what, name, obj, skip, options):
+#     if what == "module" and "experimental" in name:
+#         if what == "module" and "decorators" not in name:
+#             skip = True
+#     return skip
+
+
+# def setup(sphinx):
+#     sphinx.connect("autoapi-skip-member", skip_util_classes)
+
+
+autoapi_dirs = ["../../numba_dpex/kernel_api"]
+autoapi_type = "python"
+
+autoapi_template_dir = "_templates/autoapi"
+
+autoapi_options = [
+    "members",
+    "undoc-members",
+    "show-inheritance",
+    "show-module-summary",
+    "imported-members",
+]
+
+autoapi_keep_files = True
+autodoc_typehints = "signature"
+
+rst_prolog = """
+.. role:: summarylabel
+"""
+
+html_css_files = [
+    "css/custom.css",
+]
+
+
+def contains(seq, item):
+    return item in seq
+
+
+def prepare_jinja_env(jinja_env) -> None:
+    jinja_env.tests["contains"] = contains
+
+
+autoapi_prepare_jinja_env = prepare_jinja_env
diff --git a/docs/source/getting_started.rst b/docs/source/getting_started.rst
index f7f24d0495..769b300529 100644
--- a/docs/source/getting_started.rst
+++ b/docs/source/getting_started.rst
@@ -19,7 +19,7 @@ to get the latest production releases.
 .. code-block:: bash
 
     conda create -n numba-dpex-env                                             \
-        numba-dpex dpnp dpctl dpcpp-llvm-spirv                     \
+        numba-dpex dpnp dpctl dpcpp-llvm-spirv                                 \
         -c intel -c conda-forge
 
 To try out the bleeding edge, the latest packages built from tip of the main
@@ -28,7 +28,7 @@ source trunk can be installed from the ``dppy/label/dev`` conda channel.
 .. code-block:: bash
 
     conda create -n numba-dpex-env                                             \
-        numba-dpex dpnp dpctl dpcpp-llvm-spirv                     \
+        numba-dpex dpnp dpctl dpcpp-llvm-spirv                                 \
         -c dppy/label/dev -c intel -c conda-forge
 
 
diff --git a/docs/source/index.rst b/docs/source/index.rst
index c8ce20c0c9..1bf6f0ab62 100644
--- a/docs/source/index.rst
+++ b/docs/source/index.rst
@@ -45,8 +45,10 @@ Data Parallel Extension for Numba*
 
    overview
    getting_started
+   programming_model
    user_guide/index
-   api_reference/index
+   autoapi/index
+   useful_links
 
 .. toctree::
    :maxdepth: 2
diff --git a/docs/source/user_guide/programming_model.rst b/docs/source/programming_model.rst
similarity index 99%
rename from docs/source/user_guide/programming_model.rst
rename to docs/source/programming_model.rst
index 0543036f91..168120fc9d 100644
--- a/docs/source/user_guide/programming_model.rst
+++ b/docs/source/programming_model.rst
@@ -1,5 +1,5 @@
 .. _programming_model:
-.. include:: ./../ext_links.txt
+.. include:: ./ext_links.txt
 
 Programming Model
 =================
diff --git a/docs/source/user_guide/useful_links.rst b/docs/source/useful_links.rst
similarity index 98%
rename from docs/source/user_guide/useful_links.rst
rename to docs/source/useful_links.rst
index 9d33d29615..221cc5cdce 100644
--- a/docs/source/user_guide/useful_links.rst
+++ b/docs/source/useful_links.rst
@@ -1,5 +1,5 @@
 .. _useful_links:
-.. include:: ./../ext_links.txt
+.. include:: ./ext_links.txt
 
 Useful links
 ============
diff --git a/docs/source/user_guide/caching.rst b/docs/source/user_guide/caching.rst
deleted file mode 100644
index 59c9e48385..0000000000
--- a/docs/source/user_guide/caching.rst
+++ /dev/null
@@ -1,49 +0,0 @@
-.. include:: ./../ext_links.txt
-
-Caching Mechanism in ``numba-dpex``
-===================================
-
-Caching is done by saving the compiled kernel code, the ELF object of the
-executable code. By using the kernel code, cached kernels have minimal overhead
-because no compilation is needed.
-
-Unlike Numba*, ``numba-dpex`` does not perform an exclusive file-based caching,
-instead a Least Recently Used (LRU) caching mechanism is used. However, when a
-kernel needs to be evicted, Numba's file-based caching mechanism is invoked as
-described `here
-<https://numba.pydata.org/numba-doc/latest/developer/caching.html>`_.
-
-Algorithm
-----------
-
-The caching mechanism for ``numba-dpex`` works as follows: The cache is an LRU
-cache backed by an ordered dictionary mapped onto a doubly linked list. The tail
-of the list contains the most recently used (MRU) kernel and the head of the
-list contains the least recently used (LRU) kernel. The list  has a fixed size.
-If a new kernel arrives to be cached and if the size is already on the maximum
-limit, the algorithm evicts the LRU kernel to make room for the MRU kernel. The
-evicted item will be serialized and pickled into a file using Numba's caching
-mechanism.
-
-Everytime when a kernel needs to be retrieved from the cache, the mechanism
-will look for the kernel in the cache and will be loaded if it's already
-present. However, if the program is seeking for a kernel that has been evicted,
-the algorithm will load it from the file and enqueue in the cache. As a result,
-the amount of file operations are significantly lower than that of Numba.
-
-Settings
----------
-
-Therefore, ``numba-dpex`` employs similar environment variables as used in
-Numba, i.e. ``NUMBA_CACHE_DIR`` etc. However there are three more environment
-variables to control the caching mechanism.
-
-- In order to specify cache capacity, ``NUMBA_DPEX_CACHE_SIZE`` can be used. By
-  default, it's set to 10.
-
-- ``NUMBA_DPEX_ENABLE_CACHE`` can be used to enable/disable the caching
-  mechanism. By default it's enabled, i.e. set to 1.
-
-- In order to enable the debugging messages related to caching, the variable
-``NUMBA_DPEX_DEBUG_CACHE`` can be set to 1. All environment variables are
-defined in :file:`numba_dpex/config.py`.
diff --git a/docs/source/user_guide/index.rst b/docs/source/user_guide/index.rst
index eb61b268b3..1cf7381bac 100644
--- a/docs/source/user_guide/index.rst
+++ b/docs/source/user_guide/index.rst
@@ -1,21 +1,13 @@
 .. _index:
 .. include:: ./../ext_links.txt
 
-.. .. image:: ./../../asset/images/DPEP-large.png
-..     :width: 400px
-..     :align: center
-..     :alt: Data Parallel Extensions for Python
-
-User Guide
-==========
+Tutorials
+=========
 
 .. toctree::
     :maxdepth: 2
 
-    programming_model
     kernel_programming/index
     dpnp_offload
     debugging/index
-    caching
     config
-    useful_links
diff --git a/environment/coverage.yml b/environment/coverage.yml
index 1b2f310a19..229d28f2e6 100644
--- a/environment/coverage.yml
+++ b/environment/coverage.yml
@@ -8,7 +8,7 @@ channels:
 dependencies:
   - libffi
   - gxx_linux-64
-  - dpcpp_linux-64>=2023.2,!=2024.0.1,!=2024.0.2
+  - dpcpp_linux-64>=2024.0,!=2024.0.1,!=2024.0.2
   - numba==0.58*
   - dpctl
   - dpnp
diff --git a/environment/docs.yml b/environment/docs.yml
index 4daf1c16b0..056dd9fd06 100644
--- a/environment/docs.yml
+++ b/environment/docs.yml
@@ -8,17 +8,18 @@ channels:
 dependencies:
   - libffi
   - gxx_linux-64
-  - dpcpp_linux-64>=2023.2,!=2024.0.1,!=2024.0.2
-  - numba==0.58*
+  - dpcpp_linux-64>=2024.0,!=2024.0.1,!=2024.0.2
+  - numba==0.59*
   - scikit-build>=0.15*
   - cmake>=3.26*
-  - dpctl
-  - dpnp
+  - dpctl>=0.16*
+  - dpnp>=0.14*
   - dpcpp-llvm-spirv
   - opencl_rt
   - pip
   - pip:
     - sphinx
+    - sphinx-autoapi
     - autodoc # there is no conda package
     - recommonmark
     - sphinx-rtd-theme
@@ -27,3 +28,4 @@ dependencies:
     - sphinxcontrib.programoutput
     - pydata-sphinx-theme
     - myst-parser
+    - furo
diff --git a/environment/pre-commit.yml b/environment/pre-commit.yml
index 24d420bd2c..f6180fd98f 100644
--- a/environment/pre-commit.yml
+++ b/environment/pre-commit.yml
@@ -7,9 +7,9 @@ channels:
   - nodefaults
 dependencies:
   - libffi
-  - numba==0.58*
-  - dpctl
-  - dpnp
+  - numba==0.59*
+  - dpctl>=0.16*
+  - dpnp>=0.14*
   - dpcpp-llvm-spirv
   - opencl_rt
   - coverage
diff --git a/numba_dpex/__init__.py b/numba_dpex/__init__.py
index d754de09b9..f77566dec4 100644
--- a/numba_dpex/__init__.py
+++ b/numba_dpex/__init__.py
@@ -15,11 +15,10 @@
 import dpctl
 import llvmlite.binding as ll
 from numba import __version__ as numba_version
-from numba.np.ufunc.decorators import Vectorize
 
 from numba_dpex.core.kernel_interface.launcher import call_kernel
-from numba_dpex.vectorizers import Vectorize as DpexVectorize
 
+from .kernel_api_impl.spirv import target as spirv_kernel_target
 from .numba_patches import patch_arrayexpr_tree_to_ir, patch_is_ufunc
 
 
@@ -69,18 +68,14 @@ def parse_sem_version(version_string: str) -> Tuple[int, int, int]:
 
 
 numba_sem_version = parse_sem_version(numba_version)
-if numba_sem_version < (0, 57, 0) or numba_sem_version >= (0, 59, 0):
+if numba_sem_version < (0, 58, 0) or numba_sem_version >= (0, 59, 0):
     logging.warning(
-        "numba_dpex needs at least numba 0.57.0 but no more than 0.59.0, using "
+        "numba_dpex needs at least numba 0.58.0 but no more than 0.59.0, using "
         f"numba={numba_version} may cause unexpected behavior"
     )
 
 # Monkey patches
 patch_is_ufunc.patch()
-if numba_sem_version < (0, 58, 0):
-    from .numba_patches import patch_mk_alloc
-
-    patch_mk_alloc.patch()
 patch_arrayexpr_tree_to_ir.patch()
 
 dpctl_sem_version = parse_sem_version(dpctl.__version__)
@@ -100,18 +95,20 @@ def parse_sem_version(version_string: str) -> Tuple[int, int, int]:
 
 # Re-export types itself
 import numba_dpex.core.types as types  # noqa E402
+from numba_dpex.core import boxing  # noqa E402
 from numba_dpex.core import config  # noqa E402
-from numba_dpex.core.kernel_interface.indexers import (  # noqa E402
-    NdRange,
-    Range,
-)
+from numba_dpex.core.kernel_interface import ranges_overloads  # noqa E402
 
 # Re-export all type names
 from numba_dpex.core.types import *  # noqa E402
 from numba_dpex.dpctl_iface import _intrinsic  # noqa E402
 from numba_dpex.dpnp_iface import dpnpimpl  # noqa E402
 
-from .core.targets import dpjit_target, kernel_target  # noqa E402
+# Importing NdRange and Range into numba_dpex for
+# backward compatibility
+from numba_dpex.kernel_api import NdRange, Range  # noqa E402
+
+from .core.targets import dpjit_target  # noqa E402
 from .decorators import dpjit, func, kernel  # noqa E402
 from .ocl.stubs import (  # noqa E402
     GLOBAL_MEM_FENCE,
@@ -135,8 +132,6 @@ def parse_sem_version(version_string: str) -> Tuple[int, int, int]:
 del load_dpctl_sycl_interface
 
 
-Vectorize.target_registry.ondemand["dpex"] = lambda: DpexVectorize
-
 from numba_dpex._version import get_versions  # noqa E402
 
 __version__ = get_versions()["version"]
diff --git a/numba_dpex/core/boxing/__init__.py b/numba_dpex/core/boxing/__init__.py
new file mode 100644
index 0000000000..66ea063322
--- /dev/null
+++ b/numba_dpex/core/boxing/__init__.py
@@ -0,0 +1,10 @@
+# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Contains the ``box`` and ``unbox`` functions for numba_dpex types that are
+passable as arguments to a kernel or dpjit decorated function.
+"""
+
+from .ranges import *
+from .usm_ndarray import *
diff --git a/numba_dpex/core/types/range_types.py b/numba_dpex/core/boxing/ranges.py
similarity index 87%
rename from numba_dpex/core/types/range_types.py
rename to numba_dpex/core/boxing/ranges.py
index c88cd54294..217acebd2a 100644
--- a/numba_dpex/core/types/range_types.py
+++ b/numba_dpex/core/boxing/ranges.py
@@ -4,55 +4,12 @@
 
 from contextlib import ExitStack
 
-from numba.core import cgutils, errors, types
+from numba.core import cgutils, types
 from numba.core.datamodel import default_manager
 from numba.extending import NativeValue, box, unbox
 
-from ..kernel_interface.indexers import NdRange, Range
-
-
-class RangeType(types.Type):
-    """Numba-dpex type corresponding to
-    :class:`numba_dpex.core.kernel_interface.indexers.Range`
-    """
-
-    def __init__(self, ndim: int):
-        self._ndim = ndim
-        if ndim < 1 or ndim > 3:
-            raise errors.TypingError(
-                "RangeType can only have 1,2, or 3 dimensions"
-            )
-        super(RangeType, self).__init__(name="Range<" + str(ndim) + ">")
-
-    @property
-    def ndim(self):
-        return self._ndim
-
-    @property
-    def key(self):
-        return self._ndim
-
-
-class NdRangeType(types.Type):
-    """Numba-dpex type corresponding to
-    :class:`numba_dpex.core.kernel_interface.indexers.NdRange`
-    """
-
-    def __init__(self, ndim: int):
-        self._ndim = ndim
-        if ndim < 1 or ndim > 3:
-            raise errors.TypingError(
-                "RangeType can only have 1,2, or 3 dimensions"
-            )
-        super(NdRangeType, self).__init__(name="NdRange<" + str(ndim) + ">")
-
-    @property
-    def ndim(self):
-        return self._ndim
-
-    @property
-    def key(self):
-        return self._ndim
+from numba_dpex.core.types import NdRangeType, RangeType
+from numba_dpex.kernel_api import NdRange, Range
 
 
 @unbox(RangeType)
diff --git a/numba_dpex/core/boxing/usm_ndarray.py b/numba_dpex/core/boxing/usm_ndarray.py
new file mode 100644
index 0000000000..379362eb78
--- /dev/null
+++ b/numba_dpex/core/boxing/usm_ndarray.py
@@ -0,0 +1,126 @@
+# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+from contextlib import ExitStack
+
+from numba.core import cgutils, types
+from numba.core.datamodel import default_manager
+from numba.core.errors import NumbaNotImplementedError
+from numba.extending import NativeValue, box, unbox
+from numba.np import numpy_support
+
+from numba_dpex.core.exceptions import UnreachableError
+from numba_dpex.core.runtime import context as dpexrt
+from numba_dpex.core.types import USMNdArray
+from numba_dpex.kernel_api import NdRange, Range
+
+
+@unbox(USMNdArray)
+def unbox_dpnp_nd_array(typ, obj, c):
+    """Converts a dpctl.tensor.usm_ndarray/dpnp.ndarray object to a Numba-dpex
+    internal array structure.
+
+    Args:
+        typ : The Numba type of the PyObject
+        obj : The actual PyObject to be unboxed
+        c : The unboxing context
+
+    Returns: A NativeValue object representing an unboxed
+        dpctl.tensor.usm_ndarray/dpnp.ndarray
+    """
+    # Reusing the numba.core.base.BaseContext's make_array function to get a
+    # struct allocated. The same struct is used for numpy.ndarray
+    # and dpnp.ndarray. It is possible to do so, as the extra information
+    # specific to dpnp.ndarray such as sycl_queue is inferred statically and
+    # stored as part of the DpnpNdArray type.
+
+    # --------------- Original Numba comment from @ubox(types.Array)
+    #
+    # This is necessary because unbox_buffer() does not work on some
+    # dtypes, e.g. datetime64 and timedelta64.
+    # TODO check matching dtype.
+    #      currently, mismatching dtype will still work and causes
+    #      potential memory corruption
+    #
+    # --------------- End of Numba comment from @ubox(types.Array)
+    nativearycls = c.context.make_array(typ)
+    nativeary = nativearycls(c.context, c.builder)
+    aryptr = nativeary._getpointer()
+
+    ptr = c.builder.bitcast(aryptr, c.pyapi.voidptr)
+    # FIXME : We need to check if Numba_RT as well as DPEX RT are enabled.
+    if c.context.enable_nrt:
+        dpexrtCtx = dpexrt.DpexRTContext(c.context)
+        errcode = dpexrtCtx.arraystruct_from_python(c.pyapi, obj, ptr)
+    else:
+        raise UnreachableError
+
+    # TODO: here we have minimal typechecking by the itemsize.
+    #       need to do better
+    try:
+        expected_itemsize = numpy_support.as_dtype(typ.dtype).itemsize
+    except NumbaNotImplementedError:
+        # Don't check types that can't be `as_dtype()`-ed
+        itemsize_mismatch = cgutils.false_bit
+    else:
+        expected_itemsize = nativeary.itemsize.type(expected_itemsize)
+        itemsize_mismatch = c.builder.icmp_unsigned(
+            "!=",
+            nativeary.itemsize,
+            expected_itemsize,
+        )
+
+    failed = c.builder.or_(
+        cgutils.is_not_null(c.builder, errcode),
+        itemsize_mismatch,
+    )
+    # Handle error
+    with c.builder.if_then(failed, likely=False):
+        c.pyapi.err_set_string(
+            "PyExc_TypeError",
+            "can't unbox usm array from PyObject into "
+            "native value.  The object maybe of a "
+            "different type",
+        )
+    return NativeValue(c.builder.load(aryptr), is_error=failed)
+
+
+@box(USMNdArray)
+def box_array(typ, val, c):
+    """Boxes a NativeValue representation of USMNdArray/DpnpNdArray type into a
+    dpctl.tensor.usm_ndarray/dpnp.ndarray PyObject
+
+    Args:
+        typ: The representation of the USMNdArray/DpnpNdArray type.
+        val: A native representation of a Numba USMNdArray/DpnpNdArray type
+            object.
+        c: The boxing context.
+
+    Returns: A Pyobject for a dpctl.tensor.usm_ndarray/dpnp.ndarray boxed from
+        the Numba-dpex native value.
+    """
+    if c.context.enable_nrt:
+        np_dtype = numpy_support.as_dtype(typ.dtype)
+        dtypeptr = c.env_manager.read_const(c.env_manager.add_const(np_dtype))
+        dpexrtCtx = dpexrt.DpexRTContext(c.context)
+        newary = dpexrtCtx.usm_ndarray_to_python_acqref(
+            c.pyapi, typ, val, dtypeptr
+        )
+
+        if not newary:
+            c.pyapi.err_set_string(
+                "PyExc_TypeError",
+                "could not box native array into a dpnp.ndarray PyObject.",
+            )
+
+        # Steals NRT ref
+        # Refer:
+        #   numba.core.base.nrt -> numba.core.runtime.context -> decref
+        #   The `NRT_decref` function is generated directly as LLVM IR inside
+        #   numba.core.runtime.nrtdynmod.py
+        c.context.nrt.decref(c.builder, typ, val)
+
+        return newary
+    else:
+        raise UnreachableError
diff --git a/numba_dpex/core/caching.py b/numba_dpex/core/caching.py
index c636fb7096..7e5f80f9b7 100644
--- a/numba_dpex/core/caching.py
+++ b/numba_dpex/core/caching.py
@@ -436,9 +436,9 @@ def put(self, key, value):
                         )
                     )
                 self._cache_file.save(self._head.key, self._head.value)
-                self._evicted[
-                    self._head.key
-                ] = None  # as we are using cache files, we save memory
+                self._evicted[self._head.key] = (
+                    None  # as we are using cache files, we save memory
+                )
             else:
                 self._evicted[self._head.key] = self._head.value
             self._lookup.pop(self._head.key)
diff --git a/numba_dpex/core/datamodel/models.py b/numba_dpex/core/datamodel/models.py
index b3c5ebf71c..2085be7eb8 100644
--- a/numba_dpex/core/datamodel/models.py
+++ b/numba_dpex/core/datamodel/models.py
@@ -2,6 +2,7 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
+from llvmlite import ir as llvmir
 from numba.core import datamodel, types
 from numba.core.datamodel.models import PrimitiveModel, StructModel
 from numba.core.extending import register_model
@@ -14,6 +15,7 @@
     DpctlSyclEvent,
     DpctlSyclQueue,
     DpnpNdArray,
+    IntEnumLiteral,
     NdRangeType,
     RangeType,
     USMNdArray,
@@ -55,6 +57,17 @@ def __init__(self, dmm, fe_type):
         super(GenericPointerModel, self).__init__(dmm, fe_type, be_type)
 
 
+class IntEnumLiteralModel(PrimitiveModel):
+    """Representation of an object of LiteralIntEnum type using Numba's
+    PrimitiveModel that can be represented natively in the target in all
+    usage contexts.
+    """
+
+    def __init__(self, dmm, fe_type):
+        be_type = llvmir.IntType(fe_type.bitwidth)
+        super().__init__(dmm, fe_type, be_type)
+
+
 class USMArrayDeviceModel(StructModel):
     """A data model to represent a usm array type in the LLVM IR generated for a
     device-only kernel function.
@@ -196,9 +209,7 @@ def __init__(self, dmm, fe_type):
 
 
 class RangeModel(StructModel):
-    """The native data model for a
-    numba_dpex.core.kernel_interface.indexers.Range PyObject.
-    """
+    """The numba_dpex data model for a numba_dpex.kernel_api.Range PyObject."""
 
     def __init__(self, dmm, fe_type):
         members = [
@@ -216,8 +227,8 @@ def flattened_field_count(self):
 
 
 class NdRangeModel(StructModel):
-    """The native data model for a
-    numba_dpex.core.kernel_interface.indexers.NdRange PyObject.
+    """
+    The numba_dpex data model for a numba_dpex.kernel_api.NdRange PyObject.
     """
 
     def __init__(self, dmm, fe_type):
@@ -239,7 +250,7 @@ def flattened_field_count(self):
 
 
 def _init_data_model_manager() -> datamodel.DataModelManager:
-    """Initializes a DpexKernelTarget-specific data model manager.
+    """Initializes a data model manager used by the SPRIVTarget.
 
     SPIRV kernel functions for certain types of devices require an explicit
     address space qualifier for pointers. For OpenCL HD Graphics
@@ -254,8 +265,7 @@ def _init_data_model_manager() -> datamodel.DataModelManager:
     a dpnp.ndarray object can be passed to any other regular function.
 
     Returns:
-        DataModelManager: A numba-dpex DpexKernelTarget-specific data model
-        manager
+        DataModelManager: A numba-dpex SPIRVTarget-specific data model manager
     """
     dmm = datamodel.default_manager.copy()
     dmm.register(types.CPointer, GenericPointerModel)
@@ -273,6 +283,8 @@ def _init_data_model_manager() -> datamodel.DataModelManager:
     # model manager. The dpex_data_model_manager is used by the DpexKernelTarget
     dmm.register(DpctlSyclQueue, SyclQueueModel)
 
+    dmm.register(IntEnumLiteral, IntEnumLiteralModel)
+
     return dmm
 
 
diff --git a/numba_dpex/core/descriptor.py b/numba_dpex/core/descriptor.py
index 2524919053..da1ce40558 100644
--- a/numba_dpex/core/descriptor.py
+++ b/numba_dpex/core/descriptor.py
@@ -9,18 +9,18 @@
 from numba.core.descriptors import TargetDescriptor
 
 from numba_dpex.core import config
+from numba_dpex.kernel_api_impl.spirv.target import (
+    SPIRV_TARGET_NAME,
+    CompilationMode,
+    SPIRVTargetContext,
+    SPIRVTypingContext,
+)
 
 from .targets.dpjit_target import (
     DPEX_TARGET_NAME,
     DpexTargetContext,
     DpexTypingContext,
 )
-from .targets.kernel_target import (
-    DPEX_KERNEL_TARGET_NAME,
-    CompilationMode,
-    DpexKernelTargetContext,
-    DpexKernelTypingContext,
-)
 
 _option_mapping = options._mapping
 
@@ -77,12 +77,12 @@ class DpexKernelTarget(TargetDescriptor):
     @cached_property
     def _toplevel_target_context(self):
         """Lazily-initialized top-level target context, for all threads."""
-        return DpexKernelTargetContext(self.typing_context, self._target_name)
+        return SPIRVTargetContext(self.typing_context, self._target_name)
 
     @cached_property
     def _toplevel_typing_context(self):
         """Lazily-initialized top-level typing context, for all threads."""
-        return DpexKernelTypingContext()
+        return SPIRVTypingContext()
 
     @property
     def target_context(self):
@@ -132,7 +132,7 @@ def typing_context(self):
 
 
 # A global instance of the DpexKernelTarget
-dpex_kernel_target = DpexKernelTarget(DPEX_KERNEL_TARGET_NAME)
+dpex_kernel_target = DpexKernelTarget(SPIRV_TARGET_NAME)
 
 # A global instance of the DpexTarget
 dpex_target = DpexTarget(DPEX_TARGET_NAME)
diff --git a/numba_dpex/core/kernel_interface/ranges_overloads.py b/numba_dpex/core/kernel_interface/ranges_overloads.py
new file mode 100644
index 0000000000..f0ed884c8b
--- /dev/null
+++ b/numba_dpex/core/kernel_interface/ranges_overloads.py
@@ -0,0 +1,171 @@
+# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+from llvmlite import ir as llvmir
+from numba.core import cgutils, errors, types
+from numba.core.datamodel import default_manager
+from numba.extending import intrinsic, overload
+
+from numba_dpex.kernel_api import NdRange, Range
+
+# can't import name because of the circular import
+DPEX_TARGET_NAME = "dpex"
+
+
+@intrinsic(target=DPEX_TARGET_NAME)
+def _intrin_range_alloc(typingctx, ty_dim0, ty_dim1, ty_dim2, ty_range):
+    ty_retty = ty_range.instance_type
+    sig = ty_retty(
+        ty_dim0,
+        ty_dim1,
+        ty_dim2,
+        ty_range,
+    )
+
+    def codegen(context, builder, sig, args):
+        typ = sig.return_type
+        dim0, dim1, dim2, _ = args
+        range_struct = cgutils.create_struct_proxy(typ)(context, builder)
+        range_struct.dim0 = dim0
+
+        if not isinstance(sig.args[1], types.NoneType):
+            range_struct.dim1 = dim1
+        else:
+            range_struct.dim1 = llvmir.Constant(
+                llvmir.types.IntType(64), Range.UNDEFINED_DIMENSION
+            )
+
+        if not isinstance(sig.args[2], types.NoneType):
+            range_struct.dim2 = dim2
+        else:
+            range_struct.dim2 = llvmir.Constant(
+                llvmir.types.IntType(64), Range.UNDEFINED_DIMENSION
+            )
+
+        range_struct.ndim = llvmir.Constant(llvmir.types.IntType(64), typ.ndim)
+
+        return range_struct._getvalue()
+
+    return sig, codegen
+
+
+@intrinsic(target=DPEX_TARGET_NAME)
+def _intrin_ndrange_alloc(
+    typingctx, ty_global_range, ty_local_range, ty_ndrange
+):
+    ty_retty = ty_ndrange.instance_type
+    sig = ty_retty(
+        ty_global_range,
+        ty_local_range,
+        ty_ndrange,
+    )
+    range_datamodel = default_manager.lookup(ty_global_range)
+
+    def codegen(context, builder, sig, args):
+        typ = sig.return_type
+
+        global_range, local_range, _ = args
+        ndrange_struct = cgutils.create_struct_proxy(typ)(context, builder)
+        ndrange_struct.ndim = llvmir.Constant(
+            llvmir.types.IntType(64), typ.ndim
+        )
+        ndrange_struct.gdim0 = builder.extract_value(
+            global_range,
+            range_datamodel.get_field_position("dim0"),
+        )
+        ndrange_struct.gdim1 = builder.extract_value(
+            global_range,
+            range_datamodel.get_field_position("dim1"),
+        )
+        ndrange_struct.gdim2 = builder.extract_value(
+            global_range,
+            range_datamodel.get_field_position("dim2"),
+        )
+        ndrange_struct.ldim0 = builder.extract_value(
+            local_range,
+            range_datamodel.get_field_position("dim0"),
+        )
+        ndrange_struct.ldim1 = builder.extract_value(
+            local_range,
+            range_datamodel.get_field_position("dim1"),
+        )
+        ndrange_struct.ldim2 = builder.extract_value(
+            local_range,
+            range_datamodel.get_field_position("dim2"),
+        )
+
+        return ndrange_struct._getvalue()
+
+    return sig, codegen
+
+
+@overload(Range, target=DPEX_TARGET_NAME)
+def _ol_range_init(dim0, dim1=None, dim2=None):
+    """Numba overload of the Range constructor to make it usable inside an
+    njit and dpjit decorated function.
+
+    """
+    from numba_dpex.core.types import RangeType
+
+    ndims = 1
+    ty_optional_dims = (dim1, dim2)
+
+    # A Range should at least have the 0th dimension populated
+    if not isinstance(dim0, types.Integer):
+        raise errors.TypingError(
+            "Expected a Range's dimension should to be an Integer value, but "
+            "encountered " + dim0.name
+        )
+
+    for ty_dim in ty_optional_dims:
+        if isinstance(ty_dim, types.Integer):
+            ndims += 1
+        elif ty_dim is not None:
+            raise errors.TypingError(
+                "Expected a Range's dimension to be an Integer value, "
+                f"but {type(ty_dim)} was provided."
+            )
+
+    ret_ty = RangeType(ndims)
+
+    def impl(dim0, dim1=None, dim2=None):
+        return _intrin_range_alloc(dim0, dim1, dim2, ret_ty)
+
+    return impl
+
+
+@overload(NdRange, target=DPEX_TARGET_NAME)
+def _ol_ndrange_init(global_range, local_range):
+    """Numba overload of the NdRange constructor to make it usable inside an
+    njit and dpjit decorated function.
+
+    """
+    from numba_dpex.core.exceptions import UnmatchedNumberOfRangeDimsError
+    from numba_dpex.core.types import NdRangeType, RangeType
+
+    if not isinstance(global_range, RangeType):
+        raise errors.TypingError(
+            "Only global range values specified as a Range are "
+            "supported inside dpjit"
+        )
+
+    if not isinstance(local_range, RangeType):
+        raise errors.TypingError(
+            "Only local range values specified as a Range are "
+            "supported inside dpjit"
+        )
+
+    if not global_range.ndim == local_range.ndim:
+        raise UnmatchedNumberOfRangeDimsError(
+            kernel_name="",
+            global_ndims=global_range.ndim,
+            local_ndims=local_range.ndim,
+        )
+
+    ret_ty = NdRangeType(global_range.ndim)
+
+    def impl(global_range, local_range):
+        return _intrin_ndrange_alloc(global_range, local_range, ret_ty)
+
+    return impl
diff --git a/numba_dpex/core/kernel_interface/spirv_kernel.py b/numba_dpex/core/kernel_interface/spirv_kernel.py
index a7008c6528..9e6d240fcb 100644
--- a/numba_dpex/core/kernel_interface/spirv_kernel.py
+++ b/numba_dpex/core/kernel_interface/spirv_kernel.py
@@ -7,11 +7,11 @@
 
 from numba.core import ir
 
-from numba_dpex import spirv_generator
 from numba_dpex.core import config
 from numba_dpex.core.compiler import compile_with_dpex
 from numba_dpex.core.exceptions import UncompiledKernelError, UnreachableError
-from numba_dpex.core.targets.kernel_target import DpexKernelTargetContext
+from numba_dpex.kernel_api_impl.spirv import spirv_generator
+from numba_dpex.kernel_api_impl.spirv.target import SPIRVTargetContext
 
 from .kernel_base import KernelInterface
 
@@ -135,7 +135,7 @@ def compile(
         )
 
         func = cres.library.get_function(cres.fndesc.llvm_func_name)
-        kernel_targetctx: DpexKernelTargetContext = cres.target_context
+        kernel_targetctx: SPIRVTargetContext = cres.target_context
         kernel = kernel_targetctx.prepare_spir_kernel(func, cres.signature.args)
 
         # XXX: Setting the inline_threshold in the following way is a temporary
diff --git a/numba_dpex/core/parfors/reduction_helper.py b/numba_dpex/core/parfors/reduction_helper.py
index 047eeb2b11..8efbbe32bb 100644
--- a/numba_dpex/core/parfors/reduction_helper.py
+++ b/numba_dpex/core/parfors/reduction_helper.py
@@ -285,9 +285,9 @@ def __init__(
         self._parfor_redvars_to_redarrs = {}
         for ele1 in reductionHelperList:
             for key in ele1.redvars_to_redarrs_dict.keys():
-                self._parfor_redvars_to_redarrs[
-                    key
-                ] = ele1.redvars_to_redarrs_dict[key]
+                self._parfor_redvars_to_redarrs[key] = (
+                    ele1.redvars_to_redarrs_dict[key]
+                )
                 tmp1.append(ele1.redvars_to_redarrs_dict[key][0])
                 tmp1.append(ele1.redvars_to_redarrs_dict[key][1])
             self._final_sum_names.append(ele1.final_sum_var.name)
diff --git a/numba_dpex/core/passes/parfor_legalize_cfd_pass.py b/numba_dpex/core/passes/parfor_legalize_cfd_pass.py
index b7b9402b66..bc37d01210 100644
--- a/numba_dpex/core/passes/parfor_legalize_cfd_pass.py
+++ b/numba_dpex/core/passes/parfor_legalize_cfd_pass.py
@@ -18,7 +18,6 @@
 
 
 class ParforLegalizeCFDPassImpl:
-
     """Legalizes the compute-follows-data based device attribute for parfor
     nodes.
 
diff --git a/numba_dpex/core/pipelines/dpjit_compiler.py b/numba_dpex/core/pipelines/dpjit_compiler.py
index d4c63ee056..3e8fcfb4fd 100644
--- a/numba_dpex/core/pipelines/dpjit_compiler.py
+++ b/numba_dpex/core/pipelines/dpjit_compiler.py
@@ -130,9 +130,9 @@ class DpjitCompiler(CompilerBase):
     def define_pipelines(self):
         pms = []
         self.state.parfor_diagnostics = ExtendedParforDiagnostics()
-        self.state.metadata[
-            "parfor_diagnostics"
-        ] = self.state.parfor_diagnostics
+        self.state.metadata["parfor_diagnostics"] = (
+            self.state.parfor_diagnostics
+        )
         if not self.state.flags.force_pyobject:
             pms.append(self._pass_builder.define_nopython_pipeline(self.state))
         if self.state.status.can_fallback or self.state.flags.force_pyobject:
diff --git a/numba_dpex/core/types/__init__.py b/numba_dpex/core/types/__init__.py
index 398476797b..54b5f24bd4 100644
--- a/numba_dpex/core/types/__init__.py
+++ b/numba_dpex/core/types/__init__.py
@@ -5,6 +5,8 @@
 from .array_type import Array
 from .dpctl_types import DpctlSyclEvent, DpctlSyclQueue
 from .dpnp_ndarray_type import DpnpNdArray
+from .kernel_api.literal_intenum import IntEnumLiteral
+from .kernel_api.ranges import NdRangeType, RangeType
 from .numba_types_short_names import (
     b1,
     bool_,
@@ -26,7 +28,6 @@
     uint64,
     void,
 )
-from .range_types import NdRangeType, RangeType
 from .usm_ndarray_type import USMNdArray
 
 usm_ndarray = USMNdArray
@@ -36,8 +37,9 @@
     "DpctlSyclQueue",
     "DpctlSyclEvent",
     "DpnpNdArray",
-    "RangeType",
+    "IntEnumLiteral",
     "NdRangeType",
+    "RangeType",
     "USMNdArray",
     "none",
     "boolean",
@@ -57,6 +59,6 @@
     "f8",
     "float_",
     "double",
-    "void",
     "usm_ndarray",
+    "void",
 ]
diff --git a/numba_dpex/core/types/dpnp_ndarray_type.py b/numba_dpex/core/types/dpnp_ndarray_type.py
index 34aefb653f..fb09598243 100644
--- a/numba_dpex/core/types/dpnp_ndarray_type.py
+++ b/numba_dpex/core/types/dpnp_ndarray_type.py
@@ -3,14 +3,9 @@
 # SPDX-License-Identifier: Apache-2.0
 
 import dpnp
-from numba.core import cgutils, ir, types
-from numba.core.errors import NumbaNotImplementedError
+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 numba.np import numpy_support
-
-from numba_dpex.core.exceptions import UnreachableError
-from numba_dpex.core.runtime import context as dpexrt
 
 from .usm_ndarray_type import USMNdArray
 
@@ -197,117 +192,3 @@ def __allocate__(
             out.extend([g_np_assign, attr_assign, typ_var_assign, alloc_assign])
 
         return out
-
-
-# TODO: move this section to separate file
-# --------------- Boxing/Unboxing logic for dpnp.ndarray ----------------------#
-
-
-@unbox(USMNdArray)
-def unbox_dpnp_nd_array(typ, obj, c):
-    """Converts a dpctl.tensor.usm_ndarray/dpnp.ndarray object to a Numba-dpex
-    internal array structure.
-
-    Args:
-        typ : The Numba type of the PyObject
-        obj : The actual PyObject to be unboxed
-        c : The unboxing context
-
-    Returns: A NativeValue object representing an unboxed
-        dpctl.tensor.usm_ndarray/dpnp.ndarray
-    """
-    # Reusing the numba.core.base.BaseContext's make_array function to get a
-    # struct allocated. The same struct is used for numpy.ndarray
-    # and dpnp.ndarray. It is possible to do so, as the extra information
-    # specific to dpnp.ndarray such as sycl_queue is inferred statically and
-    # stored as part of the DpnpNdArray type.
-
-    # --------------- Original Numba comment from @ubox(types.Array)
-    #
-    # This is necessary because unbox_buffer() does not work on some
-    # dtypes, e.g. datetime64 and timedelta64.
-    # TODO check matching dtype.
-    #      currently, mismatching dtype will still work and causes
-    #      potential memory corruption
-    #
-    # --------------- End of Numba comment from @ubox(types.Array)
-    nativearycls = c.context.make_array(typ)
-    nativeary = nativearycls(c.context, c.builder)
-    aryptr = nativeary._getpointer()
-
-    ptr = c.builder.bitcast(aryptr, c.pyapi.voidptr)
-    # FIXME : We need to check if Numba_RT as well as DPEX RT are enabled.
-    if c.context.enable_nrt:
-        dpexrtCtx = dpexrt.DpexRTContext(c.context)
-        errcode = dpexrtCtx.arraystruct_from_python(c.pyapi, obj, ptr)
-    else:
-        raise UnreachableError
-
-    # TODO: here we have minimal typechecking by the itemsize.
-    #       need to do better
-    try:
-        expected_itemsize = numpy_support.as_dtype(typ.dtype).itemsize
-    except NumbaNotImplementedError:
-        # Don't check types that can't be `as_dtype()`-ed
-        itemsize_mismatch = cgutils.false_bit
-    else:
-        expected_itemsize = nativeary.itemsize.type(expected_itemsize)
-        itemsize_mismatch = c.builder.icmp_unsigned(
-            "!=",
-            nativeary.itemsize,
-            expected_itemsize,
-        )
-
-    failed = c.builder.or_(
-        cgutils.is_not_null(c.builder, errcode),
-        itemsize_mismatch,
-    )
-    # Handle error
-    with c.builder.if_then(failed, likely=False):
-        c.pyapi.err_set_string(
-            "PyExc_TypeError",
-            "can't unbox usm array from PyObject into "
-            "native value.  The object maybe of a "
-            "different type",
-        )
-    return NativeValue(c.builder.load(aryptr), is_error=failed)
-
-
-@box(USMNdArray)
-def box_array(typ, val, c):
-    """Boxes a NativeValue representation of USMNdArray/DpnpNdArray type into a
-    dpctl.tensor.usm_ndarray/dpnp.ndarray PyObject
-
-    Args:
-        typ: The representation of the USMNdArray/DpnpNdArray type.
-        val: A native representation of a Numba USMNdArray/DpnpNdArray type
-            object.
-        c: The boxing context.
-
-    Returns: A Pyobject for a dpctl.tensor.usm_ndarray/dpnp.ndarray boxed from
-        the Numba-dpex native value.
-    """
-    if c.context.enable_nrt:
-        np_dtype = numpy_support.as_dtype(typ.dtype)
-        dtypeptr = c.env_manager.read_const(c.env_manager.add_const(np_dtype))
-        dpexrtCtx = dpexrt.DpexRTContext(c.context)
-        newary = dpexrtCtx.usm_ndarray_to_python_acqref(
-            c.pyapi, typ, val, dtypeptr
-        )
-
-        if not newary:
-            c.pyapi.err_set_string(
-                "PyExc_TypeError",
-                "could not box native array into a dpnp.ndarray PyObject.",
-            )
-
-        # Steals NRT ref
-        # Refer:
-        #   numba.core.base.nrt -> numba.core.runtime.context -> decref
-        #   The `NRT_decref` function is generated directly as LLVM IR inside
-        #   numba.core.runtime.nrtdynmod.py
-        c.context.nrt.decref(c.builder, typ, val)
-
-        return newary
-    else:
-        raise UnreachableError
diff --git a/numba_dpex/tests/experimental/IntEnumLiteral/__init__.py b/numba_dpex/core/types/kernel_api/__init__.py
similarity index 100%
rename from numba_dpex/tests/experimental/IntEnumLiteral/__init__.py
rename to numba_dpex/core/types/kernel_api/__init__.py
diff --git a/numba_dpex/experimental/dpcpp_types.py b/numba_dpex/core/types/kernel_api/atomic_ref.py
similarity index 89%
rename from numba_dpex/experimental/dpcpp_types.py
rename to numba_dpex/core/types/kernel_api/atomic_ref.py
index 6cc5961f24..cdc2eb1890 100644
--- a/numba_dpex/experimental/dpcpp_types.py
+++ b/numba_dpex/core/types/kernel_api/atomic_ref.py
@@ -2,7 +2,7 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
-"""Collection of numba-dpex typing classes for kernel_iface Python classes.
+"""Collection of numba-dpex typing classes for kernel_api Python classes.
 """
 
 from numba.core.types import Type
@@ -10,7 +10,7 @@
 
 class AtomicRefType(Type):
     """numba-dpex internal type to represent a Python object of
-    :class:`numba_dpex.experimental.kernel_iface.AtomicRef`.
+    :class:`numba_dpex.kernel_api.AtomicRef`.
     """
 
     def __init__(
@@ -38,21 +38,21 @@ def __init__(
     @property
     def memory_order(self) -> int:
         """Returns the integer value for a memory order that corresponds to
-        kernel_iface.MemoryOrder.
+        kernel_api.MemoryOrder.
         """
         return self._memory_order
 
     @property
     def memory_scope(self) -> int:
         """Returns the integer value for a memory order that corresponds to
-        kernel_iface.MemoryScope.
+        kernel_api.MemoryScope.
         """
         return self._memory_scope
 
     @property
     def address_space(self) -> int:
         """Returns the integer value for a memory order that corresponds to
-        kernel_iface.AddressSpace.
+        kernel_api.AddressSpace.
         """
         return self._address_space
 
diff --git a/numba_dpex/core/types/kernel_api/index_space_ids.py b/numba_dpex/core/types/kernel_api/index_space_ids.py
new file mode 100644
index 0000000000..798aceeb2e
--- /dev/null
+++ b/numba_dpex/core/types/kernel_api/index_space_ids.py
@@ -0,0 +1,82 @@
+# SPDX-FileCopyrightText: 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Defines numba types for Item and NdItem classes"""
+
+from numba.core import errors, types
+
+
+class GroupType(types.Type):
+    """Numba-dpex type corresponding to :class:`numba_dpex.kernel_api.Group`"""
+
+    def __init__(self, ndim: int):
+        self._ndim = ndim
+        if ndim < 1 or ndim > 3:
+            raise errors.TypingError(
+                "ItemType can only have 1, 2 or 3 dimensions"
+            )
+        super().__init__(name="Group<" + str(ndim) + ">")
+
+    @property
+    def ndim(self):
+        """Returns number of dimensions"""
+        return self._ndim
+
+    @property
+    def key(self):
+        """Numba type specific overload"""
+        return self._ndim
+
+    def cast_python_value(self, args):
+        raise NotImplementedError
+
+
+class ItemType(types.Type):
+    """Numba-dpex type corresponding to :class:`numba_dpex.kernel_api.Item`"""
+
+    def __init__(self, ndim: int):
+        self._ndim = ndim
+        if ndim < 1 or ndim > 3:
+            raise errors.TypingError(
+                "ItemType can only have 1, 2 or 3 dimensions"
+            )
+        super().__init__(name="Item<" + str(ndim) + ">")
+
+    @property
+    def ndim(self):
+        """Returns number of dimensions"""
+        return self._ndim
+
+    @property
+    def key(self):
+        """Numba type specific overload"""
+        return self._ndim
+
+    def cast_python_value(self, args):
+        raise NotImplementedError
+
+
+class NdItemType(types.Type):
+    """Numba-dpex type corresponding to :class:`numba_dpex.kernel_api.NdItem`"""
+
+    def __init__(self, ndim: int):
+        self._ndim = ndim
+        if ndim < 1 or ndim > 3:
+            raise errors.TypingError(
+                "ItemType can only have 1, 2 or 3 dimensions"
+            )
+        super().__init__(name="NdItem<" + str(ndim) + ">")
+
+    @property
+    def ndim(self):
+        """Returns number of dimensions"""
+        return self._ndim
+
+    @property
+    def key(self):
+        """Numba type specific overload"""
+        return self._ndim
+
+    def cast_python_value(self, args):
+        raise NotImplementedError
diff --git a/numba_dpex/experimental/literal_intenum_type.py b/numba_dpex/core/types/kernel_api/literal_intenum.py
similarity index 96%
rename from numba_dpex/experimental/literal_intenum_type.py
rename to numba_dpex/core/types/kernel_api/literal_intenum.py
index 64769b45cb..d42fb0893f 100644
--- a/numba_dpex/experimental/literal_intenum_type.py
+++ b/numba_dpex/core/types/kernel_api/literal_intenum.py
@@ -13,7 +13,7 @@
 from numba.core.typing.typeof import typeof
 
 from numba_dpex.core.exceptions import IllegalIntEnumLiteralValueError
-from numba_dpex.experimental.flag_enum import FlagEnum
+from numba_dpex.kernel_api.flag_enum import FlagEnum
 
 
 class IntEnumLiteral(Literal, Integer):
diff --git a/numba_dpex/core/types/kernel_api/ranges.py b/numba_dpex/core/types/kernel_api/ranges.py
new file mode 100644
index 0000000000..af2f0125a7
--- /dev/null
+++ b/numba_dpex/core/types/kernel_api/ranges.py
@@ -0,0 +1,51 @@
+# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+from contextlib import ExitStack
+
+from numba.core import cgutils, errors, types
+
+
+class RangeType(types.Type):
+    """Numba-dpex type corresponding to
+    :class:`numba_dpex.kernel_api.ranges.Range`
+    """
+
+    def __init__(self, ndim: int):
+        self._ndim = ndim
+        if ndim < 1 or ndim > 3:
+            raise errors.TypingError(
+                "RangeType can only have 1,2, or 3 dimensions"
+            )
+        super(RangeType, self).__init__(name="Range<" + str(ndim) + ">")
+
+    @property
+    def ndim(self):
+        return self._ndim
+
+    @property
+    def key(self):
+        return self._ndim
+
+
+class NdRangeType(types.Type):
+    """Numba-dpex type corresponding to
+    :class:`numba_dpex.kernel_api.ranges.NdRange`
+    """
+
+    def __init__(self, ndim: int):
+        self._ndim = ndim
+        if ndim < 1 or ndim > 3:
+            raise errors.TypingError(
+                "RangeType can only have 1,2, or 3 dimensions"
+            )
+        super(NdRangeType, self).__init__(name="NdRange<" + str(ndim) + ">")
+
+    @property
+    def ndim(self):
+        return self._ndim
+
+    @property
+    def key(self):
+        return self._ndim
diff --git a/numba_dpex/core/typing/dpnpdecl.py b/numba_dpex/core/typing/dpnpdecl.py
index 355be22179..6856f8ea1f 100644
--- a/numba_dpex/core/typing/dpnpdecl.py
+++ b/numba_dpex/core/typing/dpnpdecl.py
@@ -179,12 +179,16 @@ def install_operations(cls):
         "ldexp",
         "spacing",
         "isnat",
+        "cbrt",
     ]
 )
 
-# A list of ufuncs that are in fact aliases of other ufuncs. They need to insert
-# the resolve method, but not register the ufunc itself
-_aliases = set(["bitwise_not", "mod", "abs"])
+# TODO: A list of ufuncs that are in fact aliases of other ufuncs. They need
+# to insert the resolve method, but not register the ufunc itself.
+# In a meantime let's just register them as user functions:
+# TODO: for some reason it affects "mod", but does not affect "bitwise_not" and
+# "abs". May be mod is not an alias?
+_aliases = {"bitwise_not", "abs"}
 
 all_ufuncs = sum(
     [
diff --git a/numba_dpex/core/typing/typeof.py b/numba_dpex/core/typing/typeof.py
index cce8cc78ba..78bbe10805 100644
--- a/numba_dpex/core/typing/typeof.py
+++ b/numba_dpex/core/typing/typeof.py
@@ -8,12 +8,12 @@
 from numba.extending import typeof_impl
 from numba.np import numpy_support
 
+from numba_dpex.kernel_api.ranges import NdRange, Range
 from numba_dpex.utils.constants import address_space
 
-from ..kernel_interface.indexers import NdRange, Range
 from ..types.dpctl_types import DpctlSyclEvent, DpctlSyclQueue
 from ..types.dpnp_ndarray_type import DpnpNdArray
-from ..types.range_types import NdRangeType, RangeType
+from ..types.kernel_api.ranges import NdRangeType, RangeType
 from ..types.usm_ndarray_type import USMNdArray
 
 
diff --git a/numba_dpex/core/utils/kernel_launcher.py b/numba_dpex/core/utils/kernel_launcher.py
index 4acc4da8b8..98814578ee 100644
--- a/numba_dpex/core/utils/kernel_launcher.py
+++ b/numba_dpex/core/utils/kernel_launcher.py
@@ -21,7 +21,7 @@
 from numba_dpex.core.exceptions import UnreachableError
 from numba_dpex.core.runtime.context import DpexRTContext
 from numba_dpex.core.types import USMNdArray
-from numba_dpex.core.types.range_types import NdRangeType, RangeType
+from numba_dpex.core.types.kernel_api.ranges import NdRangeType, RangeType
 from numba_dpex.dpctl_iface import libsyclinterface_bindings as sycl
 from numba_dpex.dpctl_iface._helpers import numba_type_to_dpctl_typenum
 from numba_dpex.utils import create_null_ptr
@@ -431,9 +431,11 @@ def _create_sycl_range(self, idx_range):
 
         """
         int64_range = [
-            self.builder.sext(rext, utils.LLVMTypes.int64_t)
-            if rext.type != utils.LLVMTypes.int64_t
-            else rext
+            (
+                self.builder.sext(rext, utils.LLVMTypes.int64_t)
+                if rext.type != utils.LLVMTypes.int64_t
+                else rext
+            )
             for rext in idx_range
         ]
 
diff --git a/numba_dpex/dpnp_iface/arrayobj.py b/numba_dpex/dpnp_iface/arrayobj.py
index 2e0e02e2a3..65ac816f89 100644
--- a/numba_dpex/dpnp_iface/arrayobj.py
+++ b/numba_dpex/dpnp_iface/arrayobj.py
@@ -19,8 +19,8 @@
 from numba_dpex.core.kernel_interface.arrayobj import (
     _getitem_array_generic as kernel_getitem_array_generic,
 )
-from numba_dpex.core.targets.kernel_target import DpexKernelTargetContext
 from numba_dpex.core.types import DpnpNdArray
+from numba_dpex.kernel_api_impl.spirv.target import SPIRVTargetContext
 
 from ._intrinsic import (
     impl_dpnp_empty,
@@ -1082,7 +1082,7 @@ def getitem_arraynd_intp(context, builder, sig, args):
     that when returning a view of a dpnp.ndarray the sycl::queue pointer
     member in the LLVM IR struct gets properly updated.
     """
-    getitem_call_in_kernel = isinstance(context, DpexKernelTargetContext)
+    getitem_call_in_kernel = isinstance(context, SPIRVTargetContext)
     _getitem_array_generic = np_getitem_array_generic
 
     if getitem_call_in_kernel:
diff --git a/numba_dpex/examples/vectorize.py b/numba_dpex/examples/vectorize.py
deleted file mode 100644
index 6cc3aa8db9..0000000000
--- a/numba_dpex/examples/vectorize.py
+++ /dev/null
@@ -1,66 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-from numba import float64, vectorize
-
-
-@vectorize(nopython=True)
-def ufunc_kernel(x, y):
-    return x + y
-
-
-def get_device():
-    device = None
-    try:
-        device = dpctl.select_gpu_device()
-    except:
-        try:
-            device = dpctl.select_cpu_device()
-        except:
-            raise RuntimeError("No device found")
-    return device
-
-
-def test_njit():
-    N = 10
-    dtype = np.float64
-
-    A = np.arange(N, dtype=dtype)
-    B = np.arange(N, dtype=dtype) * 10
-
-    # Use the environment variable SYCL_DEVICE_FILTER to change the default device.
-    # See https://github.com/intel/llvm/blob/sycl/sycl/doc/EnvironmentVariables.md#sycl_device_filter.
-    device = dpctl.select_default_device()
-    print("Using device ...")
-    device.print_device_info()
-
-    with dpctl.device_context(device):
-        C = ufunc_kernel(A, B)
-
-    print(C)
-
-    print("Done...")
-
-
-@vectorize([float64(float64, float64)], target="dpex")
-def vector_add(a, b):
-    return a + b
-
-
-def test_vectorize():
-    A = np.arange(10, dtype=np.float64).reshape((5, 2))
-    B = np.arange(10, dtype=np.float64).reshape((5, 2))
-
-    device = dpctl.select_default_device()
-    with dpctl.device_context(device):
-        C = vector_add(A, B)
-
-    print(C)
-
-
-if __name__ == "__main__":
-    test_njit()
-    test_vectorize()
diff --git a/numba_dpex/experimental/__init__.py b/numba_dpex/experimental/__init__.py
index b860e215a6..918987c87f 100644
--- a/numba_dpex/experimental/__init__.py
+++ b/numba_dpex/experimental/__init__.py
@@ -8,11 +8,17 @@
 
 from numba.core.imputils import Registry
 
-from ._kernel_dpcpp_spirv_overloads import _atomic_ref_overloads
+# Temporary so that Range and NdRange work in experimental call_kernel
+from numba_dpex.core.boxing import *
+from numba_dpex.kernel_api_impl.spirv.dispatcher import SPIRVKernelDispatcher
+
+from ._kernel_dpcpp_spirv_overloads import (
+    _atomic_ref_overloads,
+    _group_barrier_overloads,
+    _index_space_id_overloads,
+)
 from .decorators import device_func, kernel
-from .kernel_dispatcher import KernelDispatcher
 from .launcher import call_kernel, call_kernel_async
-from .literal_intenum_type import IntEnumLiteral
 from .models import *
 from .types import KernelDispatcherType
 
@@ -33,6 +39,5 @@ def dpex_dispatcher_const(context):
     "kernel",
     "call_kernel",
     "call_kernel_async",
-    "IntEnumLiteral",
-    "KernelDispatcher",
+    "SPIRVKernelDispatcher",
 ]
diff --git a/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_ref_overloads.py b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_ref_overloads.py
index c6074c5749..f2289d52ee 100644
--- a/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_ref_overloads.py
+++ b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_atomic_ref_overloads.py
@@ -3,7 +3,7 @@
 # SPDX-License-Identifier: Apache-2.0
 
 """
-Implements the SPIR-V overloads for the kernel_api.atomic_ref class methods.
+Implements the SPIR-V overloads for the kernel_api.AtomicRef class methods.
 """
 
 from llvmlite import ir as llvmir
@@ -12,18 +12,33 @@
 from numba.extending import intrinsic, overload, overload_method
 
 from numba_dpex.core import itanium_mangler as ext_itanium_mangler
-from numba_dpex.core.targets.kernel_target import CC_SPIR_FUNC, LLVM_SPIRV_ARGS
 from numba_dpex.core.types import USMNdArray
-from numba_dpex.experimental.flag_enum import FlagEnum
+from numba_dpex.kernel_api import (
+    AddressSpace,
+    AtomicRef,
+    MemoryOrder,
+    MemoryScope,
+)
+from numba_dpex.kernel_api.flag_enum import FlagEnum
+from numba_dpex.kernel_api_impl.spirv.target import (
+    CC_SPIR_FUNC,
+    LLVM_SPIRV_ARGS,
+)
 
-from ..dpcpp_types import AtomicRefType
-from ..kernel_iface import AddressSpace, AtomicRef, MemoryOrder, MemoryScope
+from ...core.types.kernel_api.atomic_ref import AtomicRefType
 from ..target import DPEX_KERNEL_EXP_TARGET_NAME
 from ._spv_atomic_inst_helper import (
     get_atomic_inst_name,
     get_memory_semantics_mask,
     get_scope,
 )
+from .spv_atomic_fn_declarations import (
+    _SUPPORT_CONVERGENT,
+    get_or_insert_atomic_load_fn,
+    get_or_insert_spv_atomic_compare_exchange_fn,
+    get_or_insert_spv_atomic_exchange_fn,
+    get_or_insert_spv_atomic_store_fn,
+)
 
 
 def _parse_enum_or_int_literal_(literal_int) -> int:
@@ -53,7 +68,7 @@ def _parse_enum_or_int_literal_(literal_int) -> int:
 def _intrinsic_helper(
     ty_context, ty_atomic_ref, ty_val, op_str  # pylint: disable=unused-argument
 ):
-    sig = types.void(ty_atomic_ref, ty_val)
+    sig = ty_atomic_ref.dtype(ty_atomic_ref, ty_val)
 
     def gen(context, builder, sig, args):
         atomic_ref_ty = sig.args[0]
@@ -100,19 +115,29 @@ def gen(context, builder, sig, args):
             mangled_fn_name,
         )
         func.calling_convention = CC_SPIR_FUNC
-        spirv_memory_semantics_mask = get_memory_semantics_mask(
-            atomic_ref_ty.memory_order
-        )
-        spirv_scope = get_scope(atomic_ref_ty.memory_scope)
+        if _SUPPORT_CONVERGENT:
+            func.attributes.add("convergent")
+        func.attributes.add("nounwind")
 
         fn_args = [
             builder.extract_value(args[0], data_attr_pos),
-            context.get_constant(types.int32, spirv_scope),
-            context.get_constant(types.int32, spirv_memory_semantics_mask),
+            context.get_constant(
+                types.int32, get_scope(atomic_ref_ty.memory_scope)
+            ),
+            context.get_constant(
+                types.int32,
+                get_memory_semantics_mask(atomic_ref_ty.memory_order),
+            ),
             args[1],
         ]
 
-        builder.call(func, fn_args)
+        callinst = builder.call(func, fn_args)
+
+        if _SUPPORT_CONVERGENT:
+            callinst.attributes.add("convergent")
+        callinst.attributes.add("nounwind")
+
+        return callinst
 
     return sig, gen
 
@@ -131,7 +156,7 @@ def gen(context, builder, sig, args):
         args_lst[1] = builder.fneg(args[1])
         args = tuple(args_lst)
 
-        gen_fn(context, builder, sig, args)
+        return gen_fn(context, builder, sig, args)
 
     return gen
 
@@ -209,6 +234,299 @@ def codegen(context, builder, sig, args):
     )
 
 
+@intrinsic(target=DPEX_KERNEL_EXP_TARGET_NAME)
+def _intrinsic_load(
+    ty_context, ty_atomic_ref  # pylint: disable=unused-argument
+):
+    sig = ty_atomic_ref.dtype(ty_atomic_ref)
+
+    def _intrinsic_load_gen(context, builder, sig, args):
+        atomic_ref_ty = sig.args[0]
+
+        atomic_ref_ptr = builder.extract_value(
+            args[0],
+            context.data_model_manager.lookup(atomic_ref_ty).get_field_position(
+                "ref"
+            ),
+        )
+        if sig.args[0].dtype == types.float32:
+            atomic_ref_ptr = builder.bitcast(
+                atomic_ref_ptr,
+                llvmir.PointerType(
+                    llvmir.IntType(32), addrspace=sig.args[0].address_space
+                ),
+            )
+        elif sig.args[0].dtype == types.float64:
+            atomic_ref_ptr = builder.bitcast(
+                atomic_ref_ptr,
+                llvmir.PointerType(
+                    llvmir.IntType(64), addrspace=sig.args[0].address_space
+                ),
+            )
+
+        fn_args = [
+            atomic_ref_ptr,
+            context.get_constant(
+                types.int32, get_scope(atomic_ref_ty.memory_scope)
+            ),
+            context.get_constant(
+                types.int32,
+                get_memory_semantics_mask(atomic_ref_ty.memory_order),
+            ),
+        ]
+
+        ret_val = builder.call(
+            get_or_insert_atomic_load_fn(
+                context, builder.module, atomic_ref_ty
+            ),
+            fn_args,
+        )
+
+        if _SUPPORT_CONVERGENT:
+            ret_val.attributes.add("convergent")
+        ret_val.attributes.add("nounwind")
+
+        if sig.args[0].dtype == types.float32:
+            ret_val = builder.bitcast(ret_val, llvmir.FloatType())
+        elif sig.args[0].dtype == types.float64:
+            ret_val = builder.bitcast(ret_val, llvmir.DoubleType())
+
+        return ret_val
+
+    return sig, _intrinsic_load_gen
+
+
+@intrinsic(target=DPEX_KERNEL_EXP_TARGET_NAME)
+def _intrinsic_store(
+    ty_context, ty_atomic_ref, ty_val
+):  # pylint: disable=unused-argument
+    sig = types.void(ty_atomic_ref, ty_val)
+
+    def _intrinsic_store_gen(context, builder, sig, args):
+        atomic_ref_ty = sig.args[0]
+
+        store_arg = args[1]
+        atomic_ref_ptr = builder.extract_value(
+            args[0],
+            context.data_model_manager.lookup(atomic_ref_ty).get_field_position(
+                "ref"
+            ),
+        )
+        if sig.args[0].dtype == types.float32:
+            atomic_ref_ptr = builder.bitcast(
+                atomic_ref_ptr,
+                llvmir.PointerType(
+                    llvmir.IntType(32), addrspace=sig.args[0].address_space
+                ),
+            )
+            store_arg = builder.bitcast(store_arg, llvmir.IntType(32))
+        elif sig.args[0].dtype == types.float64:
+            atomic_ref_ptr = builder.bitcast(
+                atomic_ref_ptr,
+                llvmir.PointerType(
+                    llvmir.IntType(64), addrspace=sig.args[0].address_space
+                ),
+            )
+            store_arg = builder.bitcast(store_arg, llvmir.IntType(64))
+
+        atomic_store_fn_args = [
+            atomic_ref_ptr,
+            context.get_constant(
+                types.int32, get_scope(atomic_ref_ty.memory_scope)
+            ),
+            context.get_constant(
+                types.int32,
+                get_memory_semantics_mask(atomic_ref_ty.memory_order),
+            ),
+            store_arg,
+        ]
+
+        callinst = builder.call(
+            get_or_insert_spv_atomic_store_fn(
+                context, builder.module, atomic_ref_ty
+            ),
+            atomic_store_fn_args,
+        )
+
+        if _SUPPORT_CONVERGENT:
+            callinst.attributes.add("convergent")
+        callinst.attributes.add("nounwind")
+
+    return sig, _intrinsic_store_gen
+
+
+@intrinsic(target=DPEX_KERNEL_EXP_TARGET_NAME)
+def _intrinsic_exchange(
+    ty_context, ty_atomic_ref, ty_val  # pylint: disable=unused-argument
+):
+    sig = ty_atomic_ref.dtype(ty_atomic_ref, ty_val)
+
+    def _intrinsic_exchange_gen(context, builder, sig, args):
+        atomic_ref_ty = sig.args[0]
+
+        exchange_arg = args[1]
+        atomic_ref_ptr = builder.extract_value(
+            args[0],
+            context.data_model_manager.lookup(atomic_ref_ty).get_field_position(
+                "ref"
+            ),
+        )
+        if sig.args[0].dtype == types.float32:
+            atomic_ref_ptr = builder.bitcast(
+                atomic_ref_ptr,
+                llvmir.PointerType(
+                    llvmir.IntType(32), addrspace=sig.args[0].address_space
+                ),
+            )
+            exchange_arg = builder.bitcast(exchange_arg, llvmir.IntType(32))
+        elif sig.args[0].dtype == types.float64:
+            atomic_ref_ptr = builder.bitcast(
+                atomic_ref_ptr,
+                llvmir.PointerType(
+                    llvmir.IntType(64), addrspace=sig.args[0].address_space
+                ),
+            )
+            exchange_arg = builder.bitcast(exchange_arg, llvmir.IntType(64))
+
+        atomic_exchange_fn_args = [
+            atomic_ref_ptr,
+            context.get_constant(
+                types.int32, get_scope(atomic_ref_ty.memory_scope)
+            ),
+            context.get_constant(
+                types.int32,
+                get_memory_semantics_mask(atomic_ref_ty.memory_order),
+            ),
+            exchange_arg,
+        ]
+
+        ret_val = builder.call(
+            get_or_insert_spv_atomic_exchange_fn(
+                context, builder.module, atomic_ref_ty
+            ),
+            atomic_exchange_fn_args,
+        )
+
+        if _SUPPORT_CONVERGENT:
+            ret_val.attributes.add("convergent")
+        ret_val.attributes.add("nounwind")
+
+        if sig.args[0].dtype == types.float32:
+            ret_val = builder.bitcast(ret_val, llvmir.FloatType())
+        elif sig.args[0].dtype == types.float64:
+            ret_val = builder.bitcast(ret_val, llvmir.DoubleType())
+
+        return ret_val
+
+    return sig, _intrinsic_exchange_gen
+
+
+@intrinsic(target=DPEX_KERNEL_EXP_TARGET_NAME)
+def _intrinsic_compare_exchange(
+    ty_context,  # pylint: disable=unused-argument
+    ty_atomic_ref,
+    ty_expected_ref,
+    ty_desired,
+    ty_expected_idx,
+):
+    sig = types.boolean(
+        ty_atomic_ref, ty_expected_ref, ty_desired, ty_expected_idx
+    )
+
+    def _intrinsic_compare_exchange_gen(context, builder, sig, args):
+        # get pointer to expected[expected_idx]
+        data_attr = builder.extract_value(
+            args[1],
+            context.data_model_manager.lookup(sig.args[1]).get_field_position(
+                "data"
+            ),
+        )
+        with builder.goto_entry_block():
+            ptr_to_data_attr = builder.alloca(data_attr.type)
+        builder.store(data_attr, ptr_to_data_attr)
+        expected_ref_ptr = builder.gep(
+            builder.load(ptr_to_data_attr), [args[3]]
+        )
+
+        expected_arg = builder.load(expected_ref_ptr)
+        desired_arg = args[2]
+        atomic_ref_ptr = builder.extract_value(
+            args[0],
+            context.data_model_manager.lookup(sig.args[0]).get_field_position(
+                "ref"
+            ),
+        )
+        # add conditional bitcast for atomic_ref pointer,
+        # expected[expected_idx], and desired
+        if sig.args[0].dtype == types.float32:
+            atomic_ref_ptr = builder.bitcast(
+                atomic_ref_ptr,
+                llvmir.PointerType(
+                    llvmir.IntType(32), addrspace=sig.args[0].address_space
+                ),
+            )
+            expected_arg = builder.bitcast(expected_arg, llvmir.IntType(32))
+            desired_arg = builder.bitcast(desired_arg, llvmir.IntType(32))
+        elif sig.args[0].dtype == types.float64:
+            atomic_ref_ptr = builder.bitcast(
+                atomic_ref_ptr,
+                llvmir.PointerType(
+                    llvmir.IntType(64), addrspace=sig.args[0].address_space
+                ),
+            )
+            expected_arg = builder.bitcast(expected_arg, llvmir.IntType(64))
+            desired_arg = builder.bitcast(desired_arg, llvmir.IntType(64))
+
+        atomic_cmpexchg_fn_args = [
+            atomic_ref_ptr,
+            context.get_constant(
+                types.int32, get_scope(sig.args[0].memory_scope)
+            ),
+            context.get_constant(
+                types.int32,
+                get_memory_semantics_mask(sig.args[0].memory_order),
+            ),
+            context.get_constant(
+                types.int32,
+                get_memory_semantics_mask(sig.args[0].memory_order),
+            ),
+            desired_arg,
+            expected_arg,
+        ]
+
+        ret_val = builder.call(
+            get_or_insert_spv_atomic_compare_exchange_fn(
+                context, builder.module, sig.args[0]
+            ),
+            atomic_cmpexchg_fn_args,
+        )
+
+        if _SUPPORT_CONVERGENT:
+            ret_val.attributes.add("convergent")
+        ret_val.attributes.add("nounwind")
+
+        # compare_exchange returns the old value stored in AtomicRef object.
+        # If the return value is same as expected, then compare_exchange
+        # succeeded in replacing AtomicRef object with desired.
+        # If the return value is not same as expected, then store return
+        # value in expected.
+        # In either case, return result of cmp instruction.
+        is_cmp_exchg_success = builder.icmp_signed("==", ret_val, expected_arg)
+
+        with builder.if_else(is_cmp_exchg_success) as (then, otherwise):
+            with then:
+                pass
+            with otherwise:
+                if sig.args[0].dtype == types.float32:
+                    ret_val = builder.bitcast(ret_val, llvmir.FloatType())
+                elif sig.args[0].dtype == types.float64:
+                    ret_val = builder.bitcast(ret_val, llvmir.DoubleType())
+                builder.store(ret_val, expected_ref_ptr)
+        return is_cmp_exchg_success
+
+    return sig, _intrinsic_compare_exchange_gen
+
+
 def _check_if_supported_ref(ref):
     supported = True
 
@@ -252,7 +570,7 @@ def ol_atomic_ref(
     address_space=AddressSpace.GLOBAL,
 ):
     """Overload of the constructor for the class
-    class:`numba_dpex.experimental.kernel_iface.AtomicRef`.
+    class:`numba_dpex.kernel_api.AtomicRef`.
 
     Raises:
         errors.TypingError: If the `ref` argument is not a UsmNdArray type.
@@ -330,8 +648,7 @@ def ol_atomic_ref_ctor_impl(
 
 @overload_method(AtomicRefType, "fetch_add", target=DPEX_KERNEL_EXP_TARGET_NAME)
 def ol_fetch_add(atomic_ref, val):
-    """SPIR-V overload for
-    :meth:`numba_dpex.experimental.kernel_iface.AtomicRef.fetch_add`.
+    """SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_add`.
 
     Generates the same LLVM IR instruction as dpcpp for the
     `atomic_ref::fetch_add` function.
@@ -355,8 +672,7 @@ def ol_fetch_add_impl(atomic_ref, val):
 
 @overload_method(AtomicRefType, "fetch_sub", target=DPEX_KERNEL_EXP_TARGET_NAME)
 def ol_fetch_sub(atomic_ref, val):
-    """SPIR-V overload for
-    :meth:`numba_dpex.experimental.kernel_iface.AtomicRef.fetch_sub`.
+    """SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_sub`.
 
     Generates the same LLVM IR instruction as dpcpp for the
     `atomic_ref::fetch_sub` function.
@@ -380,8 +696,7 @@ def ol_fetch_sub_impl(atomic_ref, val):
 
 @overload_method(AtomicRefType, "fetch_min", target=DPEX_KERNEL_EXP_TARGET_NAME)
 def ol_fetch_min(atomic_ref, val):
-    """SPIR-V overload for
-    :meth:`numba_dpex.experimental.kernel_iface.AtomicRef.fetch_min`.
+    """SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_min`.
 
     Generates the same LLVM IR instruction as dpcpp for the
     `atomic_ref::fetch_min` function.
@@ -405,8 +720,7 @@ def ol_fetch_min_impl(atomic_ref, val):
 
 @overload_method(AtomicRefType, "fetch_max", target=DPEX_KERNEL_EXP_TARGET_NAME)
 def ol_fetch_max(atomic_ref, val):
-    """SPIR-V overload for
-    :meth:`numba_dpex.experimental.kernel_iface.AtomicRef.fetch_max`.
+    """SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_max`.
 
     Generates the same LLVM IR instruction as dpcpp for the
     `atomic_ref::fetch_max` function.
@@ -430,8 +744,7 @@ def ol_fetch_max_impl(atomic_ref, val):
 
 @overload_method(AtomicRefType, "fetch_and", target=DPEX_KERNEL_EXP_TARGET_NAME)
 def ol_fetch_and(atomic_ref, val):
-    """SPIR-V overload for
-    :meth:`numba_dpex.experimental.kernel_iface.AtomicRef.fetch_and`.
+    """SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_and`.
 
     Generates the same LLVM IR instruction as dpcpp for the
     `atomic_ref::fetch_and` function.
@@ -460,8 +773,7 @@ def ol_fetch_and_impl(atomic_ref, val):
 
 @overload_method(AtomicRefType, "fetch_or", target=DPEX_KERNEL_EXP_TARGET_NAME)
 def ol_fetch_or(atomic_ref, val):
-    """SPIR-V overload for
-    :meth:`numba_dpex.experimental.kernel_iface.AtomicRef.fetch_or`.
+    """SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_or`.
 
     Generates the same LLVM IR instruction as dpcpp for the
     `atomic_ref::fetch_or` function.
@@ -490,8 +802,7 @@ def ol_fetch_or_impl(atomic_ref, val):
 
 @overload_method(AtomicRefType, "fetch_xor", target=DPEX_KERNEL_EXP_TARGET_NAME)
 def ol_fetch_xor(atomic_ref, val):
-    """SPIR-V overload for
-    :meth:`numba_dpex.experimental.kernel_iface.AtomicRef.fetch_xor`.
+    """SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.fetch_xor`.
 
     Generates the same LLVM IR instruction as dpcpp for the
     `atomic_ref::fetch_xor` function.
@@ -516,3 +827,116 @@ def ol_fetch_xor_impl(atomic_ref, val):
         return _intrinsic_fetch_xor(atomic_ref, val)
 
     return ol_fetch_xor_impl
+
+
+@overload_method(AtomicRefType, "load", target=DPEX_KERNEL_EXP_TARGET_NAME)
+def ol_load(atomic_ref):  # pylint: disable=unused-argument
+    """SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.load`.
+
+    Generates the same LLVM IR instruction as dpcpp for the
+    `atomic_ref::load` function.
+
+    """
+
+    def ol_load_impl(atomic_ref):
+        # pylint: disable=no-value-for-parameter
+        return _intrinsic_load(atomic_ref)
+
+    return ol_load_impl
+
+
+@overload_method(AtomicRefType, "store", target=DPEX_KERNEL_EXP_TARGET_NAME)
+def ol_store(atomic_ref, val):
+    """SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.store`.
+
+    Generates the same LLVM IR instruction as dpcpp for the
+    `atomic_ref::store` function.
+
+    Raises:
+        TypingError: When the dtype of the value stored does not match the
+        dtype of the AtomicRef type.
+    """
+
+    if atomic_ref.dtype != val:
+        raise errors.TypingError(
+            f"Type of value to store: {val} does not match the type of the "
+            f"reference: {atomic_ref.dtype} stored in the atomic ref."
+        )
+
+    def ol_store_impl(atomic_ref, val):
+        # pylint: disable=no-value-for-parameter
+        return _intrinsic_store(atomic_ref, val)
+
+    return ol_store_impl
+
+
+@overload_method(AtomicRefType, "exchange", target=DPEX_KERNEL_EXP_TARGET_NAME)
+def ol_exchange(atomic_ref, val):
+    """SPIR-V overload for :meth:`numba_dpex.kernel_api.AtomicRef.exchange`.
+
+    Generates the same LLVM IR instruction as dpcpp for the
+    `atomic_ref::exchange` function.
+
+    Raises:
+        TypingError: When the dtype of the value passed to `exchange`
+        does not match the dtype of the AtomicRef type.
+    """
+
+    if atomic_ref.dtype != val:
+        raise errors.TypingError(
+            f"Type of value to exchange: {val} does not match the type of the "
+            f"reference: {atomic_ref.dtype} stored in the atomic ref."
+        )
+
+    def ol_exchange_impl(atomic_ref, val):
+        # pylint: disable=no-value-for-parameter
+        return _intrinsic_exchange(atomic_ref, val)
+
+    return ol_exchange_impl
+
+
+@overload_method(
+    AtomicRefType,
+    "compare_exchange",
+    target=DPEX_KERNEL_EXP_TARGET_NAME,
+)
+def ol_compare_exchange(
+    atomic_ref,
+    expected_ref,
+    desired,
+    expected_idx=0,  # pylint: disable=unused-argument
+):
+    """SPIR-V overload for
+    :meth:`numba_dpex.experimental.kernel_iface.AtomicRef.compare_exchange`.
+
+    Generates the same LLVM IR instruction as dpcpp for the
+    `atomic_ref::compare_exchange_strong` function.
+
+    Raises:
+        TypingError: When the dtype of the value passed to `compare_exchange`
+        does not match the dtype of the AtomicRef type.
+    """
+
+    _check_if_supported_ref(expected_ref)
+
+    if atomic_ref.dtype != expected_ref.dtype:
+        raise errors.TypingError(
+            f"Type of value to compare_exchange: {expected_ref} does not match the "
+            f"type of the reference: {atomic_ref.dtype} stored in the atomic ref."
+        )
+
+    if atomic_ref.dtype != desired:
+        raise errors.TypingError(
+            f"Type of value to compare_exchange: {desired} does not match the "
+            f"type of the reference: {atomic_ref.dtype} stored in the atomic ref."
+        )
+
+    def ol_compare_exchange_impl(
+        atomic_ref, expected_ref, desired, expected_idx=0
+    ):
+        # pylint: disable=no-value-for-parameter
+        return _intrinsic_compare_exchange(
+            atomic_ref, expected_ref, desired, expected_idx
+        )
+
+    return ol_compare_exchange_impl
diff --git a/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads.py b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads.py
new file mode 100644
index 0000000000..85485127f2
--- /dev/null
+++ b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads.py
@@ -0,0 +1,142 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""
+Provides overloads for functions included in kernel_iface.barrier that
+generate dpcpp SPIR-V LLVM IR intrinsic function calls.
+"""
+
+from llvmlite import ir as llvmir
+from numba.core import cgutils, types
+from numba.core.errors import TypingError
+from numba.extending import intrinsic, overload
+
+from numba_dpex.core import itanium_mangler as ext_itanium_mangler
+from numba_dpex.core.types.kernel_api.index_space_ids import GroupType
+from numba_dpex.experimental.target import DPEX_KERNEL_EXP_TARGET_NAME
+from numba_dpex.kernel_api import group_barrier
+from numba_dpex.kernel_api.memory_enums import MemoryOrder, MemoryScope
+
+from ._spv_atomic_inst_helper import get_memory_semantics_mask, get_scope
+from .spv_atomic_fn_declarations import _SUPPORT_CONVERGENT
+
+
+def _get_memory_scope(fence_scope):
+    if isinstance(fence_scope, types.Literal):
+        return get_scope(fence_scope.literal_value)
+    return get_scope(fence_scope.value)
+
+
+@intrinsic
+def _intrinsic_barrier(
+    ty_context,  # pylint: disable=unused-argument
+    ty_exec_scope,  # pylint: disable=unused-argument
+    ty_mem_scope,  # pylint: disable=unused-argument
+    ty_spirv_mem_sem_mask,  # pylint: disable=unused-argument
+):
+    # Signature of `__spirv_control_barrier` call that is
+    # generated for group_barrier. It takes three arguments -
+    # exec_scope, memory_scope and memory_semantics_mask.
+    # All arguments have to be of type unsigned int32.
+    sig = types.void(types.uint32, types.uint32, types.uint32)
+
+    def _intrinsic_barrier_codegen(
+        context, builder, sig, args  # pylint: disable=unused-argument
+    ):
+        exec_scope_arg = builder.trunc(args[0], llvmir.IntType(32))
+        mem_scope_arg = builder.trunc(args[1], llvmir.IntType(32))
+        spirv_memory_semantics_mask_arg = builder.trunc(
+            args[2], llvmir.IntType(32)
+        )
+
+        fn_args = [
+            exec_scope_arg,
+            mem_scope_arg,
+            spirv_memory_semantics_mask_arg,
+        ]
+
+        mangled_fn_name = ext_itanium_mangler.mangle_ext(
+            "__spirv_ControlBarrier", [types.uint32, types.uint32, types.uint32]
+        )
+
+        spirv_fn_arg_types = [
+            llvmir.IntType(32),
+            llvmir.IntType(32),
+            llvmir.IntType(32),
+        ]
+
+        fn = cgutils.get_or_insert_function(
+            builder.module,
+            llvmir.FunctionType(llvmir.VoidType(), spirv_fn_arg_types),
+            mangled_fn_name,
+        )
+
+        if _SUPPORT_CONVERGENT:
+            fn.attributes.add("convergent")
+        fn.attributes.add("nounwind")
+        fn.calling_convention = "spir_func"
+
+        callinst = builder.call(fn, fn_args)
+
+        if _SUPPORT_CONVERGENT:
+            callinst.attributes.add("convergent")
+        callinst.attributes.add("nounwind")
+
+    return (
+        sig,
+        _intrinsic_barrier_codegen,
+    )
+
+
+@overload(
+    group_barrier,
+    prefer_literal=True,
+    target=DPEX_KERNEL_EXP_TARGET_NAME,
+)
+def ol_group_barrier(group, fence_scope=MemoryScope.WORK_GROUP):
+    """SPIR-V overload for
+    :meth:`numba_dpex.kernel_api.group_barrier`.
+
+    Generates the same LLVM IR instruction as DPC++ for the SYCL
+    `group_barrier` function.
+
+    Per SYCL spec, group_barrier must perform both control barrier and memory
+    fence operations. Hence, group_barrier requires two scopes and one memory
+    consistency specification as its three arguments.
+
+    mem_scope - scope of any memory consistency operations that are performed by
+                the barrier. By default, mem_scope is set to `work_group`.
+    exec_scope - scope that determines the set of work-items that synchronize at
+                 barrier. Set to `work_group` for group_barrier always.
+    spirv_memory_semantics_mask - Based on SYCL implementation. Always set to
+                                  use sequential consistency memory order.
+    """
+
+    if not isinstance(group, GroupType):
+        raise TypingError(
+            "Expected a group should to be a GroupType value, but "
+            f"encountered {type(group)}"
+        )
+
+    mem_scope = _get_memory_scope(fence_scope)
+    # TODO: exec_scope needs to be determined based on
+    # group argument. If group refers to a work_group then,
+    # exec_scope is MemoryScope.WORK_GROUP.
+    # If group is sub_group then, exec_scope needs to be
+    # MemoryScope.SUB_GROUP
+    exec_scope = get_scope(MemoryScope.WORK_GROUP.value)
+    spirv_memory_semantics_mask = get_memory_semantics_mask(
+        MemoryOrder.SEQ_CST.value
+    )
+
+    def _ol_group_barrier_impl(
+        group,
+        fence_scope=MemoryScope.WORK_GROUP,
+    ):  # pylint: disable=unused-argument
+        # pylint: disable=no-value-for-parameter
+        return _intrinsic_barrier(
+            exec_scope, mem_scope, spirv_memory_semantics_mask
+        )
+
+    return _ol_group_barrier_impl
diff --git a/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_index_space_id_overloads.py b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_index_space_id_overloads.py
new file mode 100644
index 0000000000..d574f6aced
--- /dev/null
+++ b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_index_space_id_overloads.py
@@ -0,0 +1,250 @@
+# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""
+Implements the SPIR-V overloads for the kernel_api.items class methods.
+"""
+
+import llvmlite.ir as llvmir
+from numba.core import cgutils, types
+from numba.core.errors import TypingError
+from numba.extending import intrinsic, overload_method
+
+from numba_dpex.core.types.kernel_api.index_space_ids import (
+    GroupType,
+    ItemType,
+    NdItemType,
+)
+from numba_dpex.kernel_api_impl.spirv.target import SPIRVTargetContext
+
+from ..target import DPEX_KERNEL_EXP_TARGET_NAME
+
+
+def spirv_name(name: str):
+    """Converts name to spirv name by adding __spirv_ prefix."""
+    return "__spirv_" + name
+
+
+def declare_spirv_const(
+    builder: llvmir.IRBuilder,
+    name: str,
+):
+    """Declares global external spirv constant"""
+    data = cgutils.add_global_variable(
+        builder.module,
+        llvmir.VectorType(llvmir.IntType(64), 3),
+        spirv_name(name),
+        addrspace=1,
+    )
+    data.linkage = "external"
+    data.global_constant = True
+    data.align = 32
+    data.storage_class = "dso_local local_unnamed_addr"
+    return data
+
+
+def _intrinsic_spirv_global_index_const(
+    ty_context,  # pylint: disable=unused-argument
+    ty_dim,  # pylint: disable=unused-argument
+    const_name: str,
+):
+    """Generates instruction to get spirv index from const_name."""
+    sig = types.int64(types.int32)
+
+    def _intrinsic_spirv_global_index_const_gen(
+        context: SPIRVTargetContext,
+        builder: llvmir.IRBuilder,
+        sig,  # pylint: disable=unused-argument
+        args,
+    ):
+        index_const = declare_spirv_const(
+            builder,
+            const_name,
+        )
+        [dim] = args
+        # TODO: llvmlite does not support gep on vector. Use this instead once
+        # supported.
+        # https://github.com/numba/llvmlite/issues/756
+        # res = builder.gep( # noqa: E800
+        #     global_invocation_id, # noqa: E800
+        #     [cgutils.int32_t(0), cgutils.int32_t(0)], # noqa: E800
+        #     inbounds=True, # noqa: E800
+        # ) # noqa: E800
+        # res = builder.load(res, align=32) # noqa: E800
+
+        res = builder.extract_element(
+            builder.load(index_const),
+            dim,
+        )
+
+        return context.cast(builder, res, types.uintp, types.intp)
+
+    return sig, _intrinsic_spirv_global_index_const_gen
+
+
+@intrinsic(target=DPEX_KERNEL_EXP_TARGET_NAME)
+def _intrinsic_spirv_global_invocation_id(
+    ty_context, ty_dim  # pylint: disable=unused-argument
+):
+    """Generates instruction to get index from BuiltInGlobalInvocationId."""
+    return _intrinsic_spirv_global_index_const(
+        ty_context, ty_dim, "BuiltInGlobalInvocationId"
+    )
+
+
+@intrinsic(target=DPEX_KERNEL_EXP_TARGET_NAME)
+def _intrinsic_spirv_local_invocation_id(
+    ty_context, ty_dim  # pylint: disable=unused-argument
+):
+    """Generates instruction to get index from BuiltInLocalInvocationId."""
+    return _intrinsic_spirv_global_index_const(
+        ty_context, ty_dim, "BuiltInLocalInvocationId"
+    )
+
+
+@intrinsic(target=DPEX_KERNEL_EXP_TARGET_NAME)
+def _intrinsic_spirv_global_size(
+    ty_context, ty_dim  # pylint: disable=unused-argument
+):
+    """Generates instruction to get index from BuiltInGlobalSize."""
+    return _intrinsic_spirv_global_index_const(
+        ty_context, ty_dim, "BuiltInGlobalSize"
+    )
+
+
+@intrinsic(target=DPEX_KERNEL_EXP_TARGET_NAME)
+def _intrinsic_spirv_workgroup_size(
+    ty_context, ty_dim  # pylint: disable=unused-argument
+):
+    """Generates instruction to get index from BuiltInWorkgroupSize."""
+    return _intrinsic_spirv_global_index_const(
+        ty_context, ty_dim, "BuiltInWorkgroupSize"
+    )
+
+
+@intrinsic(target=DPEX_KERNEL_EXP_TARGET_NAME)
+def _intrinsic_spirv_workgroup_id(
+    ty_context, ty_dim  # pylint: disable=unused-argument
+):
+    """Generates instruction to get index from BuiltInWorkgroupId."""
+    return _intrinsic_spirv_global_index_const(
+        ty_context, ty_dim, "BuiltInWorkgroupId"
+    )
+
+
+@intrinsic(target=DPEX_KERNEL_EXP_TARGET_NAME)
+def _intrinsic_spirv_numworkgroups(
+    ty_context, ty_dim  # pylint: disable=unused-argument
+):
+    """Generates instruction to get index from BuiltInNumWorkgroups."""
+    return _intrinsic_spirv_global_index_const(
+        ty_context, ty_dim, "BuiltInNumWorkgroups"
+    )
+
+
+def generate_index_overload(_type, _intrinsic):
+    """Generates overload for the index method that generates specific IR from
+    provided intrinsic."""
+
+    def ol_item_gen_index(item, dim):
+        """SPIR-V overload for :meth:`numba_dpex.kernel_api.<_type>.<method>`.
+
+        Generates the same LLVM IR instruction as dpcpp for the
+        `sycl::<type>::<method>` function.
+
+        Raises:
+            TypingError: When argument is not an integer.
+        """
+        if not isinstance(item, _type):
+            raise TypingError(
+                f"Expected an item should to be an {_type} value, but "
+                f"encountered {type(item)}"
+            )
+
+        if not isinstance(dim, types.Integer):
+            raise TypingError(
+                f"Expected an {_type}'s dim should to be an Integer value, but "
+                f"encountered {type(dim)}"
+            )
+
+        # pylint: disable=unused-argument
+        def ol_item_get_index_impl(item, dim):
+            # TODO: call in reverse index once index reversing is removed from
+            # kernel submission
+            # pylint: disable=no-value-for-parameter
+            return _intrinsic(dim)
+
+        return ol_item_get_index_impl
+
+    return ol_item_gen_index
+
+
+_index_const_overload_methods = [
+    (ItemType, "get_id", _intrinsic_spirv_global_invocation_id),
+    (ItemType, "get_range", _intrinsic_spirv_global_size),
+    (NdItemType, "get_global_id", _intrinsic_spirv_global_invocation_id),
+    (NdItemType, "get_local_id", _intrinsic_spirv_local_invocation_id),
+    (NdItemType, "get_global_range", _intrinsic_spirv_global_size),
+    (NdItemType, "get_local_range", _intrinsic_spirv_workgroup_size),
+    (GroupType, "get_group_id", _intrinsic_spirv_workgroup_id),
+    (GroupType, "get_group_range", _intrinsic_spirv_numworkgroups),
+    (GroupType, "get_local_range", _intrinsic_spirv_workgroup_size),
+]
+
+for index_overload in _index_const_overload_methods:
+    _type, method, _intrinsic = index_overload
+
+    ol_index_func = generate_index_overload(_type, _intrinsic)
+
+    overload_method(_type, method, target=DPEX_KERNEL_EXP_TARGET_NAME)(
+        ol_index_func
+    )
+
+
+@intrinsic(target=DPEX_KERNEL_EXP_TARGET_NAME)
+def _intrinsic_get_group(
+    ty_context, ty_nd_item: NdItemType  # pylint: disable=unused-argument
+):
+    """Generates group with a dimension of nd_item."""
+
+    if not isinstance(ty_nd_item, NdItemType):
+        raise TypingError(
+            f"Expected an NdItemType value, but encountered {ty_nd_item}"
+        )
+
+    ty_group = GroupType(ty_nd_item.ndim)
+    sig = ty_group(ty_nd_item)
+
+    # pylint: disable=unused-argument
+    def _intrinsic_get_group_gen(context, builder, sig, args):
+        group_struct = cgutils.create_struct_proxy(ty_group)(context, builder)
+        # pylint: disable=protected-access
+        return group_struct._getvalue()
+
+    return sig, _intrinsic_get_group_gen
+
+
+@overload_method(NdItemType, "get_group", target=DPEX_KERNEL_EXP_TARGET_NAME)
+def ol_nd_item_get_group(nd_item):
+    """SPIR-V overload for :meth:`numba_dpex.kernel_api.NdItem.get_group`.
+
+    Generates the same LLVM IR instruction as dpcpp for the
+    `sycl::nd_item::get_group` function.
+
+    Raises:
+        TypingError: When argument is not NdItem.
+    """
+    if not isinstance(nd_item, NdItemType):
+        # since it is a method overload, this error should not be reached
+        raise TypingError(
+            "Expected a nd_item should to be a NdItem value, but "
+            f"encountered {type(nd_item)}"
+        )
+
+    # pylint: disable=unused-argument
+    def ol_nd_item_get_group_impl(nd_item):
+        # pylint: disable=no-value-for-parameter
+        return _intrinsic_get_group(nd_item)
+
+    return ol_nd_item_get_group_impl
diff --git a/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_spv_atomic_inst_helper.py b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_spv_atomic_inst_helper.py
index 948589b581..575ea2d526 100644
--- a/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_spv_atomic_inst_helper.py
+++ b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_spv_atomic_inst_helper.py
@@ -8,7 +8,7 @@
 
 from numba.core import types
 
-from ..kernel_iface import MemoryOrder, MemoryScope
+from numba_dpex.kernel_api import MemoryOrder, MemoryScope
 
 
 class _SpvScope(IntEnum):
diff --git a/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/spv_atomic_fn_declarations.py b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/spv_atomic_fn_declarations.py
new file mode 100644
index 0000000000..ef41e287ab
--- /dev/null
+++ b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/spv_atomic_fn_declarations.py
@@ -0,0 +1,228 @@
+# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""
+Implements a set of helper functions to generate the LLVM IR for SPIR-V
+functions and their use inside an LLVM module.
+"""
+
+import warnings
+
+from llvmlite import ir as llvmir
+from numba.core import cgutils, types
+
+from numba_dpex.core import itanium_mangler as ext_itanium_mangler
+from numba_dpex.kernel_api_impl.spirv.target import CC_SPIR_FUNC
+
+_SUPPORT_CONVERGENT = True
+
+try:
+    llvmir.FunctionAttributes("convergent")
+except ValueError:
+    warnings.warn(
+        "convergent attribute is supported only starting llvmlite "
+        + "0.42. Not setting this attribute may result in unexpected behavior"
+        + "when using group_barrier"
+    )
+    _SUPPORT_CONVERGENT = False
+
+
+def get_or_insert_atomic_load_fn(context, module, atomic_ref_ty):
+    """
+    Gets or inserts a declaration for a __spirv_AtomicLoad call into the
+    specified LLVM IR module.
+    """
+    atomic_ref_dtype = atomic_ref_ty.dtype
+
+    if atomic_ref_dtype == types.float32:
+        atomic_ref_dtype = types.uint32
+    elif atomic_ref_dtype == types.float64:
+        atomic_ref_dtype = types.uint64
+
+    atomic_load_fn_retty = context.get_value_type(atomic_ref_dtype)
+    ptr_type = atomic_load_fn_retty.as_pointer()
+    ptr_type.addrspace = atomic_ref_ty.address_space
+
+    atomic_load_fn_arg_types = [
+        ptr_type,
+        llvmir.IntType(32),
+        llvmir.IntType(32),
+    ]
+    mangled_fn_name = ext_itanium_mangler.mangle_ext(
+        "__spirv_AtomicLoad",
+        [
+            types.CPointer(atomic_ref_dtype, addrspace=ptr_type.addrspace),
+            "__spv.Scope.Flag",
+            "__spv.MemorySemanticsMask.Flag",
+        ],
+    )
+
+    fn = cgutils.get_or_insert_function(
+        module,
+        llvmir.FunctionType(atomic_load_fn_retty, atomic_load_fn_arg_types),
+        mangled_fn_name,
+    )
+    fn.calling_convention = CC_SPIR_FUNC
+
+    if _SUPPORT_CONVERGENT:
+        fn.attributes.add("convergent")
+    fn.attributes.add("nounwind")
+
+    return fn
+
+
+def get_or_insert_spv_atomic_store_fn(context, module, atomic_ref_ty):
+    """
+    Gets or inserts a declaration for a __spirv_AtomicStore call into the
+    specified LLVM IR module.
+    """
+    atomic_ref_dtype = atomic_ref_ty.dtype
+
+    if atomic_ref_dtype == types.float32:
+        atomic_ref_dtype = types.uint32
+    elif atomic_ref_dtype == types.float64:
+        atomic_ref_dtype = types.uint64
+
+    ptr_type = context.get_value_type(atomic_ref_dtype).as_pointer()
+    ptr_type.addrspace = atomic_ref_ty.address_space
+    atomic_store_fn_retty = llvmir.VoidType()
+
+    atomic_store_fn_arg_types = [
+        ptr_type,
+        llvmir.IntType(32),
+        llvmir.IntType(32),
+        context.get_value_type(atomic_ref_dtype),
+    ]
+
+    mangled_fn_name = ext_itanium_mangler.mangle_ext(
+        "__spirv_AtomicStore",
+        [
+            types.CPointer(atomic_ref_dtype, addrspace=ptr_type.addrspace),
+            "__spv.Scope.Flag",
+            "__spv.MemorySemanticsMask.Flag",
+            atomic_ref_dtype,
+        ],
+    )
+
+    fn = cgutils.get_or_insert_function(
+        module,
+        llvmir.FunctionType(atomic_store_fn_retty, atomic_store_fn_arg_types),
+        mangled_fn_name,
+    )
+    fn.calling_convention = CC_SPIR_FUNC
+
+    if _SUPPORT_CONVERGENT:
+        fn.attributes.add("convergent")
+    fn.attributes.add("nounwind")
+
+    return fn
+
+
+def get_or_insert_spv_atomic_exchange_fn(context, module, atomic_ref_ty):
+    """
+    Gets or inserts a declaration for a __spirv_AtomicExchange call into the
+    specified LLVM IR module.
+    """
+    atomic_ref_dtype = atomic_ref_ty.dtype
+
+    if atomic_ref_dtype == types.float32:
+        atomic_ref_dtype = types.uint32
+    elif atomic_ref_dtype == types.float64:
+        atomic_ref_dtype = types.uint64
+
+    ptr_type = context.get_value_type(atomic_ref_dtype).as_pointer()
+    ptr_type.addrspace = atomic_ref_ty.address_space
+    atomic_exchange_fn_retty = context.get_value_type(atomic_ref_dtype)
+
+    atomic_exchange_fn_arg_types = [
+        ptr_type,
+        llvmir.IntType(32),
+        llvmir.IntType(32),
+        context.get_value_type(atomic_ref_dtype),
+    ]
+
+    mangled_fn_name = ext_itanium_mangler.mangle_ext(
+        "__spirv_AtomicExchange",
+        [
+            types.CPointer(atomic_ref_dtype, addrspace=ptr_type.addrspace),
+            "__spv.Scope.Flag",
+            "__spv.MemorySemanticsMask.Flag",
+            atomic_ref_dtype,
+        ],
+    )
+
+    fn = cgutils.get_or_insert_function(
+        module,
+        llvmir.FunctionType(
+            atomic_exchange_fn_retty, atomic_exchange_fn_arg_types
+        ),
+        mangled_fn_name,
+    )
+    fn.calling_convention = CC_SPIR_FUNC
+
+    if _SUPPORT_CONVERGENT:
+        fn.attributes.add("convergent")
+    fn.attributes.add("nounwind")
+
+    return fn
+
+
+def get_or_insert_spv_atomic_compare_exchange_fn(
+    context, module, atomic_ref_ty
+):
+    """
+    Gets or inserts a declaration for a __spirv_AtomicCompareExchange call into the
+    specified LLVM IR module.
+    """
+    atomic_ref_dtype = atomic_ref_ty.dtype
+
+    # Spirv spec requires arguments and return type to be of integer types.
+    # That is why the type is changed from float to int
+    # while maintaining the bit-width.
+    # During function call, bitcasting is performed
+    # to adhere to this convention.
+    if atomic_ref_dtype == types.float32:
+        atomic_ref_dtype = types.uint32
+    elif atomic_ref_dtype == types.float64:
+        atomic_ref_dtype = types.uint64
+
+    ptr_type = context.get_value_type(atomic_ref_dtype).as_pointer()
+    ptr_type.addrspace = atomic_ref_ty.address_space
+    atomic_cmpexchg_fn_retty = context.get_value_type(atomic_ref_dtype)
+
+    atomic_cmpexchg_fn_arg_types = [
+        ptr_type,
+        llvmir.IntType(32),
+        llvmir.IntType(32),
+        llvmir.IntType(32),
+        context.get_value_type(atomic_ref_dtype),
+        context.get_value_type(atomic_ref_dtype),
+    ]
+
+    mangled_fn_name = ext_itanium_mangler.mangle_ext(
+        "__spirv_AtomicCompareExchange",
+        [
+            types.CPointer(atomic_ref_dtype, addrspace=ptr_type.addrspace),
+            "__spv.Scope.Flag",
+            "__spv.MemorySemanticsMask.Flag",
+            "__spv.MemorySemanticsMask.Flag",
+            atomic_ref_dtype,
+            atomic_ref_dtype,
+        ],
+    )
+
+    fn = cgutils.get_or_insert_function(
+        module,
+        llvmir.FunctionType(
+            atomic_cmpexchg_fn_retty, atomic_cmpexchg_fn_arg_types
+        ),
+        mangled_fn_name,
+    )
+    fn.calling_convention = CC_SPIR_FUNC
+
+    if _SUPPORT_CONVERGENT:
+        fn.attributes.add("convergent")
+    fn.attributes.add("nounwind")
+
+    return fn
diff --git a/numba_dpex/experimental/decorators.py b/numba_dpex/experimental/decorators.py
index df44320572..a7529cea24 100644
--- a/numba_dpex/experimental/decorators.py
+++ b/numba_dpex/experimental/decorators.py
@@ -15,8 +15,8 @@
     target_registry,
 )
 
-from numba_dpex.core.targets.kernel_target import CompilationMode
-from numba_dpex.experimental.kernel_dispatcher import KernelDispatcher
+from numba_dpex.kernel_api_impl.spirv.dispatcher import SPIRVKernelDispatcher
+from numba_dpex.kernel_api_impl.spirv.target import CompilationMode
 
 from .target import DPEX_KERNEL_EXP_TARGET_NAME
 
@@ -74,11 +74,11 @@ def kernel(func_or_sig=None, **options):
     for sig in sigs:
         if isinstance(sig, str):
             raise NotImplementedError(
-                "Specifying signatures as string is not yet supported by numba-dpex"
+                "Specifying signatures as string is not yet supported"
             )
 
     def _kernel_dispatcher(pyfunc):
-        disp: KernelDispatcher = dispatcher(
+        disp: SPIRVKernelDispatcher = dispatcher(
             pyfunc=pyfunc,
             targetoptions=options,
         )
diff --git a/numba_dpex/experimental/kernel_iface/__init__.py b/numba_dpex/experimental/kernel_iface/__init__.py
deleted file mode 100644
index 4433756733..0000000000
--- a/numba_dpex/experimental/kernel_iface/__init__.py
+++ /dev/null
@@ -1,14 +0,0 @@
-# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-"""The kernel_iface provides a set of Python classes and functions that are
-analogous to the C++ SYCL API. The kernel_iface API is meant to allow
-prototyping SYCL-like kernels in pure Python before compiling them using
-numba_dpex.kernel.
-"""
-
-from .atomic_ref import AtomicRef
-from .memory_enums import AddressSpace, MemoryOrder, MemoryScope
-
-__all__ = ["AddressSpace", "AtomicRef", "MemoryOrder", "MemoryScope"]
diff --git a/numba_dpex/experimental/launcher.py b/numba_dpex/experimental/launcher.py
index 503a8bafe7..64f5c36f4d 100644
--- a/numba_dpex/experimental/launcher.py
+++ b/numba_dpex/experimental/launcher.py
@@ -6,6 +6,8 @@
 from either CPython or a numba_dpex.dpjit decorated function.
 """
 
+import warnings
+from inspect import signature
 from typing import NamedTuple, Union
 
 import dpctl
@@ -18,15 +20,19 @@
 
 from numba_dpex import dpjit
 from numba_dpex.core.targets.dpjit_target import DPEX_TARGET_NAME
-from numba_dpex.core.targets.kernel_target import DpexKernelTargetContext
 from numba_dpex.core.types import DpctlSyclEvent, NdRangeType, RangeType
+from numba_dpex.core.types.kernel_api.index_space_ids import (
+    ItemType,
+    NdItemType,
+)
 from numba_dpex.core.utils import kernel_launcher as kl
 from numba_dpex.dpctl_iface import libsyclinterface_bindings as sycl
 from numba_dpex.dpctl_iface.wrappers import wrap_event_reference
-from numba_dpex.experimental.kernel_dispatcher import (
-    KernelDispatcher,
-    _KernelCompileResult,
+from numba_dpex.kernel_api_impl.spirv.dispatcher import (
+    SPIRVKernelDispatcher,
+    _SPIRVKernelCompileResult,
 )
+from numba_dpex.kernel_api_impl.spirv.target import SPIRVTargetContext
 
 
 class LLRange(NamedTuple):
@@ -129,16 +135,33 @@ def _submit_kernel(  # pylint: disable=too-many-arguments
     else:
         sig = ty_return(ty_kernel_fn, ty_index_space, ty_kernel_args_tuple)
 
-    kernel_sig = types.void(*ty_kernel_args_tuple)
+    # Add Item/NdItem as a first argument to kernel arguments list. It is
+    # an empty struct so any other modifications at kernel submission are not
+    # needed.
+    if len(signature(ty_kernel_fn.dispatcher.py_func).parameters) > len(
+        ty_kernel_args_tuple
+    ):
+        if isinstance(ty_index_space, RangeType):
+            ty_item = ItemType(ty_index_space.ndim)
+        else:
+            ty_item = NdItemType(ty_index_space.ndim)
+
+        ty_kernel_args_tuple = (ty_item, *ty_kernel_args_tuple)
+    else:
+        warnings.warn(
+            "Kernels without item/nd_item will be not supported in the future",
+            DeprecationWarning,
+        )
+
     # ty_kernel_fn is type specific to exact function, so we can get function
     # directly from type and compile it. Thats why we don't need to get it in
     # codegen
-    kernel_dispatcher: KernelDispatcher = ty_kernel_fn.dispatcher
-    kcres: _KernelCompileResult = kernel_dispatcher.get_compile_result(
-        kernel_sig
+    kernel_dispatcher: SPIRVKernelDispatcher = ty_kernel_fn.dispatcher
+    kcres: _SPIRVKernelCompileResult = kernel_dispatcher.get_compile_result(
+        types.void(*ty_kernel_args_tuple)  # kernel signature
     )
     kernel_module: kl.SPIRVKernelModule = kcres.kernel_device_ir_module
-    kernel_targetctx: DpexKernelTargetContext = kernel_dispatcher.targetctx
+    kernel_targetctx: SPIRVTargetContext = kernel_dispatcher.targetctx
 
     def codegen(
         cgctx: CPUContext, builder: llvmir.IRBuilder, sig, llargs: list
diff --git a/numba_dpex/experimental/models.py b/numba_dpex/experimental/models.py
index 34c789d507..f9e7a2f53d 100644
--- a/numba_dpex/experimental/models.py
+++ b/numba_dpex/experimental/models.py
@@ -6,16 +6,19 @@
 numba_dpex.experimental module.
 """
 
-from llvmlite import ir as llvmir
 from numba.core import types
 from numba.core.datamodel import DataModelManager, models
-from numba.core.datamodel.models import PrimitiveModel, StructModel
+from numba.core.datamodel.models import StructModel
 from numba.core.extending import register_model
 
 import numba_dpex.core.datamodel.models as dpex_core_models
+from numba_dpex.core.types.kernel_api.index_space_ids import (
+    GroupType,
+    ItemType,
+    NdItemType,
+)
 
-from .dpcpp_types import AtomicRefType
-from .literal_intenum_type import IntEnumLiteral
+from ..core.types.kernel_api.atomic_ref import AtomicRefType
 from .types import KernelDispatcherType
 
 
@@ -32,15 +35,13 @@ def __init__(self, dmm, fe_type):
         super().__init__(dmm, fe_type, members)
 
 
-class IntEnumLiteralModel(PrimitiveModel):
-    """Representation of an object of LiteralIntEnum type using Numba's
-    PrimitiveModel that can be represented natively in the target in all
-    usage contexts.
-    """
+class EmptyStructModel(StructModel):
+    """Data model that does not take space. Intended to be used with types that
+    are presented only at typing stage and not represented physically."""
 
     def __init__(self, dmm, fe_type):
-        be_type = llvmir.IntType(fe_type.bitwidth)
-        super().__init__(dmm, fe_type, be_type)
+        members = []
+        super().__init__(dmm, fe_type, members)
 
 
 def _init_exp_data_model_manager() -> DataModelManager:
@@ -57,9 +58,17 @@ def _init_exp_data_model_manager() -> DataModelManager:
     dmm = dpex_core_models.dpex_data_model_manager.copy()
 
     # Register the types and data model in the DpexExpTargetContext
-    dmm.register(IntEnumLiteral, IntEnumLiteralModel)
     dmm.register(AtomicRefType, AtomicRefModel)
 
+    # Register the GroupType type
+    dmm.register(GroupType, EmptyStructModel)
+
+    # Register the ItemType type
+    dmm.register(ItemType, EmptyStructModel)
+
+    # Register the NdItemType type
+    dmm.register(NdItemType, EmptyStructModel)
+
     return dmm
 
 
@@ -67,3 +76,12 @@ def _init_exp_data_model_manager() -> DataModelManager:
 
 # Register any new type that should go into numba.core.datamodel.default_manager
 register_model(KernelDispatcherType)(models.OpaqueModel)
+
+# Register the GroupType type
+register_model(GroupType)(EmptyStructModel)
+
+# Register the ItemType type
+register_model(ItemType)(EmptyStructModel)
+
+# Register the NdItemType type
+register_model(NdItemType)(EmptyStructModel)
diff --git a/numba_dpex/experimental/target.py b/numba_dpex/experimental/target.py
index 7eb3594264..924c4d637c 100644
--- a/numba_dpex/experimental/target.py
+++ b/numba_dpex/experimental/target.py
@@ -8,21 +8,15 @@
 
 from functools import cached_property
 
-from llvmlite import ir as llvmir
-from numba.core import types
 from numba.core.descriptors import TargetDescriptor
 from numba.core.target_extension import GPU, target_registry
-from numba.core.types.scalars import IntEnumClass
 
 from numba_dpex.core.descriptor import DpexTargetOptions
-from numba_dpex.core.targets.kernel_target import (
-    DpexKernelTargetContext,
-    DpexKernelTypingContext,
-)
 from numba_dpex.experimental.models import exp_dmm
-
-from .flag_enum import FlagEnum
-from .literal_intenum_type import IntEnumLiteral
+from numba_dpex.kernel_api_impl.spirv.target import (
+    SPIRVTargetContext,
+    SPIRVTypingContext,
+)
 
 
 #  pylint: disable=R0903
@@ -35,7 +29,7 @@ class SyclDeviceExp(GPU):
 target_registry[DPEX_KERNEL_EXP_TARGET_NAME] = SyclDeviceExp
 
 
-class DpexExpKernelTypingContext(DpexKernelTypingContext):
+class DpexExpKernelTypingContext(SPIRVTypingContext):
     """Experimental typing context class extending the DpexKernelTypingContext
     by overriding super class functions for new experimental types.
 
@@ -45,42 +39,11 @@ class DpexExpKernelTypingContext(DpexKernelTypingContext):
     are stable enough to be migrated to DpexKernelTypingContext.
     """
 
-    def resolve_value_type(self, val):
-        """
-        Return the numba type of a Python value that is being used
-        as a runtime constant.
-        ValueError is raised for unsupported types.
-        """
-
-        typ = super().resolve_value_type(val)
-
-        if isinstance(typ, IntEnumClass) and issubclass(val, FlagEnum):
-            typ = IntEnumLiteral(val)
-
-        return typ
-
-    def resolve_getattr(self, typ, attr):
-        """
-        Resolve getting the attribute *attr* (a string) on the Numba type.
-        The attribute's type is returned, or None if resolution failed.
-        """
-        retty = None
-
-        if isinstance(typ, IntEnumLiteral):
-            try:
-                attrval = getattr(typ.literal_value, attr).value
-                retty = types.IntegerLiteral(attrval)
-            except ValueError:
-                pass
-        else:
-            retty = super().resolve_getattr(typ, attr)
-        return retty
-
 
 #  pylint: disable=W0223
 # FIXME: Remove the pylint disablement once we add an override for
 # get_executable
-class DpexExpKernelTargetContext(DpexKernelTargetContext):
+class DpexExpKernelTargetContext(SPIRVTargetContext):
     """Experimental target context class extending the DpexKernelTargetContext
     by overriding super class functions for new experimental types.
 
@@ -95,22 +58,6 @@ def __init__(self, typingctx, target=DPEX_KERNEL_EXP_TARGET_NAME):
         super().__init__(typingctx, target)
         self.data_model_manager = exp_dmm
 
-    def get_getattr(self, typ, attr):
-        """
-        Overrides the get_getattr function to provide an implementation for
-        getattr call on an IntegerEnumLiteral type.
-        """
-
-        if isinstance(typ, IntEnumLiteral):
-            #  pylint: disable=W0613
-            def enum_literal_getattr_imp(context, builder, typ, val, attr):
-                enum_attr_value = getattr(typ.literal_value, attr).value
-                return llvmir.Constant(llvmir.IntType(64), enum_attr_value)
-
-            return enum_literal_getattr_imp
-
-        return super().get_getattr(typ, attr)
-
 
 class DpexExpKernelTarget(TargetDescriptor):
     """
diff --git a/numba_dpex/experimental/typeof.py b/numba_dpex/experimental/typeof.py
index e90bfccf0a..108e9d3b09 100644
--- a/numba_dpex/experimental/typeof.py
+++ b/numba_dpex/experimental/typeof.py
@@ -9,8 +9,14 @@
 
 from numba.extending import typeof_impl
 
-from .dpcpp_types import AtomicRefType
-from .kernel_iface import AtomicRef
+from numba_dpex.core.types.kernel_api.index_space_ids import (
+    GroupType,
+    ItemType,
+    NdItemType,
+)
+from numba_dpex.kernel_api import AtomicRef, Group, Item, NdItem
+
+from ..core.types.kernel_api.atomic_ref import AtomicRefType
 
 
 @typeof_impl.register(AtomicRef)
@@ -33,3 +39,48 @@ def typeof_atomic_ref(val: AtomicRef, ctx) -> AtomicRefType:
         memory_scope=val.memory_scope.value,
         address_space=val.address_space.value,
     )
+
+
+@typeof_impl.register(Group)
+def typeof_group(val: Group, c):
+    """Registers the type inference implementation function for a
+    numba_dpex.kernel_api.Group PyObject.
+
+    Args:
+        val : An instance of numba_dpex.kernel_api.Group.
+        c : Unused argument used to be consistent with Numba API.
+
+    Returns: A numba_dpex.experimental.core.types.kernel_api.items.GroupType
+        instance.
+    """
+    return GroupType(val.ndim)
+
+
+@typeof_impl.register(Item)
+def typeof_item(val: Item, c):
+    """Registers the type inference implementation function for a
+    numba_dpex.kernel_api.Item PyObject.
+
+    Args:
+        val : An instance of numba_dpex.kernel_api.Item.
+        c : Unused argument used to be consistent with Numba API.
+
+    Returns: A numba_dpex.experimental.core.types.kernel_api.items.ItemType
+        instance.
+    """
+    return ItemType(val.ndim)
+
+
+@typeof_impl.register(NdItem)
+def typeof_nditem(val, c):
+    """Registers the type inference implementation function for a
+    numba_dpex.kernel_api.NdItem PyObject.
+
+    Args:
+        val : An instance of numba_dpex.kernel_api.NdItem.
+        c : Unused argument used to be consistent with Numba API.
+
+    Returns: A numba_dpex.experimental.core.types.kernel_api.items.NdItemType
+        instance.
+    """
+    return NdItemType(val.ndim)
diff --git a/numba_dpex/kernel_api/__init__.py b/numba_dpex/kernel_api/__init__.py
new file mode 100644
index 0000000000..32c9f933ef
--- /dev/null
+++ b/numba_dpex/kernel_api/__init__.py
@@ -0,0 +1,31 @@
+# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""
+The kernel_api module provides a set of Python classes and functions that are
+analogous to the C++ SYCL API. The kernel_api module is meant to allow
+prototyping SYCL-like kernels in pure Python before compiling them using
+numba_dpex.
+"""
+
+from .atomic_ref import AtomicRef
+from .barrier import group_barrier
+from .index_space_ids import Group, Item, NdItem
+from .launcher import call_kernel
+from .memory_enums import AddressSpace, MemoryOrder, MemoryScope
+from .ranges import NdRange, Range
+
+__all__ = [
+    "AddressSpace",
+    "AtomicRef",
+    "MemoryOrder",
+    "MemoryScope",
+    "NdRange",
+    "Range",
+    "Group",
+    "NdItem",
+    "Item",
+    "group_barrier",
+    "call_kernel",
+]
diff --git a/numba_dpex/experimental/kernel_iface/atomic_ref.py b/numba_dpex/kernel_api/atomic_ref.py
similarity index 64%
rename from numba_dpex/experimental/kernel_iface/atomic_ref.py
rename to numba_dpex/kernel_api/atomic_ref.py
index 2090e39757..7d9e426aa1 100644
--- a/numba_dpex/experimental/kernel_iface/atomic_ref.py
+++ b/numba_dpex/kernel_api/atomic_ref.py
@@ -68,7 +68,7 @@ def fetch_add(self, val):
         Returns: The original value of the object referenced by the AtomicRef.
 
         """
-        old = self._ref[self._index]
+        old = self._ref[self._index].copy()
         self._ref[self._index] += val
         return old
 
@@ -78,13 +78,13 @@ def fetch_sub(self, val):
         Returns the original value of the object.
 
         Args:
-            val : Value to be subtracted to the object referenced by the
+            val : Value to be subtracted from the object referenced by the
             AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
 
         """
-        old = self._ref[self._index]
+        old = self._ref[self._index].copy()
         self._ref[self._index] -= val
         return old
 
@@ -94,13 +94,13 @@ def fetch_min(self, val):
         referenced object. Returns the original value of the object.
 
         Args:
-            val : Value to be compared against to the object referenced by the
+            val : Value to be compared against the object referenced by the
             AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
 
         """
-        old = self._ref[self._index]
+        old = self._ref[self._index].copy()
         self._ref[self._index] = min(old, val)
         return old
 
@@ -110,13 +110,13 @@ def fetch_max(self, val):
         referenced object. Returns the original value of the object.
 
         Args:
-            val : Value to be compared against to the object referenced by the
+            val : Value to be compared against the object referenced by the
             AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
 
         """
-        old = self._ref[self._index]
+        old = self._ref[self._index].copy()
         self._ref[self._index] = max(old, val)
         return old
 
@@ -126,13 +126,13 @@ def fetch_and(self, val):
         referenced object. Returns the original value of the object.
 
         Args:
-            val : Value to be bitwise ANDed against to the object referenced by
+            val : Value to be bitwise ANDed against the object referenced by
             the AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
 
         """
-        old = self._ref[self._index]
+        old = self._ref[self._index].copy()
         self._ref[self._index] &= val
         return old
 
@@ -142,13 +142,13 @@ def fetch_or(self, val):
         referenced object. Returns the original value of the object.
 
         Args:
-            val : Value to be bitwise ORed against to the object referenced by
+            val : Value to be bitwise ORed against the object referenced by
             the AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
 
         """
-        old = self._ref[self._index]
+        old = self._ref[self._index].copy()
         self._ref[self._index] |= val
         return old
 
@@ -158,12 +158,71 @@ def fetch_xor(self, val):
         referenced object. Returns the original value of the object.
 
         Args:
-            val : Value to be bitwise XORed against to the object referenced by
+            val : Value to be bitwise XORed against the object referenced by
             the AtomicRef.
 
         Returns: The original value of the object referenced by the AtomicRef.
 
         """
-        old = self._ref[self._index]
+        old = self._ref[self._index].copy()
         self._ref[self._index] ^= val
         return old
+
+    def load(self):
+        """Loads the value of the object referenced by the AtomicRef.
+
+        Returns: The value of the object referenced by the AtomicRef.
+
+        """
+        return self._ref[self._index]
+
+    def store(self, val):
+        """Stores operand ``val`` to the object referenced by the AtomicRef.
+
+        Args:
+            val : Value to be stored in the object referenced by
+            the AtomicRef.
+
+        """
+        self._ref[self._index] = val
+
+    def exchange(self, val):
+        """Replaces the value of the object referenced by the AtomicRef
+        with value of ``val``. Returns the original value of the referenced object.
+
+        Args:
+            val : Value to be exchanged against the object referenced by
+            the AtomicRef.
+
+        Returns: The original value of the object referenced by the AtomicRef.
+
+        """
+        old = self._ref[self._index].copy()
+        self._ref[self._index] = val
+        return old
+
+    def compare_exchange(self, expected, desired, expected_idx=0):
+        """Compares the value of the object referenced by the AtomicRef
+        against the value of ``expected[expected_idx]``.
+        If the values are equal, replaces the value of the
+        referenced object with the value of ``desired``.
+        Otherwise assigns the original value of the
+        referenced object to ``expected[expected_idx]``.
+
+        Args:
+            expected : Array containing the expected value of the
+            object referenced by the AtomicRef.
+            desired : Value that replaces the value of the object
+            referenced by the AtomicRef.
+            expected_idx: Offset in `expected` array where the expected
+            value of the object referenced by the AtomicRef is present.
+
+        Returns: Returns ``True`` if the comparison operation and
+        replacement operation were successful.
+
+        """
+        if self._ref[self._index] == expected[expected_idx]:
+            self._ref[self._index] = desired
+            return True
+        expected[expected_idx] = self._ref[self._index]
+        return False
diff --git a/numba_dpex/kernel_api/barrier.py b/numba_dpex/kernel_api/barrier.py
new file mode 100644
index 0000000000..744262403d
--- /dev/null
+++ b/numba_dpex/kernel_api/barrier.py
@@ -0,0 +1,33 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Python functions that simulate SYCL's barrier primitives.
+"""
+
+from .index_space_ids import Group
+from .memory_enums import MemoryScope
+
+
+def group_barrier(group: Group, fence_scope=MemoryScope.WORK_GROUP):
+    """Performs a barrier operation across all work-items in a work group.
+
+    The function is modeled after the ``sycl::group_barrier`` function. It
+    synchronizes work within a group of work items. All the work-items
+    of the group must execute the barrier construct before any work-item
+    continues execution beyond the barrier.
+
+    The ``group_barrier`` performs mem-fence operations ensuring that memory
+    accesses issued before the barrier are not re-ordered with those issued
+    after the barrier: all work-items in group g execute a release fence prior
+    to synchronizing at the barrier, all work-items in group g execute an
+    acquire fence afterwards, and there is an implicit synchronization of these
+    fences as if provided by an explicit atomic operation on an atomic object.
+
+    Args:
+        fence_scope (optional): scope of any memory consistency
+        operations that are performed by the barrier.
+    """
+
+    # TODO: A pure Python simulation of a group_barrier will be added later.
+    raise NotImplementedError
diff --git a/numba_dpex/experimental/flag_enum.py b/numba_dpex/kernel_api/flag_enum.py
similarity index 70%
rename from numba_dpex/experimental/flag_enum.py
rename to numba_dpex/kernel_api/flag_enum.py
index e800c0addf..e9e6a98dfe 100644
--- a/numba_dpex/experimental/flag_enum.py
+++ b/numba_dpex/kernel_api/flag_enum.py
@@ -2,7 +2,8 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
-"""Provides a FlagEnum class to help distinguish IntEnum types that numba-dpex
+"""
+Provides a FlagEnum class to help distinguish IntEnum types that numba_dpex
 intends to use as Integer literal types inside the compiler type inferring
 infrastructure.
 """
@@ -10,13 +11,13 @@
 
 
 class FlagEnum(IntEnum):
-    """Helper class to distinguish IntEnum types that numba-dpex should consider
+    """Helper class to distinguish IntEnum types that numba_dpex should consider
     as Numba Literal types.
     """
 
     @classmethod
     def basetype(cls) -> int:
-        """Returns an dummy int object that helps numba-dpex infer the type of
+        """Returns an dummy int object that helps numba_dpex infer the type of
         an instance of a FlagEnum class.
 
         Returns:
diff --git a/numba_dpex/kernel_api/index_space_ids.py b/numba_dpex/kernel_api/index_space_ids.py
new file mode 100644
index 0000000000..4e10bfc688
--- /dev/null
+++ b/numba_dpex/kernel_api/index_space_ids.py
@@ -0,0 +1,237 @@
+# SPDX-FileCopyrightText: 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Implements a mock Python classes to represent ``sycl::item`` and
+``sycl::nd_item``for prototyping numba_dpex kernel functions before they are JIT
+compiled.
+"""
+
+from .ranges import Range
+
+
+class Group:
+    """Analogue to the ``sycl::group`` type."""
+
+    def __init__(
+        self,
+        global_range: Range,
+        local_range: Range,
+        group_range: Range,
+        index: list,
+    ):
+        self._global_range = global_range
+        self._local_range = local_range
+        self._group_range = group_range
+        self._index = index
+        self._leader = False
+
+    def get_group_id(self, dim):
+        """Returns the index of the work-group within the global nd-range for
+        specified dimension.
+
+        Since the work-items in a work-group have a defined position within the
+        global nd-range, the returned group id can be used along with the local
+        id to uniquely identify the work-item in the global nd-range.
+        """
+        if dim > len(self._index) - 1:
+            raise ValueError(
+                "Dimension value is out of bounds for the group index"
+            )
+        return self._index[dim]
+
+    def get_group_linear_id(self):
+        """Returns a linearized version of the work-group index."""
+        if len(self._index) == 1:
+            return self._index[0]
+        if len(self._index) == 2:
+            return self._index[0] * self._group_range[1] + self._index[1]
+        return (
+            (self._index[0] * self._group_range[1] * self._group_range[2])
+            + (self._index[1] * self._group_range[2])
+            + (self._index[2])
+        )
+
+    def get_group_range(self, dim):
+        """Returns a the extent of the range representing the number of groups
+        in the nd-range for a specified dimension.
+        """
+        return self._group_range[dim]
+
+    def get_group_linear_range(self):
+        """Return the total number of work-groups in the nd_range."""
+        num_wg = 1
+        for ext in self._group_range:
+            num_wg *= ext
+
+        return num_wg
+
+    def get_local_range(self, dim):
+        """Returns the extent of the SYCL range representing all dimensions
+        of the local range for a specified dimension. This local range may
+        have been provided by the programmer, or chosen by the SYCL runtime.
+        """
+        return self._local_range[dim]
+
+    def get_local_linear_range(self):
+        """Return the total number of work-items in the work-group."""
+        num_wi = 1
+        for ext in self._local_range:
+            num_wi *= ext
+
+        return num_wi
+
+    @property
+    def leader(self):
+        """Return true for exactly one work-item in the work-group, if the
+        calling work-item is the leader of the work-group, and false for all
+        other work-items in the work-group.
+
+        The leader of the work-group is determined during construction of the
+        work-group, and is invariant for the lifetime of the work-group. The
+        leader of the work-group is guaranteed to be the work-item with a
+        local id of 0.
+
+
+        Returns:
+            bool: If the work item is the designated leader of the
+        """
+        return self._leader
+
+    @leader.setter
+    def leader(self, work_item_id):
+        """Sets the leader attribute for the group."""
+        self._leader = work_item_id
+
+
+class Item:
+    """Analogue to the ``sycl::item`` type. Identifies an instance of the
+    function object executing at each point in an Range.
+    """
+
+    def __init__(self, extent: Range, index: list):
+        self._extent = extent
+        self._index = index
+
+    def get_linear_id(self):
+        """Get the linear id associated with this item for all dimensions.
+        Original implementation could be found at ``sycl::item_base`` class.
+
+        Returns:
+            int: The linear id.
+        """
+        if len(self._extent) == 1:
+            return self._index[0]
+        if len(self._extent) == 2:
+            return self._index[0] * self._extent[1] + self._index[1]
+        return (
+            (self._index[0] * self._extent[1] * self._extent[2])
+            + (self._index[1] * self._extent[2])
+            + (self._index[2])
+        )
+
+    def get_id(self, idx):
+        """Get the id for a specific dimension.
+
+        Returns:
+            int: The id
+        """
+        return self._index[idx]
+
+    def get_range(self, idx):
+        """Get the range size for a specific dimension.
+
+        Returns:
+            int: The size
+        """
+        return self._extent[idx]
+
+    @property
+    def ndim(self) -> int:
+        """Returns the rank of a Item object.
+
+        Returns:
+            int: Number of dimensions in the Item object
+        """
+        return self._extent.ndim
+
+
+class NdItem:
+    """Analogue to the ``sycl::nd_item`` type. Identifies an instance of the
+    function object executing at each point in an NdRange.
+    """
+
+    # TODO: define group type
+    def __init__(self, global_item: Item, local_item: Item, group: Group):
+        # TODO: assert offset and dimensions
+        self._global_item = global_item
+        self._local_item = local_item
+        self._group = group
+        if self.get_local_linear_id() == 0:
+            self._group.leader = True
+
+    def get_global_id(self, idx):
+        """Get the global id for a specific dimension.
+
+        Returns:
+            int: The global id
+        """
+        return self._global_item.get_id(idx)
+
+    def get_global_linear_id(self):
+        """Get the global linear id associated with this item for all
+        dimensions.
+
+        Returns:
+            int: The global linear id.
+        """
+        return self._global_item.get_linear_id()
+
+    def get_local_id(self, idx):
+        """Get the local id for a specific dimension.
+
+        Returns:
+            int: The local id
+        """
+        return self._local_item.get_id(idx)
+
+    def get_local_linear_id(self):
+        """Get the local linear id associated with this item for all
+        dimensions.
+
+        Returns:
+            int: The local linear id.
+        """
+        return self._local_item.get_linear_id()
+
+    def get_global_range(self, idx):
+        """Get the global range size for a specific dimension.
+
+        Returns:
+            int: The size
+        """
+        return self._global_item.get_range(idx)
+
+    def get_local_range(self, idx):
+        """Get the local range size for a specific dimension.
+
+        Returns:
+            int: The size
+        """
+        return self._local_item.get_range(idx)
+
+    def get_group(self):
+        """Returns the group.
+
+        Returns:
+            A group object."""
+        return self._group
+
+    @property
+    def ndim(self) -> int:
+        """Returns the rank of a NdItem object.
+
+        Returns:
+            int: Number of dimensions in the NdItem object
+        """
+        return self._global_item.ndim
diff --git a/numba_dpex/kernel_api/launcher.py b/numba_dpex/kernel_api/launcher.py
new file mode 100644
index 0000000000..98a293e52b
--- /dev/null
+++ b/numba_dpex/kernel_api/launcher.py
@@ -0,0 +1,125 @@
+# SPDX-FileCopyrightText: 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Implementation of mock kernel launcher functions
+"""
+
+from inspect import signature
+from itertools import product
+
+from .index_space_ids import Group, Item, NdItem
+from .ranges import NdRange, Range
+
+
+def _range_kernel_launcher(kernel_fn, index_range, *kernel_args):
+    """Executes a function that mocks a range kernel.
+
+    Converts the range into a set of index tuple that represent an element in
+    the iteration domain over which the kernel will be executed. Then the
+    kernel is called sequentially over that set of indices after each index
+    value is used to construct an Item object.
+
+    Args:
+        kernel_fn : A callable function object
+        index_range (numba_dpex.Range): An instance of a Range object
+
+    Raises:
+        ValueError: If the number of passed in kernel arguments is not the
+        number of function parameters subtracted by one. The first kernel
+        argument is expected to be an Item object.
+    """
+
+    range_sets = [range(ir) for ir in index_range]
+    index_tuples = list(product(*range_sets))
+
+    for idx in index_tuples:
+        it = Item(extent=index_range, index=idx)
+
+        if len(signature(kernel_fn).parameters) - len(kernel_args) != 1:
+            raise ValueError(
+                "Required number of kernel function arguments do not "
+                "match provided number of kernel args"
+            )
+
+        kernel_fn(it, *kernel_args)
+
+
+def _ndrange_kernel_launcher(kernel_fn, index_range, *kernel_args):
+    """Executes a function that mocks a range kernel.
+
+    Args:
+        kernel_fn : A callable function object
+        index_range (numba_dpex.NdRange): An instance of a NdRange object
+
+    Raises:
+        ValueError: If the number of passed in kernel arguments is not the
+        number of function parameters subtracted by one. The first kernel
+        argument is expected to be an Item object.
+    """
+    group_range = tuple(
+        gr // lr
+        for gr, lr in zip(index_range.global_range, index_range.local_range)
+    )
+    local_range_sets = [range(ir) for ir in index_range.local_range]
+    group_range_sets = [range(gr) for gr in group_range]
+    local_index_tuples = list(product(*local_range_sets))
+    group_index_tuples = list(product(*group_range_sets))
+
+    # Loop over the groups (parallel loop)
+    for gidx in group_index_tuples:
+        # loop over work items in the group (parallel loop)
+        for lidx in local_index_tuples:
+            global_id = []
+            # to calculate global indices
+            for dim, gidx_val in enumerate(gidx):
+                global_id.append(
+                    gidx_val * index_range.local_range[dim] + lidx[dim]
+                )
+            # Every NdItem has its own global Item, local Item and Group
+            nditem = NdItem(
+                global_item=Item(
+                    extent=index_range.global_range, index=global_id
+                ),
+                local_item=Item(extent=index_range.local_range, index=lidx),
+                group=Group(
+                    index_range.global_range,
+                    index_range.local_range,
+                    group_range,
+                    gidx,
+                ),
+            )
+
+            if len(signature(kernel_fn).parameters) - len(kernel_args) != 1:
+                raise ValueError(
+                    "Required number of kernel function arguments do not "
+                    "match provided number of kernel args"
+                )
+
+            kernel_fn(nditem, *kernel_args)
+
+
+def call_kernel(kernel_fn, index_range, *kernel_args):
+    """Mocks the launching of a kernel function over either a Range or NdRange.
+
+    Args:
+        kernel_fn : A callable function object
+        index_range (numba_dpex.Range): An instance of a Range object
+
+    Raises:
+        ValueError: If the first positional argument is not callable
+        ValueError: If the second positional argument is not a Range or an
+        Ndrange object
+    """
+    if not callable(kernel_fn):
+        raise ValueError(
+            "Expected the first positional argument to be a function object"
+        )
+    if isinstance(index_range, Range):
+        _range_kernel_launcher(kernel_fn, index_range, *kernel_args)
+    elif isinstance(index_range, NdRange):
+        _ndrange_kernel_launcher(kernel_fn, index_range, *kernel_args)
+    else:
+        raise ValueError(
+            "Expected second positional argument to be Range or NdRange object"
+        )
diff --git a/numba_dpex/experimental/kernel_iface/memory_enums.py b/numba_dpex/kernel_api/memory_enums.py
similarity index 97%
rename from numba_dpex/experimental/kernel_iface/memory_enums.py
rename to numba_dpex/kernel_api/memory_enums.py
index c65342f008..ef738daf2b 100644
--- a/numba_dpex/experimental/kernel_iface/memory_enums.py
+++ b/numba_dpex/kernel_api/memory_enums.py
@@ -6,7 +6,7 @@
 memory enum classes.
 """
 
-from numba_dpex.experimental.flag_enum import FlagEnum
+from numba_dpex.kernel_api.flag_enum import FlagEnum
 
 
 class MemoryOrder(FlagEnum):
diff --git a/numba_dpex/core/kernel_interface/indexers.py b/numba_dpex/kernel_api/ranges.py
similarity index 56%
rename from numba_dpex/core/kernel_interface/indexers.py
rename to numba_dpex/kernel_api/ranges.py
index a5914c14f4..85cea812af 100644
--- a/numba_dpex/core/kernel_interface/indexers.py
+++ b/numba_dpex/kernel_api/ranges.py
@@ -2,15 +2,16 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
-from collections.abc import Iterable
+"""Defines types to define the range of execution of a kernel. The types are
+designed along the lines of classes defined in the SYCL 2020 spec section 4.9.
+"""
 
-from llvmlite import ir as llvmir
-from numba.core import cgutils, errors, types
-from numba.core.datamodel import default_manager
-from numba.extending import intrinsic, overload
+from collections.abc import Iterable
 
-# can't import name because of the circular import
-DPEX_TARGET_NAME = "dpex"
+from numba_dpex.core.exceptions import (
+    UnmatchedNumberOfRangeDimsError,
+    UnsupportedGroupWorkItemSizeError,
+)
 
 
 class Range(tuple):
@@ -79,10 +80,10 @@ def size(self):
         n = len(self)
         if n > 2:
             return self[0] * self[1] * self[2]
-        elif n > 1:
+        if n > 1:
             return self[0] * self[1]
-        else:
-            return self[0]
+
+        return self[0]
 
     @property
     def ndim(self) -> int:
@@ -172,6 +173,22 @@ def __init__(self, global_size, local_size):
                 + "must be of either type Range or Iterable of int's."
             )
 
+        if len(self._local_range) != len(self._global_range):
+            raise UnmatchedNumberOfRangeDimsError(
+                kernel_name="",
+                global_ndims=len(self._global_range),
+                local_ndims=len(self._local_range),
+            )
+
+        for i, _ in enumerate(self._global_range):
+            if self._global_range[i] % self._local_range[i] != 0:
+                raise UnsupportedGroupWorkItemSizeError(
+                    kernel_name="",
+                    dim=i,
+                    work_groups=self._global_range[i],
+                    work_items=self._local_range[i],
+                )
+
     @property
     def global_range(self):
         """Accessor for global_range.
@@ -230,163 +247,5 @@ def __eq__(self, other):
                 self.global_range == other.global_range
                 and self.local_range == other.local_range
             )
-        else:
-            return False
-
-
-@intrinsic(target=DPEX_TARGET_NAME)
-def _intrin_range_alloc(typingctx, ty_dim0, ty_dim1, ty_dim2, ty_range):
-    ty_retty = ty_range.instance_type
-    sig = ty_retty(
-        ty_dim0,
-        ty_dim1,
-        ty_dim2,
-        ty_range,
-    )
-
-    def codegen(context, builder, sig, args):
-        typ = sig.return_type
-        dim0, dim1, dim2, _ = args
-        range_struct = cgutils.create_struct_proxy(typ)(context, builder)
-        range_struct.dim0 = dim0
-
-        if not isinstance(sig.args[1], types.NoneType):
-            range_struct.dim1 = dim1
-        else:
-            range_struct.dim1 = llvmir.Constant(
-                llvmir.types.IntType(64), Range.UNDEFINED_DIMENSION
-            )
-
-        if not isinstance(sig.args[2], types.NoneType):
-            range_struct.dim2 = dim2
-        else:
-            range_struct.dim2 = llvmir.Constant(
-                llvmir.types.IntType(64), Range.UNDEFINED_DIMENSION
-            )
-
-        range_struct.ndim = llvmir.Constant(llvmir.types.IntType(64), typ.ndim)
-
-        return range_struct._getvalue()
-
-    return sig, codegen
-
-
-@intrinsic(target=DPEX_TARGET_NAME)
-def _intrin_ndrange_alloc(
-    typingctx, ty_global_range, ty_local_range, ty_ndrange
-):
-    ty_retty = ty_ndrange.instance_type
-    sig = ty_retty(
-        ty_global_range,
-        ty_local_range,
-        ty_ndrange,
-    )
-    range_datamodel = default_manager.lookup(ty_global_range)
-
-    def codegen(context, builder, sig, args):
-        typ = sig.return_type
-
-        global_range, local_range, _ = args
-        ndrange_struct = cgutils.create_struct_proxy(typ)(context, builder)
-        ndrange_struct.ndim = llvmir.Constant(
-            llvmir.types.IntType(64), typ.ndim
-        )
-        ndrange_struct.gdim0 = builder.extract_value(
-            global_range,
-            range_datamodel.get_field_position("dim0"),
-        )
-        ndrange_struct.gdim1 = builder.extract_value(
-            global_range,
-            range_datamodel.get_field_position("dim1"),
-        )
-        ndrange_struct.gdim2 = builder.extract_value(
-            global_range,
-            range_datamodel.get_field_position("dim2"),
-        )
-        ndrange_struct.ldim0 = builder.extract_value(
-            local_range,
-            range_datamodel.get_field_position("dim0"),
-        )
-        ndrange_struct.ldim1 = builder.extract_value(
-            local_range,
-            range_datamodel.get_field_position("dim1"),
-        )
-        ndrange_struct.ldim2 = builder.extract_value(
-            local_range,
-            range_datamodel.get_field_position("dim2"),
-        )
-
-        return ndrange_struct._getvalue()
-
-    return sig, codegen
-
-
-@overload(Range, target=DPEX_TARGET_NAME)
-def _ol_range_init(dim0, dim1=None, dim2=None):
-    """Numba overload of the Range constructor to make it usable inside an
-    njit and dpjit decorated function.
-
-    """
-    from numba_dpex.core.types import RangeType
-
-    ndims = 1
-    ty_optional_dims = (dim1, dim2)
-
-    # A Range should at least have the 0th dimension populated
-    if not isinstance(dim0, types.Integer):
-        raise errors.TypingError(
-            "Expected a Range's dimension should to be an Integer value, but "
-            "encountered " + dim0.name
-        )
-
-    for ty_dim in ty_optional_dims:
-        if isinstance(ty_dim, types.Integer):
-            ndims += 1
-        elif ty_dim is not None:
-            raise errors.TypingError(
-                "Expected a Range's dimension to be an Integer value, "
-                f"but {type(ty_dim)} was provided."
-            )
-
-    ret_ty = RangeType(ndims)
-
-    def impl(dim0, dim1=None, dim2=None):
-        return _intrin_range_alloc(dim0, dim1, dim2, ret_ty)
-
-    return impl
-
-
-@overload(NdRange, target=DPEX_TARGET_NAME)
-def _ol_ndrange_init(global_range, local_range):
-    """Numba overload of the NdRange constructor to make it usable inside an
-    njit and dpjit decorated function.
-
-    """
-    from numba_dpex.core.exceptions import UnmatchedNumberOfRangeDimsError
-    from numba_dpex.core.types import NdRangeType, RangeType
-
-    if not isinstance(global_range, RangeType):
-        raise errors.TypingError(
-            "Only global range values specified as a Range are "
-            "supported inside dpjit"
-        )
-
-    if not isinstance(local_range, RangeType):
-        raise errors.TypingError(
-            "Only local range values specified as a Range are "
-            "supported inside dpjit"
-        )
-
-    if not global_range.ndim == local_range.ndim:
-        raise UnmatchedNumberOfRangeDimsError(
-            kernel_name="",
-            global_ndims=global_range.ndim,
-            local_ndims=local_range.ndim,
-        )
-
-    ret_ty = NdRangeType(global_range.ndim)
-
-    def impl(global_range, local_range):
-        return _intrin_ndrange_alloc(global_range, local_range, ret_ty)
 
-    return impl
+        return False
diff --git a/numba_dpex/kernel_api_impl/__init__.py b/numba_dpex/kernel_api_impl/__init__.py
new file mode 100644
index 0000000000..d0aa0cd6c2
--- /dev/null
+++ b/numba_dpex/kernel_api_impl/__init__.py
@@ -0,0 +1,7 @@
+# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""The module stores the numba_dpex backends implementing the target-specific
+code generation for the kernel_api Python functions.
+"""
diff --git a/numba_dpex/kernel_api_impl/spirv/__init__.py b/numba_dpex/kernel_api_impl/spirv/__init__.py
new file mode 100644
index 0000000000..1b3ce5213b
--- /dev/null
+++ b/numba_dpex/kernel_api_impl/spirv/__init__.py
@@ -0,0 +1,6 @@
+# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""A SPIR-V backend to compile the numba_dpex.kernel_api functions to SPIR-V.
+"""
diff --git a/numba_dpex/core/codegen.py b/numba_dpex/kernel_api_impl/spirv/codegen.py
similarity index 86%
rename from numba_dpex/core/codegen.py
rename to numba_dpex/kernel_api_impl/spirv/codegen.py
index f21be104ba..a0fc7739d0 100644
--- a/numba_dpex/core/codegen.py
+++ b/numba_dpex/kernel_api_impl/spirv/codegen.py
@@ -2,6 +2,9 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
+"""Code generator for a LLVM module for SPIR-V kernels.
+"""
+
 import warnings
 
 from llvmlite import binding as ll
@@ -26,6 +29,10 @@
 
 
 class SPIRVCodeLibrary(CPUCodeLibrary):
+    """A Numba code library that stores a spir_kernel function and all the
+    internally defined spir_func functions called from the spir_kernel function.
+    """
+
     def _optimize_functions(self, ll_module):
         pass
 
@@ -36,8 +43,8 @@ def inline_threshold(self):
         """
         if hasattr(self, "_inline_threshold"):
             return self._inline_threshold
-        else:
-            return 0
+
+        return 0
 
     @inline_threshold.setter
     def inline_threshold(self, value: int):
@@ -105,6 +112,7 @@ def get_asm_str(self):
 
     @property
     def final_module(self):
+        """Return the final SPIR-V module after it has been finalized."""
         return self._final_module
 
 
@@ -116,7 +124,7 @@ class JITSPIRVCodegen(CPUCodegen):
     _library_class = SPIRVCodeLibrary
 
     def _init(self, llvm_module):
-        assert list(llvm_module.global_variables) == [], "Module isn't empty"
+        assert not list(llvm_module.global_variables), "Module isn't empty"
         self._data_layout = SPIR_DATA_LAYOUT[utils.MACHINE_BITS]
         self._target_data = ll.create_target_data(self._data_layout)
         self._tm_features = (
@@ -130,10 +138,15 @@ def _create_empty_module(self, name):
             ir_module.data_layout = self._data_layout
         return ir_module
 
-    def _module_pass_manager(self):
+    def create_empty_spirv_module(self, name):
+        """Public method to create an empty LLVM Module with SPIR-V layout."""
+
+        return self._create_empty_module(name)
+
+    def _module_pass_manager(self, **kwargs):
         raise NotImplementedError
 
-    def _function_pass_manager(self, llvm_module):
+    def _function_pass_manager(self, llvm_module, **kwargs):
         raise NotImplementedError
 
     def _add_module(self, module):
diff --git a/numba_dpex/experimental/kernel_dispatcher.py b/numba_dpex/kernel_api_impl/spirv/dispatcher.py
similarity index 93%
rename from numba_dpex/experimental/kernel_dispatcher.py
rename to numba_dpex/kernel_api_impl/spirv/dispatcher.py
index 13a539c970..056c9ffc11 100644
--- a/numba_dpex/experimental/kernel_dispatcher.py
+++ b/numba_dpex/kernel_api_impl/spirv/dispatcher.py
@@ -25,8 +25,7 @@
 from numba.core.types import void
 from numba.core.typing.typeof import Purpose, typeof
 
-from numba_dpex import config, numba_sem_version, spirv_generator
-from numba_dpex.core.codegen import SPIRVCodeLibrary
+from numba_dpex import config, numba_sem_version
 from numba_dpex.core.exceptions import (
     ExecutionQueueInferenceError,
     InvalidKernelSpecializationError,
@@ -34,21 +33,25 @@
     UnsupportedKernelArgumentError,
 )
 from numba_dpex.core.pipelines import kernel_compiler
-from numba_dpex.core.targets.kernel_target import (
-    CompilationMode,
-    DpexKernelTargetContext,
-)
 from numba_dpex.core.types import USMNdArray
 from numba_dpex.core.utils import kernel_launcher as kl
+from numba_dpex.experimental.target import (
+    DPEX_KERNEL_EXP_TARGET_NAME,
+    dpex_exp_kernel_target,
+)
+from numba_dpex.kernel_api_impl.spirv import spirv_generator
+from numba_dpex.kernel_api_impl.spirv.codegen import SPIRVCodeLibrary
+from numba_dpex.kernel_api_impl.spirv.target import (
+    CompilationMode,
+    SPIRVTargetContext,
+)
 
-from .target import DPEX_KERNEL_EXP_TARGET_NAME, dpex_exp_kernel_target
-
-_KernelCompileResult = namedtuple(
+_SPIRVKernelCompileResult = namedtuple(
     "_KernelCompileResult", CompileResult._fields + ("kernel_device_ir_module",)
 )
 
 
-class _KernelCompiler(_FunctionCompiler):
+class _SPIRVKernelCompiler(_FunctionCompiler):
     """A special compiler class used to compile numba_dpex.kernel decorated
     functions.
     """
@@ -155,7 +158,7 @@ def _compile_to_spirv(
         self,
         kernel_library: SPIRVCodeLibrary,
         kernel_fndesc: PythonFunctionDescriptor,
-        kernel_targetctx: DpexKernelTargetContext,
+        kernel_targetctx: SPIRVTargetContext,
     ):
         kernel_func: ValueRef = kernel_library.get_function(
             kernel_fndesc.llvm_func_name
@@ -188,7 +191,7 @@ def _compile_to_spirv(
             kernel_name=kernel_fn.name, kernel_bitcode=kernel_spirv_module
         )
 
-    def compile(self, args, return_type) -> _KernelCompileResult:
+    def compile(self, args, return_type) -> _SPIRVKernelCompileResult:
         status, kcres = self._compile_cached(args, return_type)
         if status:
             return kcres
@@ -197,7 +200,7 @@ def compile(self, args, return_type) -> _KernelCompileResult:
 
     def _compile_cached(
         self, args, return_type: types.Type
-    ) -> Tuple[bool, _KernelCompileResult]:
+    ) -> Tuple[bool, _SPIRVKernelCompileResult]:
         """Compiles the kernel function to bitcode and generates a host-callable
         wrapper to submit the kernel to a SYCL queue.
 
@@ -271,16 +274,16 @@ def _compile_cached(
                     "w",
                     encoding="UTF-8",
                 ) as fptr:
-                    fptr.write(cres.library.final_module)
+                    fptr.write(str(cres.library.final_module))
 
         except errors.TypingError as err:
             self._failed_cache[key] = err
             return False, err
 
-        return True, _KernelCompileResult(*kcres_attrs)
+        return True, _SPIRVKernelCompileResult(*kcres_attrs)
 
 
-class KernelDispatcher(Dispatcher):
+class SPIRVKernelDispatcher(Dispatcher):
     """Dispatcher class designed to compile kernel decorated functions. The
     dispatcher inherits the Numba Dispatcher class, but has a different
     compilation strategy. Instead of compiling a kernel decorated function to
@@ -311,6 +314,7 @@ def __init__(
         self._kernel_name = pyfunc.__name__
 
         if numba_sem_version < (0, 59, 0):
+            # pylint: disable=unexpected-keyword-arg
             super().__init__(
                 py_func=pyfunc,
                 locals=local_vars_to_numba_types,
@@ -325,7 +329,7 @@ def __init__(
                 targetoptions=targetoptions,
                 pipeline_class=pipeline_class,
             )
-        self._compiler = _KernelCompiler(
+        self._compiler = _SPIRVKernelCompiler(
             pyfunc,
             self.targetdescr,
             targetoptions,
@@ -426,8 +430,8 @@ def cb_llvm(dur):
                     },
                 ):
                     try:
-                        compiler: _KernelCompiler = self._compiler
-                        kcres: _KernelCompileResult = compiler.compile(
+                        compiler: _SPIRVKernelCompiler = self._compiler
+                        kcres: _SPIRVKernelCompileResult = compiler.compile(
                             args, return_type
                         )
                     except errors.ForceLiteralArg as err:
@@ -463,4 +467,4 @@ def __call__(self, *args, **kw_args):
 
 
 _dpex_target = target_registry[DPEX_KERNEL_EXP_TARGET_NAME]
-dispatcher_registry[_dpex_target] = KernelDispatcher
+dispatcher_registry[_dpex_target] = SPIRVKernelDispatcher
diff --git a/numba_dpex/spirv_generator.py b/numba_dpex/kernel_api_impl/spirv/spirv_generator.py
similarity index 84%
rename from numba_dpex/spirv_generator.py
rename to numba_dpex/kernel_api_impl/spirv/spirv_generator.py
index 7f5c9931c6..1171faca4a 100644
--- a/numba_dpex/spirv_generator.py
+++ b/numba_dpex/kernel_api_impl/spirv/spirv_generator.py
@@ -13,76 +13,50 @@
 
 from numba_dpex.core import config
 from numba_dpex.core.exceptions import InternalError
-from numba_dpex.core.targets.kernel_target import LLVM_SPIRV_ARGS
+from numba_dpex.kernel_api_impl.spirv.target import LLVM_SPIRV_ARGS
+
+try:
+    import dpcpp_llvm_spirv as dls
+except ImportError as err:
+    raise ImportError("Cannot import dpcpp-llvm-spirv package") from err
 
 
 def run_cmd(args, error_message=None):
+    """
+    Helper to run an external command and provide a meaningful error message.
+    """
     try:
         check_output(
             args,
             stderr=STDOUT,
         )
-    except CalledProcessError as err:
+    except CalledProcessError as cper:
         if error_message is None:
             error_message = f"Error during call to {args[0]}"
         raise InternalError(
             f"{error_message}:\n\t"
-            + "\t".join(err.output.decode("utf-8").splitlines(True))
-        )
-
-
-class _SpirvGenerator:
-    """Generates a SPIR-V binary from supplied LLVM IR."""
-
-    def generate(self, llvm_spirv_args, ipath, opath):
-        """
-        Generate a spirv module from llvm bitcode.
-
-        We use llvm-spirv tool to translate a llvm bitcode into spirv.
-
-        Args:
-            llvm_spirv_args: Args to be provided to llvm-spirv tool.
-            ipath: Input file path of the llvm bitcode.
-            opath: Output file path of the generated spirv.
-        """
-        llvm_spirv_flags = []
-        if config.DEBUG:
-            llvm_spirv_flags.append("--spirv-debug-info-version=ocl-100")
+            + "\t".join(cper.output.decode("utf-8").splitlines(True))
+        ) from cper
 
-        llvm_spirv_tool = self._llvm_spirv()
 
-        if config.DEBUG:
-            print(f"Use llvm-spirv: {llvm_spirv_tool}")
-
-        run_cmd(
-            [llvm_spirv_tool, *llvm_spirv_args, "-o", opath, ipath],
-            error_message="Error during lowering LLVM IR to SPIRV",
-        )
+class Module:
+    """
+    Abstracts a SPIR-V binary module that is created by calling
+    ``llvm-spirv`` translator on a LLVM IR binary module.
+    """
 
     @staticmethod
     def _llvm_spirv():
         """Return path to llvm-spirv executable."""
 
-        try:
-            import dpcpp_llvm_spirv as dls
-        except ImportError:
-            raise ImportError("Cannot import dpcpp-llvm-spirv package")
+        return dls.get_llvm_spirv_path()
 
-        result = dls.get_llvm_spirv_path()
-        return result
-
-
-class Module(object):
     def __init__(self, context, llvmir, llvmbc):
-        """
-        Setup
-        """
         self._tmpdir = tempfile.mkdtemp()
         self._tempfiles = []
-        self._generator = _SpirvGenerator()
         self._finalized = False
         self.context = context
-
+        self._llvmfile = None
         self._llvmir = llvmir
         self._llvmbc = llvmbc
 
@@ -94,12 +68,35 @@ def __del__(self):
         os.rmdir(self._tmpdir)
 
     def _track_temp_file(self, name):
-        path = os.path.join(
-            self._tmpdir, "{0}-{1}".format(len(self._tempfiles), name)
-        )
+        path = os.path.join(self._tmpdir, f"{len(self._tempfiles)}-{name}")
         self._tempfiles.append(path)
         return path
 
+    def _generate_spirv(self, llvm_spirv_args, ipath, opath):
+        """
+        Generate a spirv module from llvm bitcode.
+
+        We use llvm-spirv tool to translate a llvm bitcode into spirv.
+
+        Args:
+            llvm_spirv_args: Args to be provided to llvm-spirv tool.
+            ipath: Input file path of the llvm bitcode.
+            opath: Output file path of the generated spirv.
+        """
+        llvm_spirv_flags = []
+        if config.DEBUG:
+            llvm_spirv_flags.append("--spirv-debug-info-version=ocl-100")
+
+        llvm_spirv_tool = self._llvm_spirv()
+
+        if config.DEBUG:
+            print(f"Use llvm-spirv: {llvm_spirv_tool}")
+
+        run_cmd(
+            [llvm_spirv_tool, *llvm_spirv_args, "-o", opath, ipath],
+            error_message="Error during lowering LLVM IR to SPIRV",
+        )
+
     def load_llvm(self):
         """
         Load LLVM with "SPIR-V friendly" SPIR 2.0 spec
@@ -134,7 +131,7 @@ def finalize(self):
 
         if config.SAVE_IR_FILES != 0:
             # Dump the llvmir and llvmbc in file
-            with open("generated_llvm.ir", "w") as f:
+            with open("generated_llvm.ir", "w", encoding="utf-8") as f:
                 f.write(self._llvmir)
             with open("generated_llvm.bc", "wb") as f:
                 f.write(self._llvmbc)
@@ -146,7 +143,7 @@ def finalize(self):
             print("generated_llvm.bc")
             print("".center(80, "="))
 
-        self._generator.generate(
+        self._generate_spirv(
             llvm_spirv_args=llvm_spirv_args,
             ipath=self._llvmfile,
             opath=spirv_path,
diff --git a/numba_dpex/core/targets/kernel_target.py b/numba_dpex/kernel_api_impl/spirv/target.py
similarity index 77%
rename from numba_dpex/core/targets/kernel_target.py
rename to numba_dpex/kernel_api_impl/spirv/target.py
index 33d32fc9e1..4e51b9b8fd 100644
--- a/numba_dpex/core/targets/kernel_target.py
+++ b/numba_dpex/kernel_api_impl/spirv/target.py
@@ -2,6 +2,9 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
+"""Implements a SPIR-V code generation-specific target and typing context.
+"""
+
 from enum import IntEnum
 from functools import cached_property
 
@@ -9,21 +12,27 @@
 from llvmlite import binding as ll
 from llvmlite import ir as llvmir
 from numba import typeof
-from numba.core import cgutils, funcdesc, types, typing, utils
+from numba.core import cgutils, funcdesc
+from numba.core import types as nb_types
+from numba.core import typing, utils
 from numba.core.base import BaseContext
 from numba.core.callconv import MinimalCallConv
-from numba.core.registry import cpu_target
 from numba.core.target_extension import GPU, target_registry
 from numba.core.types import Array as NpArrayType
+from numba.core.types.scalars import IntEnumClass
+from numba.core.typing import cmathdecl, enumdecl, npydecl
 
 from numba_dpex.core.datamodel.models import _init_data_model_manager
 from numba_dpex.core.exceptions import UnsupportedKernelArgumentError
 from numba_dpex.core.typeconv import to_usm_ndarray
-from numba_dpex.core.types import USMNdArray
+from numba_dpex.core.types import IntEnumLiteral, USMNdArray
+from numba_dpex.core.typing import dpnpdecl
 from numba_dpex.core.utils import get_info_from_suai
+from numba_dpex.kernel_api.flag_enum import FlagEnum
+from numba_dpex.ocl.mathimpl import lower_ocl_impl, sig_mapper
 from numba_dpex.utils import address_space, calling_conv
 
-from .. import codegen
+from . import codegen
 
 CC_SPIR_KERNEL = "spir_kernel"
 CC_SPIR_FUNC = "spir_func"
@@ -52,7 +61,7 @@ class CompilationMode(IntEnum):
     DEVICE_FUNC = 2
 
 
-class DpexKernelTypingContext(typing.BaseContext):
+class SPIRVTypingContext(typing.BaseContext):
     """Custom typing context to support kernel compilation.
 
     The customized typing context provides two features required to compile
@@ -64,6 +73,37 @@ class DpexKernelTypingContext(typing.BaseContext):
 
     """
 
+    def resolve_value_type(self, val):
+        """
+        Return the numba type of a Python value that is being used
+        as a runtime constant.
+        ValueError is raised for unsupported types.
+        """
+
+        typ = super().resolve_value_type(val)
+
+        if isinstance(typ, IntEnumClass) and issubclass(val, FlagEnum):
+            typ = IntEnumLiteral(val)
+
+        return typ
+
+    def resolve_getattr(self, typ, attr):
+        """
+        Resolve getting the attribute *attr* (a string) on the Numba type.
+        The attribute's type is returned, or None if resolution failed.
+        """
+        retty = None
+
+        if isinstance(typ, IntEnumLiteral):
+            try:
+                attrval = getattr(typ.literal_value, attr).value
+                retty = nb_types.IntegerLiteral(attrval)
+            except ValueError:
+                pass
+        else:
+            retty = super().resolve_getattr(typ, attr)
+        return retty
+
     def resolve_argument_type(self, val):
         """Return the Numba type of a Python value used as a function argument.
 
@@ -100,20 +140,17 @@ def resolve_argument_type(self, val):
             try:
                 suai_attrs = get_info_from_suai(val)
                 return to_usm_ndarray(suai_attrs)
-            except Exception:
+            except Exception as err:
                 raise UnsupportedKernelArgumentError(
                     type=str(type(val)), value=val
-                )
+                ) from err
 
         return super().resolve_argument_type(val)
 
     def load_additional_registries(self):
         """Register the OpenCL API and math and other functions."""
-        from numba.core.typing import cmathdecl, enumdecl, npydecl
-
-        from numba_dpex.core.typing import dpnpdecl
-
-        from ...ocl import mathdecl, ocldecl
+        # pylint: disable=import-outside-toplevel
+        from numba_dpex.ocl import mathdecl, ocldecl
 
         self.install_registry(ocldecl.registry)
         self.install_registry(mathdecl.registry)
@@ -124,20 +161,19 @@ def load_additional_registries(self):
         self.install_registry(enumdecl.registry)
 
 
-class SyclDevice(GPU):
-    """Mark the hardware target as SYCL Device."""
-
-    pass
+# pylint: disable=too-few-public-methods
+class SPIRVDevice(GPU):
+    """Mark the hardware target as device that supports SPIR-V bitcode."""
 
 
-DPEX_KERNEL_TARGET_NAME = "dpex_kernel"
+SPIRV_TARGET_NAME = "spirv"
 
-target_registry[DPEX_KERNEL_TARGET_NAME] = SyclDevice
+target_registry[SPIRV_TARGET_NAME] = SPIRVDevice
 
 
-class DpexKernelTargetContext(BaseContext):
+class SPIRVTargetContext(BaseContext):
     """A target context inheriting Numba's ``BaseContext`` that is customized
-    for generating SYCL kernels.
+    for generating SPIR-V kernels.
 
     A customized target context for generating SPIR-V kernels. The class defines
     helper functions to generates SPIR-V kernels as LLVM IR using the required
@@ -213,12 +249,14 @@ def _generate_spir_kernel_wrapper(self, func, argtypes):
         wrapperfnty = llvmir.FunctionType(
             llvmir.VoidType(), arginfo.argument_types
         )
-        wrapper_module = self.create_module("dpex.kernel.wrapper")
+        wrapper_module = self._internal_codegen.create_empty_spirv_module(
+            "dpex.kernel.wrapper"
+        )
         wrappername = func.name.replace("dpex_fn", "dpex_kernel")
         argtys = list(arginfo.argument_types)
         fnty = llvmir.FunctionType(
             llvmir.IntType(32),
-            [self.call_conv.get_return_type(types.pyobject)] + argtys,
+            [self.call_conv.get_return_type(nb_types.pyobject)] + argtys,
         )
         func = llvmir.Function(wrapper_module, fnty, name=func.name)
         func.calling_convention = CC_SPIR_FUNC
@@ -228,8 +266,8 @@ def _generate_spir_kernel_wrapper(self, func, argtypes):
         callargs = arginfo.from_arguments(builder, wrapper.args)
 
         # XXX handle error status
-        status, _ = self.call_conv.call_function(
-            builder, func, types.void, argtypes, callargs
+        self.call_conv.call_function(
+            builder, func, nb_types.void, argtypes, callargs
         )
         builder.ret_void()
 
@@ -243,13 +281,16 @@ def _generate_spir_kernel_wrapper(self, func, argtypes):
         module.get_function(func.name).linkage = "internal"
         return wrapper
 
-    def __init__(self, typingctx, target=DPEX_KERNEL_TARGET_NAME):
+    def __init__(self, typingctx, target=SPIRV_TARGET_NAME):
         super().__init__(typingctx, target)
 
     def init(self):
         """Called by the super().__init__ constructor to initalize the child
         class.
         """
+        # pylint: disable=import-outside-toplevel
+        from numba_dpex.dpnp_iface.dpnp_ufunc_db import _lazy_init_dpnp_db
+
         self._internal_codegen = codegen.JITSPIRVCodegen("numba_dpex.kernel")
         self._target_data = ll.create_target_data(
             codegen.SPIR_DATA_LAYOUT[utils.MACHINE_BITS]
@@ -257,9 +298,7 @@ def init(self):
 
         # Override data model manager to SPIR model
         self.data_model_manager = _init_data_model_manager()
-        self.extra_compile_options = dict()
-
-        from numba_dpex.dpnp_iface.dpnp_ufunc_db import _lazy_init_dpnp_db
+        self.extra_compile_options = {}
 
         _lazy_init_dpnp_db()
 
@@ -269,12 +308,29 @@ def init(self):
 
         self.ufunc_db = _dpnp_ufunc_db
 
+    def get_getattr(self, typ, attr):
+        """
+        Overrides the get_getattr function to provide an implementation for
+        getattr call on an IntegerEnumLiteral type.
+        """
+
+        if isinstance(typ, IntEnumLiteral):
+            #  pylint: disable=W0613
+            def enum_literal_getattr_imp(context, builder, typ, val, attr):
+                enum_attr_value = getattr(typ.literal_value, attr).value
+                return llvmir.Constant(llvmir.IntType(64), enum_attr_value)
+
+            return enum_literal_getattr_imp
+
+        return super().get_getattr(typ, attr)
+
     def create_module(self, name):
-        return self._internal_codegen._create_empty_module(name)
+        return self._internal_codegen.create_empty_spirv_module(name)
 
     def replace_dpnp_ufunc_with_ocl_intrinsics(self):
-        from numba_dpex.ocl.mathimpl import lower_ocl_impl, sig_mapper
-
+        """Replaces the implementation in the ufunc_db for specific math
+        functions for which a SPIR-V intrinsic should be used.
+        """
         ufuncs = [
             ("fabs", dpnp.fabs),
             ("exp", dpnp.exp),
@@ -324,10 +380,10 @@ def load_additional_registries(self):
         target context.
 
         """
+        # pylint: disable=import-outside-toplevel
+        from numba_dpex import printimpl
         from numba_dpex.dpnp_iface import dpnpimpl
-
-        from ... import printimpl
-        from ...ocl import mathimpl, oclimpl
+        from numba_dpex.ocl import mathimpl, oclimpl
 
         self.insert_func_defn(oclimpl.registry.functions)
         self.insert_func_defn(mathimpl.registry.functions)
@@ -338,27 +394,40 @@ def load_additional_registries(self):
 
     @cached_property
     def call_conv(self):
-        return DpexCallConv(self)
+        """
+        Return the CallConv object used by the SPIRVTargetContext.
+        """
+        return SPIRVCallConv(self)
 
     def codegen(self):
+        """Return the CodeGen object used by the SPIRVTargetContext."""
         return self._internal_codegen
 
     @property
     def target_data(self):
         return self._target_data
 
-    def mangler(self, name, argtypes, abi_tags=(), uid=None):
+    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", argtypes, abi_tags=abi_tags, uid=uid
+            name + "dpex_fn", types, abi_tags=abi_tags, uid=uid
         )
 
     def prepare_spir_kernel(self, func, argtypes):
+        """Generates a wrapper function with \"spir_kernel\" calling conv that
+        calls the compiled \"spir_func\" generated by numba_dpex for a kernel
+        decorated function.
+        """
         func.linkage = "linkonce_odr"
         func.module.data_layout = codegen.SPIR_DATA_LAYOUT[self.address_size]
         wrapper = self._generate_spir_kernel_wrapper(func, argtypes)
         return wrapper
 
     def set_spir_func_calling_conv(self, func):
+        """Sets the calling convetion of the provided LLVM func to spir_func"""
         # Adapt to SPIR
         func.calling_convention = CC_SPIR_FUNC
         func.linkage = "linkonce_odr"
@@ -385,9 +454,7 @@ def declare_function(self, module, fndesc):
         )
         if not self.enable_debuginfo:
             fn.attributes.add("alwaysinline")
-        ret = super(DpexKernelTargetContext, self).declare_function(
-            module, fndesc
-        )
+        ret = super().declare_function(module, fndesc)
         ret.calling_convention = calling_conv.CC_SPIR_FUNC
         return ret
 
@@ -439,20 +506,26 @@ def populate_array(self, arr, **kwargs):
         """
         Populate array structure.
         """
+        # pylint: disable=import-outside-toplevel
         from numba_dpex.core.kernel_interface import arrayobj
 
         return arrayobj.populate_array(arr, **kwargs)
 
+    def get_executable(self, func, fndesc, env):
+        """Not implemented for SPIRVTargetContext"""
+        raise NotImplementedError("Not implemented for SPIRVTargetContext")
+
 
-class DpexCallConv(MinimalCallConv):
+class SPIRVCallConv(MinimalCallConv):
     """Custom calling convention class used by numba-dpex.
 
     numba_dpex's calling convention derives from
     :class:`numba.core.callconv import MinimalCallConv`. The
-    :class:`DpexCallConv` overrides :func:`call_function`.
+    :class:`SPIRVCallConv` overrides :func:`call_function`.
 
     """
 
+    # pylint:disable=too-many-arguments
     def call_function(self, builder, callee, resty, argtys, args, env=None):
         """Call the Numba-compiled *callee*."""
         assert env is None
diff --git a/numba_dpex/numba_patches/patch_mk_alloc.py b/numba_dpex/numba_patches/patch_mk_alloc.py
deleted file mode 100644
index 0f4f8f195e..0000000000
--- a/numba_dpex/numba_patches/patch_mk_alloc.py
+++ /dev/null
@@ -1,175 +0,0 @@
-# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-
-def patch():
-    """
-    Patches the numba.core.ir_utils.mk_alloc function to support non-NumPy array
-    types.
-
-    The patch extends the ir_utils.mk_alloc function that is used by
-    numba.parfors.parfor to allocate an empty temp array inside a parfor
-    section. Without the patch Numba always allocates the temp using
-    numpy.empty, with the patch the function checks if the array type to be
-    allocated has a __allocate__ method, and if so uses that method to allocate
-    the temporary array.
-
-    TODO: Already merged into Numba main and should be removed when numba-dpex
-    is ported to Numba 0.58
-    """
-    import numpy
-    from numba.core import ir, ir_utils, types
-    from numba.core.ir_utils import (
-        convert_size_to_var,
-        get_np_ufunc_typ,
-        mk_unique_var,
-    )
-    from numba.parfors import parfor
-
-    def _mk_alloc(
-        typingctx, typemap, calltypes, lhs, size_var, dtype, scope, loc, lhs_typ
-    ):
-        """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.
-        """
-        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,
-            )
-
-        # g_np_var = Global(numpy) # noqa: E800
-        g_np_var = ir.Var(scope, mk_unique_var("$np_g_var"), loc)
-        if typemap:
-            typemap[g_np_var.name] = types.misc.Module(numpy)
-        g_np = ir.Global("np", numpy, loc)
-        g_np_assign = ir.Assign(g_np, g_np_var, loc)
-        # attr call: empty_attr = getattr(g_np_var, empty)
-        empty_attr_call = ir.Expr.getattr(g_np_var, "empty", loc)
-        attr_var = ir.Var(scope, mk_unique_var("$empty_attr_attr"), loc)
-        if typemap:
-            typemap[attr_var.name] = get_np_ufunc_typ(numpy.empty)
-        attr_assign = ir.Assign(empty_attr_call, attr_var, loc)
-        # Assume str(dtype) returns a valid type
-        dtype_str = str(dtype)
-        # alloc call: lhs = empty_attr(size_var, typ_var)
-        typ_var = ir.Var(scope, mk_unique_var("$np_typ_var"), loc)
-        if typemap:
-            typemap[typ_var.name] = types.functions.NumberClass(dtype)
-        # If dtype is a datetime/timedelta with a unit,
-        # then it won't return a valid type and instead can be created
-        # with a string. i.e. "datetime64[ns]")
-        if (
-            isinstance(dtype, (types.NPDatetime, types.NPTimedelta))
-            and dtype.unit != ""
-        ):
-            typename_const = ir.Const(dtype_str, loc)
-            typ_var_assign = ir.Assign(typename_const, typ_var, loc)
-        else:
-            if dtype_str == "bool":
-                # empty doesn't like 'bool' sometimes (e.g. kmeans example)
-                dtype_str = "bool_"
-            np_typ_getattr = ir.Expr.getattr(g_np_var, dtype_str, loc)
-            typ_var_assign = ir.Assign(np_typ_getattr, typ_var, loc)
-        alloc_call = ir.Expr.call(attr_var, [size_var, typ_var], (), loc)
-
-        if calltypes:
-            cac = typemap[attr_var.name].get_call_type(
-                typingctx, [size_typ, types.functions.NumberClass(dtype)], {}
-            )
-            # By default, all calls to "empty" are typed as returning a standard
-            # NumPy ndarray.  If we are allocating a ndarray subclass here then
-            # just change the return type to be that of the subclass.
-            cac._return_type = (
-                lhs_typ.copy(layout="C") if lhs_typ.layout == "F" else lhs_typ
-            )
-            calltypes[alloc_call] = cac
-        if lhs_typ.layout == "F":
-            empty_c_typ = lhs_typ.copy(layout="C")
-            empty_c_var = ir.Var(scope, mk_unique_var("$empty_c_var"), loc)
-            if typemap:
-                typemap[empty_c_var.name] = lhs_typ.copy(layout="C")
-            empty_c_assign = ir.Assign(alloc_call, empty_c_var, loc)
-
-            # attr call: asfortranarray = getattr(g_np_var, asfortranarray)
-            asfortranarray_attr_call = ir.Expr.getattr(
-                g_np_var, "asfortranarray", loc
-            )
-            afa_attr_var = ir.Var(
-                scope, mk_unique_var("$asfortran_array_attr"), loc
-            )
-            if typemap:
-                typemap[afa_attr_var.name] = get_np_ufunc_typ(
-                    numpy.asfortranarray
-                )
-            afa_attr_assign = ir.Assign(
-                asfortranarray_attr_call, afa_attr_var, loc
-            )
-            # call asfortranarray
-            asfortranarray_call = ir.Expr.call(
-                afa_attr_var, [empty_c_var], (), loc
-            )
-            if calltypes:
-                calltypes[asfortranarray_call] = typemap[
-                    afa_attr_var.name
-                ].get_call_type(typingctx, [empty_c_typ], {})
-
-            asfortranarray_assign = ir.Assign(asfortranarray_call, lhs, loc)
-
-            out.extend(
-                [
-                    g_np_assign,
-                    attr_assign,
-                    typ_var_assign,
-                    empty_c_assign,
-                    afa_attr_assign,
-                    asfortranarray_assign,
-                ]
-            )
-        else:
-            alloc_assign = ir.Assign(alloc_call, lhs, loc)
-            out.extend([g_np_assign, attr_assign, typ_var_assign, alloc_assign])
-
-        return out
-
-    ir_utils.mk_alloc = _mk_alloc
-    parfor.mk_alloc = _mk_alloc
diff --git a/numba_dpex/ocl/_declare_function.py b/numba_dpex/ocl/_declare_function.py
index 2bb63b9557..4a38e7faa7 100644
--- a/numba_dpex/ocl/_declare_function.py
+++ b/numba_dpex/ocl/_declare_function.py
@@ -42,7 +42,7 @@ def _declare_function(context, builder, name, sig, cargs, mangler=mangle_c):
     fnty = llvmir.FunctionType(llretty, llargs)
     mangled = mangler(name, cargs)
     fn = cgutils.get_or_insert_function(mod, fnty, mangled)
-    from numba_dpex import kernel_target
+    from numba_dpex import spirv_kernel_target
 
-    fn.calling_convention = kernel_target.CC_SPIR_FUNC
+    fn.calling_convention = spirv_kernel_target.CC_SPIR_FUNC
     return fn
diff --git a/numba_dpex/ocl/oclimpl.py b/numba_dpex/ocl/oclimpl.py
index 727ed632da..fe7cbca6d6 100644
--- a/numba_dpex/ocl/oclimpl.py
+++ b/numba_dpex/ocl/oclimpl.py
@@ -12,10 +12,10 @@
 from numba.core.imputils import Registry
 from numba.core.typing.npydecl import parse_dtype
 
-from numba_dpex import kernel_target
+from numba_dpex import spirv_kernel_target
 from numba_dpex.core import config
-from numba_dpex.core.codegen import SPIR_DATA_LAYOUT
 from numba_dpex.core.types import Array
+from numba_dpex.kernel_api_impl.spirv.codegen import SPIR_DATA_LAYOUT
 from numba_dpex.ocl.atomics import atomic_helper
 from numba_dpex.utils import address_space
 
@@ -171,7 +171,7 @@ def native_atomic_add(context, builder, sig, args):
     ptr = cgutils.get_item_pointer(context, builder, aryty, lary, indices)
 
     if dtype == types.float32 or dtype == types.float64:
-        context.extra_compile_options[kernel_target.LLVM_SPIRV_ARGS] = [
+        context.extra_compile_options[spirv_kernel_target.LLVM_SPIRV_ARGS] = [
             "--spirv-ext=+SPV_EXT_shader_atomic_float_add"
         ]
         name = "__spirv_AtomicFAddEXT"
@@ -208,7 +208,7 @@ def native_atomic_add(context, builder, sig, args):
 
     fnty = llvmir.FunctionType(retty, spirv_fn_arg_types)
     fn = cgutils.get_or_insert_function(builder.module, fnty, mangled_fn_name)
-    fn.calling_convention = kernel_target.CC_SPIR_FUNC
+    fn.calling_convention = spirv_kernel_target.CC_SPIR_FUNC
 
     sycl_memory_order = atomic_helper.sycl_memory_order.relaxed
     sycl_memory_scope = atomic_helper.sycl_memory_scope.device
diff --git a/numba_dpex/vectorizers.py b/numba_dpex/vectorizers.py
deleted file mode 100644
index 819adeef3d..0000000000
--- a/numba_dpex/vectorizers.py
+++ /dev/null
@@ -1,207 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-"""Provide a Dpex target for Numba's ``vectorize`` decorator."""
-
-import warnings
-
-import dpctl
-import numpy as np
-from numba.np.ufunc import deviceufunc
-
-import numba_dpex as dpex
-from numba_dpex.utils import (
-    as_usm_obj,
-    copy_to_numpy_from_usm_obj,
-    has_usm_memory,
-)
-
-vectorizer_stager_source = """
-def __vectorized_{name}({args}, __out__):
-    __tid__ = __dpex__.get_global_id(0)
-    if __tid__ < __out__.shape[0]:
-        __out__[__tid__] = __core__({argitems})
-"""
-
-
-class Vectorize(deviceufunc.DeviceVectorize):
-    def _compile_core(self, sig):
-        devfn = dpex.func(sig)(self.pyfunc)
-        return devfn, devfn.cres.signature.return_type
-
-    def _get_globals(self, corefn):
-        glbl = self.pyfunc.__globals__.copy()
-        glbl.update({"__dpex__": dpex, "__core__": corefn})
-        return glbl
-
-    def _compile_kernel(self, fnobj, sig):
-        return dpex.kernel(sig)(fnobj)
-
-    def build_ufunc(self):
-        return UFuncDispatcher(self.kernelmap)
-
-    @property
-    def _kernel_template(self):
-        return vectorizer_stager_source
-
-
-class UFuncDispatcher(object):
-    """
-    Invoke the Dpex ufunc specialization for the given inputs.
-    """
-
-    def __init__(self, types_to_retty_kernels):
-        self.functions = types_to_retty_kernels
-
-    def __call__(self, *args, **kws):
-        """
-        Call the kernel launching mechanism
-        Args:
-            *args (np.ndarray): NumPy arrays
-            **kws (optional):
-                queue (dpctl._sycl_queue.SyclQueue): SYCL queue.
-                out (np.ndarray): Output array.
-        """
-        return UFuncMechanism.call(self.functions, args, kws)
-
-    def reduce(self, arg, queue=0):
-        raise NotImplementedError
-
-
-class UFuncMechanism(deviceufunc.UFuncMechanism):
-    """
-    Mechanism to process Input to a SYCL kernel and launch that kernel
-    """
-
-    @classmethod
-    def call(cls, typemap, args, kws):
-        """
-        Perform the entire ufunc call mechanism.
-
-        Args:
-            typemap (dict): Signature mapped to kernel.
-            args: Arguments to the @vectorize function.
-            kws (optional): Optional keywords. Not supported.
-
-        """
-        # Handle keywords
-        queue = dpctl.get_current_queue()
-        out = kws.pop("out", None)
-
-        if kws:
-            warnings.warn("unrecognized keywords: %s" % ", ".join(kws))
-
-        # Begin call resolution
-        cr = cls(typemap, args)
-        args = cr.get_arguments()
-        resty, func = cr.get_function()
-
-        outshape = args[0].shape
-
-        # Adjust output value
-        if out is not None and cr.is_device_array(out):
-            out = cr.as_device_array(out)
-
-        def attempt_ravel(a):
-            if cr.SUPPORT_DEVICE_SLICING:
-                raise NotImplementedError
-
-            try:
-                # Call the `.ravel()` method
-                return a.ravel()
-            except NotImplementedError:
-                # If it is not a device array
-                if not cr.is_device_array(a):
-                    raise
-                # For device array, retry ravel on the host by first
-                # copying it back.
-                else:
-                    hostary = cr.to_host(a, queue).ravel()
-                    return cr.to_device(hostary, queue)
-
-        if args[0].ndim > 1:
-            args = [attempt_ravel(a) for a in args]
-
-        # Prepare argument on the device
-        devarys = []
-        any_device = True
-        for a in args:
-            if cr.is_device_array(a):
-                devarys.append(a)
-            else:
-                dev_a = cr.to_device(a, queue=queue)
-                devarys.append(dev_a)
-
-        # Launch
-        shape = args[0].shape
-        if out is None:
-            # No output is provided
-            devout = cr.device_array(shape, resty, queue=queue)
-
-            devarys.extend([devout])
-            cr.launch(func, shape[0], queue, devarys)
-
-            if any_device:
-                # If any of the arguments are on device,
-                # Keep output on the device
-                return devout.reshape(outshape)
-            else:
-                # Otherwise, transfer output back to host
-                raise ValueError("copy_to_host() is not yet supported")
-
-        elif cr.is_device_array(out):
-            # If output is provided and it is a device array,
-            # Return device array
-            if out.ndim > 1:
-                out = attempt_ravel(out)
-            devout = out
-            devarys.extend([devout])
-            cr.launch(func, shape[0], queue, devarys)
-            return devout.reshape(outshape)
-
-        else:
-            # If output is provided and it is a host array,
-            # Return host array
-            assert out.shape == shape
-            assert out.dtype == resty
-            devout = cr.device_array(shape, resty, queue=queue)
-            devarys.extend([devout])
-            cr.launch(func, shape[0], queue, devarys)
-            return devout.reshape(outshape)
-
-    def as_device_array(self, obj):
-        return obj
-
-    def is_device_array(self, obj):
-        ret = has_usm_memory(obj)
-        return ret is not None
-
-    def is_host_array(self, obj):
-        return not self.is_device_array(obj)
-
-    def to_device(self, hostary, queue):
-        usm_mem = as_usm_obj(hostary, queue=queue, usm_type="shared")
-        usm_backed_ndary = np.ndarray(
-            hostary.shape, buffer=usm_mem, dtype=hostary.dtype
-        )
-        return usm_backed_ndary
-
-    def to_host(self, devary, queue):
-        hostary = np.empty(devary.shape, dtype=devary.dtype)
-        devary_memview = memoryview(devary)
-        devary_memview = devary_memview.cast("B")
-        copy_to_numpy_from_usm_obj(devary_memview, hostary)
-
-    def launch(self, func, count, queue, args):
-        func[dpex.Range(count)](*args)
-
-    def device_array(self, shape, dtype, queue):
-        size = np.prod(shape)
-        itemsize = dtype.itemsize
-        usm_mem = dpctl.memory.MemoryUSMShared(size * itemsize, queue=queue)
-        usm_backed_ndary = np.ndarray(shape, buffer=usm_mem, dtype=dtype)
-        return usm_backed_ndary
-
-    def broadcast_device(self, ary, shape):
-        raise NotImplementedError("device broadcast_device NIY")
diff --git a/pyproject.toml b/pyproject.toml
index 10d25af87c..94bdc408cb 100644
--- a/pyproject.toml
+++ b/pyproject.toml
@@ -57,7 +57,7 @@ requires = [
     "scikit-build>=0.13",
     "cmake>=3.18",
     "ninja",
-    "numba>=0.57",
+    "numba>=0.58",
     "versioneer-518"
 ]
 build-backend = "setuptools.build_meta"
diff --git a/scripts/run_examples.sh b/scripts/run_examples.sh
index b4dff888f4..9255628863 100755
--- a/scripts/run_examples.sh
+++ b/scripts/run_examples.sh
@@ -25,7 +25,6 @@ run_checks() {
   check numba_dpex/examples/sum_reduction.py
   check numba_dpex/examples/sum_reduction_recursive_ocl.py
   # check numba_dpex/examples/usm_ndarray.py  # See https://github.com/IntelPython/numba-dpex/issues/436
-  check numba_dpex/examples/vectorize.py
 
   check numba_dpex/examples/auto_offload_examples/sum-1d.py
   check numba_dpex/examples/auto_offload_examples/sum-2d.py
diff --git a/scripts/set_examples_to_doc.py b/scripts/set_examples_to_doc.py
index 82d6483275..c43b7b6d4c 100644
--- a/scripts/set_examples_to_doc.py
+++ b/scripts/set_examples_to_doc.py
@@ -1,5 +1,6 @@
 """ This script is needed to convert gdb scripts from commands to documentation
 """
+
 import os
 
 
diff --git a/setup.py b/setup.py
index 36f900770c..5f6abc0085 100644
--- a/setup.py
+++ b/setup.py
@@ -80,7 +80,7 @@ def to_cmake_format(version: str):
     platforms=["Linux", "Windows"],
     author="Intel Corporation",
     url="https://github.com/IntelPython/numba-dpex",
-    install_requires=["numba >={0:s}".format("0.57"), "dpctl", "packaging"],
+    install_requires=["numba >={0:s}".format("0.58"), "dpctl", "packaging"],
     packages=find_packages("."),
     include_package_data=True,
     zip_safe=False,


Test patch:
diff --git a/numba_dpex/tests/_helper.py b/numba_dpex/tests/_helper.py
index e5382eac69..3eca02ae06 100644
--- a/numba_dpex/tests/_helper.py
+++ b/numba_dpex/tests/_helper.py
@@ -5,6 +5,7 @@
 # SPDX-License-Identifier: Apache-2.0
 
 import contextlib
+import inspect
 import shutil
 from functools import cache
 
@@ -12,8 +13,7 @@
 import dpnp
 import pytest
 
-from numba_dpex import dpjit, numba_sem_version
-from numba_dpex.core import config
+from numba_dpex import config, dpjit, numba_sem_version
 
 
 @cache
@@ -179,6 +179,14 @@ def get_complex_dtypes(device=None):
     return dtypes
 
 
+def get_int_dtypes(device=None):
+    """
+    Build a list of integer types supported by DPNP based on device capabilities.
+    """
+
+    return [dpnp.int32, dpnp.int64]
+
+
 def get_float_dtypes(no_float16=True, device=None):
     """
     Build a list of floating types supported by DPNP based on device capabilities.
@@ -227,7 +235,7 @@ def get_all_dtypes(
 
     # add integer types
     if not no_int:
-        dtypes.extend([dpnp.int32, dpnp.int64])
+        dtypes.extend(get_int_dtypes(device=dev))
 
     # add floating types
     if not no_float:
@@ -276,3 +284,20 @@ def skip_if_dtype_not_supported(dt, q_or_dev):
         pytest.skip(
             f"{dev.name} does not support half precision floating point type"
         )
+
+
+def num_required_arguments(func):
+    """Returns number of required arguments of the functions. Does not work
+    with kwargs arguments."""
+    if func == dpnp.true_divide:
+        func = dpnp.divide
+
+    sig = inspect.signature(func)
+    params = sig.parameters
+    required_args = [
+        p
+        for p in params
+        if params[p].default == inspect._empty and p != "kwargs"
+    ]
+
+    return len(required_args)
diff --git a/numba_dpex/tests/core/types/IntEnumLiteral/__init__.py b/numba_dpex/tests/core/types/IntEnumLiteral/__init__.py
new file mode 100644
index 0000000000..ee20629a35
--- /dev/null
+++ b/numba_dpex/tests/core/types/IntEnumLiteral/__init__.py
@@ -0,0 +1,3 @@
+# SPDX-FileCopyrightText: 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
diff --git a/numba_dpex/tests/experimental/IntEnumLiteral/test_compilation.py b/numba_dpex/tests/core/types/IntEnumLiteral/test_compilation.py
similarity index 93%
rename from numba_dpex/tests/experimental/IntEnumLiteral/test_compilation.py
rename to numba_dpex/tests/core/types/IntEnumLiteral/test_compilation.py
index 31a0f6be56..9362ba20ff 100644
--- a/numba_dpex/tests/experimental/IntEnumLiteral/test_compilation.py
+++ b/numba_dpex/tests/core/types/IntEnumLiteral/test_compilation.py
@@ -6,7 +6,7 @@
 
 import numba_dpex.experimental as exp_dpex
 from numba_dpex import Range
-from numba_dpex.experimental.flag_enum import FlagEnum
+from numba_dpex.kernel_api.flag_enum import FlagEnum
 
 
 class MockFlags(FlagEnum):
diff --git a/numba_dpex/tests/experimental/IntEnumLiteral/test_type_creation.py b/numba_dpex/tests/core/types/IntEnumLiteral/test_type_creation.py
similarity index 86%
rename from numba_dpex/tests/experimental/IntEnumLiteral/test_type_creation.py
rename to numba_dpex/tests/core/types/IntEnumLiteral/test_type_creation.py
index 3e974b973d..3027dda904 100644
--- a/numba_dpex/tests/experimental/IntEnumLiteral/test_type_creation.py
+++ b/numba_dpex/tests/core/types/IntEnumLiteral/test_type_creation.py
@@ -7,8 +7,8 @@
 import pytest
 
 from numba_dpex.core.exceptions import IllegalIntEnumLiteralValueError
-from numba_dpex.experimental import IntEnumLiteral
-from numba_dpex.experimental.flag_enum import FlagEnum
+from numba_dpex.core.types import IntEnumLiteral
+from numba_dpex.kernel_api.flag_enum import FlagEnum
 
 
 def test_intenumliteral_creation():
diff --git a/numba_dpex/tests/experimental/IntEnumLiteral/test_type_registration.py b/numba_dpex/tests/core/types/IntEnumLiteral/test_type_registration.py
similarity index 77%
rename from numba_dpex/tests/experimental/IntEnumLiteral/test_type_registration.py
rename to numba_dpex/tests/core/types/IntEnumLiteral/test_type_registration.py
index 3bcbef8691..62dd1e4db9 100644
--- a/numba_dpex/tests/experimental/IntEnumLiteral/test_type_registration.py
+++ b/numba_dpex/tests/core/types/IntEnumLiteral/test_type_registration.py
@@ -6,9 +6,8 @@
 from numba.core.datamodel import default_manager
 
 from numba_dpex.core.datamodel.models import dpex_data_model_manager
-from numba_dpex.experimental import IntEnumLiteral
-from numba_dpex.experimental.flag_enum import FlagEnum
-from numba_dpex.experimental.models import exp_dmm
+from numba_dpex.core.types import IntEnumLiteral
+from numba_dpex.kernel_api.flag_enum import FlagEnum
 
 
 def test_data_model_registration():
@@ -24,11 +23,8 @@ class DummyFlags(FlagEnum):
     with pytest.raises(KeyError):
         default_manager.lookup(dummy)
 
-    with pytest.raises(KeyError):
-        dpex_data_model_manager.lookup(dummy)
-
     try:
-        exp_dmm.lookup(dummy)
+        dpex_data_model_manager.lookup(dummy)
     except:
         pytest.fail(
             "IntEnumLiteral type lookup failed in experimental "
diff --git a/numba_dpex/tests/core/types/range_types/test_data_model.py b/numba_dpex/tests/core/types/range_types/test_data_model.py
index 9dd4aad5c7..c00e234bb0 100644
--- a/numba_dpex/tests/core/types/range_types/test_data_model.py
+++ b/numba_dpex/tests/core/types/range_types/test_data_model.py
@@ -12,7 +12,7 @@
     dpex_data_model_manager,
 )
 from numba_dpex.core.descriptor import dpex_kernel_target
-from numba_dpex.core.types.range_types import NdRangeType, RangeType
+from numba_dpex.core.types.kernel_api.ranges import NdRangeType, RangeType
 
 rfields = ["ndim", "dim0", "dim1", "dim2"]
 ndrfields = ["ndim", "gdim0", "gdim1", "gdim2", "ldim0", "ldim1", "ldim2"]
diff --git a/numba_dpex/tests/core/types/test_box_unbox.py b/numba_dpex/tests/core/types/test_box_unbox.py
index a792414f94..42684f62d5 100644
--- a/numba_dpex/tests/core/types/test_box_unbox.py
+++ b/numba_dpex/tests/core/types/test_box_unbox.py
@@ -15,7 +15,7 @@
 import numpy
 import pytest
 
-import numba_dpex as dpx
+import numba_dpex as dpex
 
 constructors_with_parent = [
     dpctl.SyclEvent,
@@ -30,9 +30,10 @@
 ranges = [(10,), (10, 10), (10, 10, 10)]
 
 constructors_without_parent = [
-    functools.partial(dpx.Range, *r) for r in ranges
+    functools.partial(dpex.Range, *r) for r in ranges
 ] + [
-    functools.partial(dpx.NdRange, dpx.Range(*r), dpx.Range(*r)) for r in ranges
+    functools.partial(dpex.NdRange, dpex.Range(*r), dpex.Range(*r))
+    for r in ranges
 ]
 
 
@@ -51,12 +52,12 @@ def constructor_np(request):
     return request.param
 
 
-@dpx.dpjit
+@dpex.dpjit
 def unbox_box(a):
     return a
 
 
-@dpx.dpjit
+@dpex.dpjit
 def unbox(a):
     return None
 
diff --git a/numba_dpex/tests/dpjit_tests/parfors/test_builtin_ops.py b/numba_dpex/tests/dpjit_tests/parfors/test_builtin_ops.py
index 8777d0e241..2f87c5f117 100644
--- a/numba_dpex/tests/dpjit_tests/parfors/test_builtin_ops.py
+++ b/numba_dpex/tests/dpjit_tests/parfors/test_builtin_ops.py
@@ -2,68 +2,350 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
+from itertools import product
+
 import dpctl
 import dpnp
 import numpy
 import pytest
 
 from numba_dpex import dpjit
-from numba_dpex.tests._helper import get_all_dtypes
+from numba_dpex.tests._helper import (
+    get_float_dtypes,
+    get_int_dtypes,
+    has_opencl_gpu,
+    is_windows,
+    num_required_arguments,
+)
 
 
 def parfor_add(a, b):
+    "Same as a + b."
     return a + b
 
 
-def parfor_sub(a, b):
-    return a - b
+def parfor_and_(a, b):
+    "Same as a & b."
+    return a & b
 
 
-def parfor_mult(a, b):
-    return a * b
+def parfor_floordiv(a, b):
+    "Same as a // b."
+    return a // b
 
 
-def parfor_divide(a, b):
-    return a / b
+def parfor_inv(a):
+    "Same as ~a."
+    return ~a
+
+
+def parfor_lshift(a, b):
+    "Same as a << b."
+    return a << b
 
 
-def parfor_modulus(a, b):
+def parfor_mod(a, b):
+    "Same as a % b."
     return a % b
 
 
-def parfor_exponent(a, b):
+def parfor_mul(a, b):
+    "Same as a * b."
+    return a * b
+
+
+def parfor_matmul(a, b):
+    "Same as a @ b."
+    return a @ b
+
+
+def parfor_neg(a):
+    "Same as -a."
+    return -a
+
+
+def parfor_or_(a, b):
+    "Same as a | b."
+    return a | b
+
+
+def parfor_pos(a):
+    "Same as +a."
+    return +a
+
+
+def parfor_pow(a, b):
+    "Same as a ** b."
     return a**b
 
 
+def parfor_rshift(a, b):
+    "Same as a >> b."
+    return a >> b
+
+
+def parfor_sub(a, b):
+    "Same as a - b."
+    return a - b
+
+
+def parfor_truediv(a, b):
+    "Same as a / b."
+    return a / b
+
+
+def parfor_xor(a, b):
+    "Same as a ^ b."
+    return a ^ b
+
+
+# Inplace operations
+def parfor_iadd(a, b):
+    "Same as a += b."
+    a += b
+    return a
+
+
+def parfor_iand(a, b):
+    "Same as a &= b."
+    a &= b
+    return a
+
+
+def parfor_ifloordiv(a, b):
+    "Same as a //= b."
+    a //= b
+    return a
+
+
+def parfor_ilshift(a, b):
+    "Same as a <<= b."
+    a <<= b
+    return a
+
+
+def parfor_imod(a, b):
+    "Same as a %= b."
+    a %= b
+    return a
+
+
+def parfor_imul(a, b):
+    "Same as a *= b."
+    a *= b
+    return a
+
+
+def parfor_imatmul(a, b):
+    "Same as a @= b."
+    a @= b
+    return a
+
+
+def parfor_ior(a, b):
+    "Same as a |= b."
+    a |= b
+    return a
+
+
+def parfor_ipow(a, b):
+    "Same as a **= b."
+    a **= b
+    return a
+
+
+def parfor_irshift(a, b):
+    "Same as a >>= b."
+    a >>= b
+    return a
+
+
+def parfor_isub(a, b):
+    "Same as a -= b."
+    a -= b
+    return a
+
+
+def parfor_itruediv(a, b):
+    "Same as a /= b."
+    a /= b
+    return a
+
+
+def parfor_ixor(a, b):
+    "Same as a ^= b."
+    a ^= b
+    return a
+
+
+# Comparison operations
+
+
+def parfor_lt(a, b):
+    "Same as a < b."
+    return a < b
+
+
+def parfor_le(a, b):
+    "Same as a <= b."
+    return a <= b
+
+
+def parfor_eq(a, b):
+    "Same as a == b."
+    return a == b
+
+
+def parfor_ne(a, b):
+    "Same as a != b."
+    return a != b
+
+
+def parfor_ge(a, b):
+    "Same as a >= b."
+    return a >= b
+
+
+def parfor_gt(a, b):
+    "Same as a > b."
+    return a > b
+
+
+# logical
+
+
+def parfor_not_(a):
+    "Same as not a."
+    return not a
+
+
 shapes = [100, (25, 4)]
-dtypes = get_all_dtypes(
-    no_bool=True, no_float16=True, no_none=True, no_complex=True
-)
+int_dtypes = set(get_int_dtypes())
+float_dtypes = set(get_float_dtypes())
 usm_types = ["device", "host", "shared"]
-funcs = [
+
+funcs = {
     parfor_add,
+    parfor_and_,
+    parfor_floordiv,
+    parfor_inv,
+    parfor_lshift,
+    parfor_mod,
+    parfor_mul,
+    parfor_matmul,
+    parfor_neg,
+    parfor_or_,
+    parfor_pos,
+    parfor_pow,
+    parfor_rshift,
     parfor_sub,
-    parfor_mult,
-    parfor_divide,
-    parfor_modulus,
-    parfor_exponent,
-]
+    parfor_truediv,
+    parfor_xor,
+}
 
-# TODO: fails for integer because it is being cast to float64 internally?
-if dpnp.float64 not in dtypes:
-    funcs.remove(parfor_divide)
-    funcs.remove(parfor_exponent)
+inplace_funcs = {
+    parfor_iand,
+    parfor_iadd,
+    parfor_iand,
+    parfor_ifloordiv,
+    parfor_ilshift,
+    parfor_imod,
+    parfor_imul,
+    parfor_imatmul,
+    parfor_ior,
+    parfor_ipow,
+    parfor_irshift,
+    parfor_isub,
+    parfor_itruediv,
+    parfor_ixor,
+}
 
+comparison_funcs = {
+    parfor_lt,
+    parfor_le,
+    parfor_eq,
+    parfor_ne,
+    parfor_ge,
+    parfor_gt,
+}
 
-def parfor_floor(a, b):
-    return a // b
+all_funcs = funcs | inplace_funcs | comparison_funcs
+
+not_float_funcs = {
+    parfor_inv,
+    parfor_and_,
+    parfor_floordiv,
+    parfor_lshift,
+    parfor_mod,
+    parfor_or_,
+    parfor_rshift,
+    parfor_xor,
+    parfor_iand,
+    parfor_ifloordiv,
+    parfor_ilshift,
+    parfor_imod,
+    parfor_ior,
+    parfor_irshift,
+    parfor_ixor,
+}
+
+not_int_funcs = {
+    parfor_itruediv,
+}
+
+not_bool_funcs = {
+    parfor_neg,
+    parfor_pos,
+    parfor_sub,
+    parfor_lshift,
+    parfor_rshift,
+    parfor_mod,
+    parfor_pow,
+    parfor_floordiv,
+    parfor_isub,
+    parfor_ilshift,
+    parfor_irshift,
+    parfor_imod,
+    parfor_ipow,
+    parfor_ifloordiv,
+} | not_int_funcs
+
+unsupported_funcs = {
+    parfor_matmul,
+    parfor_imatmul,
+    parfor_not_,  # not supported by numpy
+}
+
+
+if has_opencl_gpu() or is_windows():
+    unsupported_funcs |= {parfor_ipow}
+
+
+func_dtypes = (
+    set(product(all_funcs - not_float_funcs, float_dtypes))
+    | set(product(all_funcs - not_int_funcs, int_dtypes))
+    | set(product((all_funcs | {parfor_not_}) - not_bool_funcs, {dpnp.bool}))
+)
 
 
 @pytest.mark.parametrize("shape", shapes)
-@pytest.mark.parametrize("dtype", dtypes)
 @pytest.mark.parametrize("usm_type", usm_types)
-@pytest.mark.parametrize("func", funcs)
+@pytest.mark.parametrize(
+    "func, dtype",
+    sorted(list(func_dtypes), key=lambda a: (str(a[0]), str(a[1]))),
+)
 def test_built_in_operators(shape, dtype, usm_type, func):
+    if func in unsupported_funcs:
+        pytest.xfail(reason="not supported")
+
+    if dpnp.float64 not in float_dtypes and func in [
+        parfor_truediv,
+        parfor_pow,
+    ]:
+        pytest.xfail(
+            f"{func} does not work on without double precision support"
+        )
+
     queue = dpctl.SyclQueue()
     a = dpnp.zeros(
         shape=shape, dtype=dtype, usm_type=usm_type, sycl_queue=queue
@@ -71,7 +353,10 @@ def test_built_in_operators(shape, dtype, usm_type, func):
     b = dpnp.ones(shape=shape, dtype=dtype, usm_type=usm_type, sycl_queue=queue)
     try:
         op = dpjit(func)
-        c = op(a, b)
+        if num_required_arguments(func) == 1:
+            c = op(a)
+        else:
+            c = op(a, b)
         del op
     except Exception:
         pytest.fail("Failed to compile.")
@@ -81,13 +366,16 @@ def test_built_in_operators(shape, dtype, usm_type, func):
     else:
         assert c.shape == shape
 
-    if func != parfor_divide:
+    if func not in {parfor_truediv} | comparison_funcs:
         assert c.dtype == dtype
     assert c.usm_type == usm_type
 
     assert c.sycl_device.filter_string == queue.sycl_device.filter_string
 
-    expected = dpnp.asnumpy(func(a, b))
+    if num_required_arguments(func) == 1:
+        expected = dpnp.asnumpy(func(a))
+    else:
+        expected = dpnp.asnumpy(func(a, b))
     nc = dpnp.asnumpy(c)
 
     numpy.allclose(nc, expected)
diff --git a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_bitwise_ops.py b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_bitwise_ops.py
deleted file mode 100644
index ab10c83e2f..0000000000
--- a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_bitwise_ops.py
+++ /dev/null
@@ -1,92 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl.tensor as dpt
-import dpnp
-import numpy as np
-import pytest
-
-from numba_dpex import dpjit
-
-list_of_binary_ops = [
-    "bitwise_and",
-    "bitwise_or",
-    "bitwise_xor",
-    "left_shift",
-    "right_shift",
-]
-
-
-@pytest.fixture(params=list_of_binary_ops)
-def binary_op(request):
-    return request.param
-
-
-list_of_unary_ops = [
-    "bitwise_not",
-    "invert",
-]
-
-
-@pytest.fixture(params=list_of_unary_ops)
-def unary_op(request):
-    return request.param
-
-
-list_of_dtypes = [
-    dpnp.int32,
-    dpnp.int64,
-]
-
-
-@pytest.fixture(params=list_of_dtypes)
-def input_arrays(request):
-    # The size of input and out arrays to be used
-    N = 2048
-    a = dpnp.array(dpnp.random.random(N), request.param)
-    b = dpnp.array(dpnp.random.random(N), request.param)
-    return a, b
-
-
-def test_binary_ops(binary_op, input_arrays):
-    a, b = input_arrays
-    binop = getattr(dpnp, binary_op)
-    actual = dpnp.empty(shape=a.shape, dtype=a.dtype)
-    expected = dpnp.empty(shape=a.shape, dtype=a.dtype)
-
-    @dpjit
-    def f(a, b):
-        return binop(a, b)
-
-    actual = f(a, b)
-
-    expected = binop(a, b)
-    np.testing.assert_allclose(
-        dpt.asnumpy(actual._array_obj),
-        dpt.asnumpy(expected._array_obj),
-        rtol=1e-5,
-        atol=0,
-    )
-
-
-def test_unary_ops(unary_op, input_arrays):
-    a = input_arrays[0]
-    uop = getattr(dpnp, unary_op)
-    actual = np.empty(shape=a.shape, dtype=a.dtype)
-    expected = np.empty(shape=a.shape, dtype=a.dtype)
-
-    @dpjit
-    def f(a):
-        return uop(a)
-
-    actual = f(a)
-
-    expected = uop(a)
-
-    np.testing.assert_allclose(
-        dpt.asnumpy(actual._array_obj),
-        dpt.asnumpy(expected._array_obj),
-        rtol=1e-5,
-        atol=0,
-    )
diff --git a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_logic_ops.py b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_logic_ops.py
deleted file mode 100644
index 9263cc83b4..0000000000
--- a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_logic_ops.py
+++ /dev/null
@@ -1,99 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import dpctl.tensor as dpt
-import dpnp
-import numpy as np
-import pytest
-
-from numba_dpex import dpjit
-from numba_dpex.tests._helper import get_all_dtypes
-
-""" Following cases, dpnp raises NotImplementedError"""
-
-list_of_binary_ops = [
-    "greater",
-    "greater_equal",
-    "less",
-    "less_equal",
-    "not_equal",
-    "equal",
-    "logical_and",
-    "logical_or",
-    "logical_xor",
-]
-
-
-@pytest.fixture(params=list_of_binary_ops)
-def binary_op(request):
-    return request.param
-
-
-list_of_unary_ops = [
-    "isinf",
-    "isfinite",
-    "isnan",
-]
-
-
-@pytest.fixture(params=list_of_unary_ops)
-def unary_op(request):
-    return request.param
-
-
-list_of_dtypes = get_all_dtypes(
-    no_bool=True, no_float16=True, no_none=True, no_complex=True
-)
-
-
-@pytest.fixture(params=list_of_dtypes)
-def input_arrays(request):
-    # The size of input and out arrays to be used
-    N = 2048
-    a = dpnp.array(dpnp.random.random(N), request.param)
-    b = dpnp.array(dpnp.random.random(N), request.param)
-    return a, b
-
-
-def test_binary_ops(binary_op, input_arrays):
-    a, b = input_arrays
-    binop = getattr(dpnp, binary_op)
-    actual = dpnp.empty(shape=a.shape, dtype=a.dtype)
-    expected = dpnp.empty(shape=a.shape, dtype=a.dtype)
-
-    @dpjit
-    def f(a, b):
-        return binop(a, b)
-
-    actual = f(a, b)
-
-    expected = binop(a, b)
-    np.testing.assert_allclose(
-        dpt.asnumpy(actual._array_obj),
-        dpt.asnumpy(expected._array_obj),
-        rtol=1e-5,
-        atol=0,
-    )
-
-
-def test_unary_ops(unary_op, input_arrays):
-    a = input_arrays[0]
-    uop = getattr(dpnp, unary_op)
-    actual = dpnp.empty(shape=a.shape, dtype=a.dtype)
-    expected = dpnp.empty(shape=a.shape, dtype=a.dtype)
-
-    @dpjit
-    def f(a):
-        return uop(a)
-
-    actual = f(a)
-
-    expected = uop(a)
-    np.testing.assert_allclose(
-        dpt.asnumpy(actual._array_obj),
-        dpt.asnumpy(expected._array_obj),
-        rtol=1e-5,
-        atol=0,
-    )
diff --git a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_transcedental_functions.py b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_transcedental_functions.py
deleted file mode 100644
index 71b6e02d30..0000000000
--- a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_transcedental_functions.py
+++ /dev/null
@@ -1,126 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import dpctl.tensor as dpt
-import dpnp
-import numpy as np
-import pytest
-
-from numba_dpex import dpjit
-from numba_dpex.tests._helper import get_all_dtypes, is_gen12
-
-"""dpnp raise error on : mod, abs and remainder(float32)"""
-list_of_binary_ops = [
-    "add",
-    "subtract",
-    "multiply",
-    "divide",
-    "power",
-    #   "remainder",
-    #   "mod",
-    "hypot",
-    "maximum",
-    "minimum",
-    "fmax",
-    "fmin",
-]
-
-
-@pytest.fixture(params=list_of_binary_ops)
-def binary_op(request):
-    return request.param
-
-
-list_of_unary_ops = [
-    "negative",
-    #   "abs",
-    "absolute",
-    "fabs",
-    "sign",
-    "conj",
-    "exp",
-    "exp2",
-    "log",
-    "log2",
-    "log10",
-    "expm1",
-    "log1p",
-    "sqrt",
-    "square",
-    "reciprocal",
-    "conjugate",
-    "floor",
-    "ceil",
-    "trunc",
-    "erf",
-]
-
-
-@pytest.fixture(params=list_of_unary_ops)
-def unary_op(request):
-    return request.param
-
-
-list_of_dtypes = get_all_dtypes(
-    no_bool=True, no_int=True, no_float16=True, no_none=True, no_complex=True
-)
-
-
-@pytest.fixture(params=list_of_dtypes)
-def input_arrays(request):
-    # The size of input and out arrays to be used
-    N = 2048
-    a = dpnp.array(dpnp.random.random(N), request.param)
-    b = dpnp.array(dpnp.random.random(N), request.param)
-    return a, b
-
-
-def test_binary_ops(binary_op, input_arrays):
-    a, b = input_arrays
-    binop = getattr(dpnp, binary_op)
-    actual = dpnp.empty(shape=a.shape, dtype=a.dtype)
-    expected = dpnp.empty(shape=a.shape, dtype=a.dtype)
-
-    @dpjit
-    def f(a, b):
-        return binop(a, b)
-
-    actual = f(a, b)
-
-    expected = binop(a, b)
-
-    np.testing.assert_allclose(
-        dpt.asnumpy(actual._array_obj),
-        dpt.asnumpy(expected._array_obj),
-        rtol=1e-5,
-        atol=0,
-    )
-
-
-def test_unary_ops(unary_op, input_arrays):
-    skip_ops = ["abs", "sign", "log", "log2", "log10", "expm1"]
-    device = dpctl.SyclDevice()
-    if unary_op in skip_ops and is_gen12(device.filter_string):
-        pytest.skip()
-
-    a = input_arrays[0]
-    uop = getattr(dpnp, unary_op)
-    actual = dpnp.empty(shape=a.shape, dtype=a.dtype)
-    expected = dpnp.empty(shape=a.shape, dtype=a.dtype)
-
-    @dpjit
-    def f(a):
-        return uop(a)
-
-    actual = f(a)
-
-    expected = uop(a)
-
-    np.testing.assert_allclose(
-        dpt.asnumpy(actual._array_obj),
-        dpt.asnumpy(expected._array_obj),
-        rtol=1e-5,
-        atol=0,
-    )
diff --git a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_trigonometric_functions.py b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_trigonometric_functions.py
deleted file mode 100644
index f5fcc1aa8d..0000000000
--- a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_trigonometric_functions.py
+++ /dev/null
@@ -1,91 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import dpctl.tensor as dpt
-import dpnp
-import numpy as np
-import pytest
-
-from numba_dpex import dpjit
-from numba_dpex.tests._helper import get_all_dtypes, is_gen12
-
-list_of_trig_ops = [
-    "sin",
-    "cos",
-    "tan",
-    "arcsin",
-    "arccos",
-    "arctan",
-    "arctan2",
-    "sinh",
-    "cosh",
-    "tanh",
-    "arcsinh",
-    "arccosh",
-    "arctanh",
-    "deg2rad",
-    "rad2deg",
-    "degrees",
-    "radians",
-]
-
-list_of_dtypes = get_all_dtypes(
-    no_bool=True, no_int=True, no_float16=True, no_none=True, no_complex=True
-)
-
-# TODO: fails for float32 because it uses cast to float64 internally?
-if dpnp.float64 not in list_of_dtypes:
-    list_of_trig_ops.remove("arctan2")
-
-
-@pytest.fixture(params=list_of_trig_ops)
-def trig_op(request):
-    return request.param
-
-
-@pytest.fixture(params=list_of_dtypes)
-def input_arrays(request):
-    # The size of input and out arrays to be used
-    N = 2048
-
-    a = dpnp.array(dpnp.random.random(N), request.param)
-    b = dpnp.array(dpnp.random.random(N), request.param)
-    return a, b
-
-
-def test_trigonometric_fn(trig_op, input_arrays):
-    filter_str = dpctl.SyclDevice().filter_string
-    # FIXME: Why does archcosh fail on Gen12 discrete graphics card?
-    if trig_op == "arccosh" and is_gen12(filter_str):
-        pytest.skip()
-
-    a, b = input_arrays
-    trig_fn = getattr(dpnp, trig_op)
-    actual = dpnp.empty(shape=a.shape, dtype=a.dtype)
-    expected = dpnp.empty(shape=a.shape, dtype=a.dtype)
-
-    if trig_op == "arctan2":
-
-        @dpjit
-        def f(a, b):
-            return trig_fn(a, b)
-
-        actual = f(a, b)
-        expected = trig_fn(a, b)
-    else:
-
-        @dpjit
-        def f(a):
-            return trig_fn(a)
-
-        actual = f(a)
-        expected = trig_fn(a)
-
-    np.testing.assert_allclose(
-        dpt.asnumpy(actual._array_obj),
-        dpt.asnumpy(expected._array_obj),
-        rtol=1e-5,
-        atol=0,
-    )
diff --git a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_ufuncs.py b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_ufuncs.py
new file mode 100644
index 0000000000..fe2e0afc91
--- /dev/null
+++ b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_ufuncs.py
@@ -0,0 +1,179 @@
+# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+from itertools import product
+
+import dpctl
+import dpctl.tensor as dpt
+import dpnp
+import numpy as np
+import pytest
+
+from numba_dpex import dpjit
+from numba_dpex.core.typing.dpnpdecl import (
+    _bit_twiddling_functions,
+    _comparison_functions,
+    _floating_functions,
+    _logic_functions,
+    _math_operations,
+    _trigonometric_functions,
+    _unsupported,
+)
+from numba_dpex.tests._helper import (
+    get_float_dtypes,
+    get_int_dtypes,
+    has_opencl_gpu,
+    is_gen12,
+    num_required_arguments,
+)
+
+int_ops = {
+    "gcd",
+    "lcm",
+    "mod",
+    "divmod",
+} | set(_bit_twiddling_functions)
+
+
+float_ops = {
+    "fmod",
+    "reciprocal",
+}
+
+all_ufuncs = set(
+    _math_operations
+    + _trigonometric_functions
+    + _bit_twiddling_functions
+    + _comparison_functions
+    + _floating_functions
+    + _logic_functions
+) - {
+    "frexp"  # there is no dpnp.frexp
+}
+
+unary_ops = {
+    op for op in all_ufuncs if num_required_arguments(getattr(dpnp, op)) == 1
+} | {"frexp"}
+
+binary_ops = {
+    op for op in all_ufuncs if num_required_arguments(getattr(dpnp, op)) == 2
+}
+
+float_ops |= {
+    "copysign",
+    "divide",
+    "hypot",
+    "power",
+    "true_divide",
+    "arccos",
+    "arccosh",
+    "arcsin",
+    "arcsinh",
+    "arctan",
+    "arctanh",
+    "ceil",
+    "cos",
+    "cosh",
+    "deg2rad",
+    "degrees",
+    "erf",
+    "exp",
+    "exp2",
+    "expm1",
+    "fabs",
+    "floor",
+    "log",
+    "log10",
+    "log1p",
+    "log2",
+    "rad2deg",
+    "radians",
+    "sin",
+    "sinh",
+    "sqrt",
+    "tan",
+    "tanh",
+    "trunc",
+}
+
+int_dtypes = get_int_dtypes()
+float_dtypes = get_float_dtypes(no_float16=True)
+
+unary_op_dtypes = set(product(unary_ops - int_ops, float_dtypes)) | set(
+    product(unary_ops - float_ops, int_dtypes)
+)
+binary_op_dtypes = set(product(binary_ops - int_ops, float_dtypes)) | set(
+    product(binary_ops - float_ops, int_dtypes)
+)
+
+N = 1024
+
+
+@pytest.mark.parametrize(
+    "binary_op, dtype",
+    sorted(list(binary_op_dtypes), key=lambda a: (a[0], str(a[1]))),
+)
+def test_binary_ops(binary_op, dtype):
+    # TODO: fails for float32 because it uses cast to float64 internally?
+    if binary_op == "arctan2" and dpnp.float64 not in float_dtypes:
+        pytest.xfail("arctan2 requires float64 support")
+
+    a = dpnp.array(dpnp.random.random(N), dtype)
+    b = dpnp.array(dpnp.random.random(N), dtype)
+
+    binop = getattr(dpnp, binary_op)
+    actual = dpnp.empty(shape=a.shape, dtype=a.dtype)
+    expected = dpnp.empty(shape=a.shape, dtype=a.dtype)
+
+    if binary_op in _unsupported:
+        pytest.xfail(reason="not supported")
+
+    @dpjit
+    def f(a, b):
+        return binop(a, b)
+
+    actual = f(a, b)
+
+    expected = binop(a, b)
+
+    np.testing.assert_allclose(
+        dpt.asnumpy(actual._array_obj),
+        dpt.asnumpy(expected._array_obj),
+        rtol=1e-5,
+        atol=0,
+    )
+
+
+@pytest.mark.parametrize(
+    "unary_op, dtype",
+    sorted(list(unary_op_dtypes), key=lambda a: (a[0], str(a[1]))),
+)
+def test_unary_ops(unary_op, dtype):
+    if unary_op in _unsupported:
+        pytest.xfail(reason="not supported")
+
+    xfail_ops = ["sign", "log", "log2", "log10", "expm1", "arccosh"]
+    if unary_op in xfail_ops and is_gen12(dpctl.SyclDevice().filter_string):
+        pytest.xfail(f"{unary_op} does not work on gen12")
+
+    a = dpnp.array(dpnp.random.random(N), dtype)
+
+    uop = getattr(dpnp, unary_op)
+    actual = dpnp.empty(shape=a.shape, dtype=a.dtype)
+    expected = dpnp.empty(shape=a.shape, dtype=a.dtype)
+
+    @dpjit
+    def f(a):
+        return uop(a)
+
+    actual = f(a)
+
+    expected = uop(a)
+
+    np.testing.assert_allclose(
+        dpt.asnumpy(actual._array_obj),
+        dpt.asnumpy(expected._array_obj),
+        rtol=1e-5,
+        atol=0,
+    )
diff --git a/numba_dpex/tests/experimental/codegen/test_intenum_literal_codegen.py b/numba_dpex/tests/experimental/codegen/test_intenum_literal_codegen.py
index 77233e7bd9..8512bd8016 100644
--- a/numba_dpex/tests/experimental/codegen/test_intenum_literal_codegen.py
+++ b/numba_dpex/tests/experimental/codegen/test_intenum_literal_codegen.py
@@ -9,7 +9,7 @@
 
 import numba_dpex.experimental as exp_dpex
 from numba_dpex import DpctlSyclQueue, DpnpNdArray, int64
-from numba_dpex.experimental.flag_enum import FlagEnum
+from numba_dpex.kernel_api.flag_enum import FlagEnum
 
 
 def test_compilation_as_literal_constant():
diff --git a/numba_dpex/tests/experimental/kernel_iface/__init__.py b/numba_dpex/tests/experimental/kernel_api_overloads/__init__.py
similarity index 100%
rename from numba_dpex/tests/experimental/kernel_iface/__init__.py
rename to numba_dpex/tests/experimental/kernel_api_overloads/__init__.py
diff --git a/numba_dpex/tests/experimental/kernel_iface/spv_overloads/__init__.py b/numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/__init__.py
similarity index 100%
rename from numba_dpex/tests/experimental/kernel_iface/spv_overloads/__init__.py
rename to numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/__init__.py
diff --git a/numba_dpex/tests/experimental/kernel_iface/spv_overloads/test_atomic_fetch_phi.py b/numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_atomic_fetch_phi.py
similarity index 81%
rename from numba_dpex/tests/experimental/kernel_iface/spv_overloads/test_atomic_fetch_phi.py
rename to numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_atomic_fetch_phi.py
index 4dd02066a5..30f268c87a 100644
--- a/numba_dpex/tests/experimental/kernel_iface/spv_overloads/test_atomic_fetch_phi.py
+++ b/numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_atomic_fetch_phi.py
@@ -8,7 +8,7 @@
 
 import numba_dpex as dpex
 import numba_dpex.experimental as dpex_exp
-from numba_dpex.experimental.kernel_iface import AtomicRef
+from numba_dpex.kernel_api import AtomicRef
 from numba_dpex.tests._helper import get_all_dtypes
 
 list_of_supported_dtypes = get_all_dtypes(
@@ -72,6 +72,36 @@ def _kernel(a, b, ref_index):
         assert b[ref_index] == b[ref_index + 1]
 
 
+def test_fetch_phi_retval(fetch_phi_fn):
+    """A test for all fetch_phi atomic functions."""
+
+    @dpex_exp.kernel
+    def _kernel(a, b, c):
+        i = dpex.get_global_id(0)
+        v = AtomicRef(b, index=i)
+        c[i] = getattr(v, fetch_phi_fn)(a[i])
+
+    N = 10
+    a = dpnp.arange(N, dtype=dpnp.int32)
+    b = dpnp.ones(N, dtype=dpnp.int32)
+    c = dpnp.zeros(N, dtype=dpnp.int32)
+    a_copy = dpnp.copy(a)
+    b_copy = dpnp.copy(b)
+    c_copy = dpnp.copy(c)
+
+    dpex_exp.call_kernel(_kernel, dpex.Range(10), a, b, c)
+
+    # Verify if the value returned by fetch_phi kernel
+    # stored into `c` is same as the value returned
+    # by fetch_phi python stored into `c_copy`
+    for i in range(a.size):
+        v = AtomicRef(b_copy, index=i)
+        c_copy[i] = getattr(v, fetch_phi_fn)(a_copy[i])
+
+    for i in range(a.size):
+        assert c[i] == c_copy[i]
+
+
 def test_fetch_phi_diff_types(fetch_phi_fn):
     """A negative test that verifies that a TypingError is raised if
     AtomicRef type and value to be added are of different types.
diff --git a/numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_atomic_load_store_cmp_exchg.py b/numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_atomic_load_store_cmp_exchg.py
new file mode 100644
index 0000000000..9d34d99319
--- /dev/null
+++ b/numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_atomic_load_store_cmp_exchg.py
@@ -0,0 +1,123 @@
+# SPDX-FileCopyrightText: 2023 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import dpctl
+import dpnp
+import pytest
+from numba.core.errors import TypingError
+
+import numba_dpex as dpex
+import numba_dpex.experimental as dpex_exp
+from numba_dpex.kernel_api import AtomicRef
+from numba_dpex.tests._helper import get_all_dtypes
+
+list_of_supported_dtypes = get_all_dtypes(
+    no_bool=True, no_float16=True, no_none=True, no_complex=True
+)
+
+
+@pytest.fixture(params=["store", "exchange"])
+def store_exchange_fn(request):
+    return request.param
+
+
+@pytest.mark.parametrize("supported_dtype", list_of_supported_dtypes)
+def test_load_store_fn(supported_dtype):
+    """A test for load/store atomic functions."""
+
+    @dpex_exp.kernel
+    def _kernel(item, a, b):
+        i = item.get_id(0)
+        a_ref = AtomicRef(a, index=i)
+        b_ref = AtomicRef(b, index=i)
+        val = b_ref.load()
+        a_ref.store(val)
+
+    N = 10
+    a = dpnp.zeros(2 * N, dtype=supported_dtype)
+    b = dpnp.arange(N, dtype=supported_dtype)
+
+    dpex_exp.call_kernel(_kernel, dpex.Range(b.size), a, b)
+    # Verify that `b[i]` loaded and stored into a[i] by kernel
+    # matches the `b[i]` loaded stored into a[i] using Python
+    for i in range(b.size):
+        a_ref = AtomicRef(a, index=i + b.size)
+        b_ref = AtomicRef(b, index=i)
+        a_ref.store(b_ref.load())
+
+    for i in range(b.size):
+        assert a[i] == a[i + b.size]
+
+
+@pytest.mark.parametrize("supported_dtype", list_of_supported_dtypes)
+def test_exchange_fn(supported_dtype):
+    """A test for exchange atomic function."""
+
+    @dpex_exp.kernel
+    def _kernel(item, a, b):
+        i = item.get_id(0)
+        v = AtomicRef(a, index=i)
+        b[i] = v.exchange(b[i])
+
+    N = 10
+    a_orig = dpnp.zeros(2 * N, dtype=supported_dtype)
+    b_orig = dpnp.arange(N, dtype=supported_dtype)
+
+    a_copy = dpnp.copy(a_orig)
+    b_copy = dpnp.copy(b_orig)
+
+    dpex_exp.call_kernel(_kernel, dpex.Range(b_orig.size), a_copy, b_copy)
+
+    # Values in `b` have been exchanged
+    # with values in `a`.
+    # Test if `a_copy` is same as `b_orig`
+    # and `b_copy` is same as `a_orig`
+    for i in range(b_orig.size):
+        assert a_copy[i] == b_orig[i]
+        assert b_copy[i] == a_orig[i]
+
+
+@pytest.mark.parametrize("supported_dtype", list_of_supported_dtypes)
+def test_compare_exchange_fns(supported_dtype):
+    """A test for compare exchange atomic functions."""
+
+    @dpex_exp.kernel
+    def _kernel(b):
+        b_ref = AtomicRef(b, index=1)
+        b[0] = b_ref.compare_exchange(
+            expected_ref=b, desired=b[3], expected_idx=2
+        )
+
+    b = dpnp.arange(4, dtype=supported_dtype)
+
+    dpex_exp.call_kernel(_kernel, dpex.Range(1), b)
+
+    # check for failure
+    assert b[0] == 0
+    assert b[2] == b[1]
+
+    dpex_exp.call_kernel(_kernel, dpex.Range(1), b)
+
+    # check for success
+    assert b[0] == 1
+    assert b[1] == b[3]
+
+
+def test_store_exchange_diff_types(store_exchange_fn):
+    """A negative test that verifies that a TypingError is raised if
+    AtomicRef type and value are of different types.
+    """
+
+    @dpex_exp.kernel
+    def _kernel(item, a, b):
+        i = item.get_id(0)
+        v = AtomicRef(b, index=0)
+        getattr(v, store_exchange_fn)(a[i])
+
+    N = 10
+    a = dpnp.ones(N, dtype=dpnp.float32)
+    b = dpnp.zeros(N, dtype=dpnp.int32)
+
+    with pytest.raises(TypingError):
+        dpex_exp.call_kernel(_kernel, dpex.Range(10), a, b)
diff --git a/numba_dpex/tests/experimental/kernel_iface/spv_overloads/test_atomic_ref.py b/numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_atomic_ref.py
similarity index 94%
rename from numba_dpex/tests/experimental/kernel_iface/spv_overloads/test_atomic_ref.py
rename to numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_atomic_ref.py
index 60f31de0ea..fb40f9311d 100644
--- a/numba_dpex/tests/experimental/kernel_iface/spv_overloads/test_atomic_ref.py
+++ b/numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_atomic_ref.py
@@ -8,7 +8,7 @@
 
 import numba_dpex as dpex
 import numba_dpex.experimental as dpex_exp
-from numba_dpex.experimental.kernel_iface import AddressSpace, AtomicRef
+from numba_dpex.kernel_api import AddressSpace, AtomicRef
 
 
 def test_atomic_ref_compilation():
diff --git a/numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_barriers.py b/numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_barriers.py
new file mode 100644
index 0000000000..61be6cb320
--- /dev/null
+++ b/numba_dpex/tests/experimental/kernel_api_overloads/spv_overloads/test_barriers.py
@@ -0,0 +1,30 @@
+import dpnp
+
+import numba_dpex as dpex
+import numba_dpex.experimental as dpex_exp
+from numba_dpex.kernel_api import MemoryScope, NdItem, group_barrier
+from numba_dpex.tests._helper import skip_windows
+
+
+# TODO: https://github.com/IntelPython/numba-dpex/issues/1308
+@skip_windows
+def test_group_barrier():
+    """A test for group_barrier function."""
+
+    @dpex_exp.kernel
+    def _kernel(nd_item: NdItem, a):
+        i = nd_item.get_global_id(0)
+
+        a[i] += 1
+        group_barrier(nd_item.get_group(), MemoryScope.DEVICE)
+
+        if i == 0:
+            for idx in range(1, a.size):
+                a[0] += a[idx]
+
+    N = 16
+    a = dpnp.ones(N, dtype=dpnp.int32)
+
+    dpex_exp.call_kernel(_kernel, dpex.NdRange((N,), (N,)), a)
+
+    assert a[0] == N * 2
diff --git a/numba_dpex/tests/experimental/kernel_iface/test_memory_enum_compilation.py b/numba_dpex/tests/experimental/kernel_api_overloads/test_memory_enum_compilation.py
similarity index 95%
rename from numba_dpex/tests/experimental/kernel_iface/test_memory_enum_compilation.py
rename to numba_dpex/tests/experimental/kernel_api_overloads/test_memory_enum_compilation.py
index a9dae36678..991e256784 100644
--- a/numba_dpex/tests/experimental/kernel_iface/test_memory_enum_compilation.py
+++ b/numba_dpex/tests/experimental/kernel_api_overloads/test_memory_enum_compilation.py
@@ -6,11 +6,7 @@
 
 import numba_dpex.experimental as exp_dpex
 from numba_dpex import Range
-from numba_dpex.experimental.kernel_iface import (
-    AddressSpace,
-    MemoryOrder,
-    MemoryScope,
-)
+from numba_dpex.kernel_api import AddressSpace, MemoryOrder, MemoryScope
 
 
 def test_compilation_of_memory_order():
diff --git a/numba_dpex/tests/experimental/test_index_space_ids.py b/numba_dpex/tests/experimental/test_index_space_ids.py
new file mode 100644
index 0000000000..54ceca8807
--- /dev/null
+++ b/numba_dpex/tests/experimental/test_index_space_ids.py
@@ -0,0 +1,270 @@
+# SPDX-FileCopyrightText: 2022 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Tests different input array type support for the kernel."""
+
+import dpnp
+import numpy as np
+import pytest
+from numba.core.errors import TypingError
+
+import numba_dpex as dpex
+import numba_dpex.experimental as dpex_exp
+from numba_dpex.kernel_api import Item, NdItem, NdRange
+from numba_dpex.kernel_api import call_kernel as kapi_call_kernel
+from numba_dpex.tests._helper import skip_windows
+
+_SIZE = 16
+_GROUP_SIZE = 4
+
+
+@dpex_exp.kernel
+def set_ones_no_item(a):
+    a[0] = 1
+
+
+@dpex_exp.kernel
+def set_ones_item(item: Item, a):
+    i = item.get_id(0)
+    a[i] = 1
+
+
+@dpex_exp.kernel
+def set_last_one_item(item: Item, a):
+    i = item.get_range(0) - 1
+    a[i] = 1
+
+
+@dpex_exp.kernel
+def set_last_one_nd_item(item: NdItem, a):
+    if item.get_global_id(0) == 0:
+        i = item.get_global_range(0) - 1
+        a[0] = i
+        a[i] = 1
+
+
+@dpex_exp.kernel
+def set_last_group_one_nd_item(item: NdItem, a):
+    if item.get_global_id(0) == 0:
+        i = item.get_local_range(0) - 1
+        a[0] = i
+        a[i] = 1
+
+
+@dpex_exp.kernel
+def set_ones_nd_item(nd_item: NdItem, a):
+    i = nd_item.get_global_id(0)
+    a[i] = 1
+
+
+@dpex_exp.kernel
+def set_local_ones_nd_item(nd_item: NdItem, a):
+    i = nd_item.get_local_id(0)
+    a[i] = 1
+
+
+def _get_group_id_driver(nditem: NdItem, a):
+    i = nditem.get_global_id(0)
+    g = nditem.get_group()
+    a[i] = g.get_group_id(0)
+
+
+def _get_group_range_driver(nditem: NdItem, a):
+    i = nditem.get_global_id(0)
+    g = nditem.get_group()
+    a[i] = g.get_group_range(0)
+
+
+def _get_group_local_range_driver(nditem: NdItem, a):
+    i = nditem.get_global_id(0)
+    g = nditem.get_group()
+    a[i] = g.get_local_range(0)
+
+
+def test_item_get_id():
+    a = dpnp.zeros(_SIZE, dtype=dpnp.float32)
+    dpex_exp.call_kernel(set_ones_item, dpex.Range(a.size), a)
+
+    assert np.array_equal(a.asnumpy(), np.ones(a.size, dtype=np.float32))
+
+
+def test_item_get_range():
+    a = dpnp.zeros(_SIZE, dtype=dpnp.float32)
+    dpex_exp.call_kernel(set_last_one_item, dpex.Range(a.size), a)
+
+    want = np.zeros(a.size, dtype=np.float32)
+    want[-1] = 1
+
+    assert np.array_equal(a.asnumpy(), want)
+
+
+# TODO: https://github.com/IntelPython/numba-dpex/issues/1308
+@skip_windows
+def test_nd_item_get_global_range():
+    a = dpnp.zeros(_SIZE, dtype=dpnp.float32)
+    dpex_exp.call_kernel(
+        set_last_one_nd_item, dpex.NdRange((a.size,), (_GROUP_SIZE,)), a
+    )
+
+    want = np.zeros(a.size, dtype=np.float32)
+    want[-1] = 1
+    want[0] = a.size - 1
+
+    assert np.array_equal(a.asnumpy(), want)
+
+
+# TODO: https://github.com/IntelPython/numba-dpex/issues/1308
+@skip_windows
+def test_nd_item_get_local_range():
+    a = dpnp.zeros(_SIZE, dtype=dpnp.float32)
+    dpex_exp.call_kernel(
+        set_last_group_one_nd_item, dpex.NdRange((a.size,), (_GROUP_SIZE,)), a
+    )
+
+    want = np.zeros(a.size, dtype=np.float32)
+    want[_GROUP_SIZE - 1] = 1
+    want[0] = _GROUP_SIZE - 1
+
+    assert np.array_equal(a.asnumpy(), want)
+
+
+# TODO: https://github.com/IntelPython/numba-dpex/issues/1308
+@skip_windows
+def test_nd_item_get_global_id():
+    a = dpnp.zeros(_SIZE, dtype=dpnp.float32)
+    dpex_exp.call_kernel(
+        set_ones_nd_item, dpex.NdRange((a.size,), (_GROUP_SIZE,)), a
+    )
+
+    assert np.array_equal(a.asnumpy(), np.ones(a.size, dtype=np.float32))
+
+
+# TODO: https://github.com/IntelPython/numba-dpex/issues/1308
+@skip_windows
+def test_nd_item_get_local_id():
+    a = dpnp.zeros(_SIZE, dtype=dpnp.float32)
+
+    dpex_exp.call_kernel(
+        set_local_ones_nd_item, dpex.NdRange((a.size,), (_GROUP_SIZE,)), a
+    )
+
+    assert np.array_equal(
+        a.asnumpy(),
+        np.array(
+            [1] * _GROUP_SIZE + [0] * (a.size - _GROUP_SIZE),
+            dtype=np.float32,
+        ),
+    )
+
+
+def test_error_item_get_global_id():
+    a = dpnp.zeros(_SIZE, dtype=dpnp.float32)
+
+    with pytest.raises(TypingError):
+        dpex_exp.call_kernel(set_ones_nd_item, dpex.Range(a.size), a)
+
+
+def test_no_item():
+    a = dpnp.zeros(_SIZE, dtype=dpnp.float32)
+    dpex_exp.call_kernel(set_ones_no_item, dpex.Range(a.size), a)
+
+    assert np.array_equal(
+        a.asnumpy(), np.array([1] + [0] * (a.size - 1), dtype=np.float32)
+    )
+
+
+# TODO: https://github.com/IntelPython/numba-dpex/issues/1308
+@skip_windows
+def test_get_group_id():
+    global_size = 100
+    group_size = 20
+    num_groups = global_size // group_size
+
+    a = dpnp.empty(global_size, dtype=dpnp.int32)
+    ka = dpnp.empty(global_size, dtype=dpnp.int32)
+    expected = np.empty(global_size, dtype=np.int32)
+    ndrange = NdRange((global_size,), (group_size,))
+    dpex_exp.call_kernel(dpex_exp.kernel(_get_group_id_driver), ndrange, a)
+    kapi_call_kernel(_get_group_id_driver, ndrange, ka)
+
+    for gid in range(num_groups):
+        for lid in range(group_size):
+            expected[gid * group_size + lid] = gid
+
+    assert np.array_equal(a.asnumpy(), expected)
+    assert np.array_equal(ka.asnumpy(), expected)
+
+
+# TODO: https://github.com/IntelPython/numba-dpex/issues/1308
+@skip_windows
+def test_get_group_range():
+    global_size = 100
+    group_size = 20
+    num_groups = global_size // group_size
+
+    a = dpnp.empty(global_size, dtype=dpnp.int32)
+    ka = dpnp.empty(global_size, dtype=dpnp.int32)
+    expected = np.empty(global_size, dtype=np.int32)
+    ndrange = NdRange((global_size,), (group_size,))
+    dpex_exp.call_kernel(dpex_exp.kernel(_get_group_range_driver), ndrange, a)
+    kapi_call_kernel(_get_group_range_driver, ndrange, ka)
+
+    for gid in range(num_groups):
+        for lid in range(group_size):
+            expected[gid * group_size + lid] = num_groups
+
+    assert np.array_equal(a.asnumpy(), expected)
+    assert np.array_equal(ka.asnumpy(), expected)
+
+
+# TODO: https://github.com/IntelPython/numba-dpex/issues/1308
+@skip_windows
+def test_get_group_local_range():
+    global_size = 100
+    group_size = 20
+    num_groups = global_size // group_size
+
+    a = dpnp.empty(global_size, dtype=dpnp.int32)
+    ka = dpnp.empty(global_size, dtype=dpnp.int32)
+    expected = np.empty(global_size, dtype=np.int32)
+    ndrange = NdRange((global_size,), (group_size,))
+    dpex_exp.call_kernel(
+        dpex_exp.kernel(_get_group_local_range_driver), ndrange, a
+    )
+    kapi_call_kernel(_get_group_local_range_driver, ndrange, ka)
+
+    for gid in range(num_groups):
+        for lid in range(group_size):
+            expected[gid * group_size + lid] = group_size
+
+    assert np.array_equal(a.asnumpy(), expected)
+    assert np.array_equal(ka.asnumpy(), expected)
+
+
+I_SIZE, J_SIZE, K_SIZE = 2, 3, 4
+
+
+@dpex_exp.kernel
+def set_3d_ones_item(item: Item, a):
+    i = item.get_id(0)
+    j = item.get_id(1)
+    k = item.get_id(2)
+
+    # Since we have different sizes for each dimention, wrong order will result
+    # that some indexes will be set twice and some won't be set.
+    index = i + I_SIZE * (j + J_SIZE * k)
+
+    a[index] = 1
+
+
+# TODO: CI tests failing for some reason... Works fine locally on cpu and gpu
+@pytest.mark.skip
+def test_index_order():
+    a = dpnp.zeros(I_SIZE * J_SIZE * K_SIZE, dtype=dpnp.int32)
+
+    dpex_exp.call_kernel(
+        set_3d_ones_item, dpex.Range(I_SIZE, J_SIZE, K_SIZE), a
+    )
+
+    assert np.array_equal(a.asnumpy(), np.ones(a.size, dtype=np.int32))
diff --git a/numba_dpex/tests/experimental/test_kernel_specialization.py b/numba_dpex/tests/experimental/test_kernel_specialization.py
index 2d1fbbaa02..45275f8f44 100644
--- a/numba_dpex/tests/experimental/test_kernel_specialization.py
+++ b/numba_dpex/tests/experimental/test_kernel_specialization.py
@@ -11,7 +11,7 @@
 import numba_dpex.experimental as dpex_exp
 from numba_dpex import DpnpNdArray, float32, int64
 from numba_dpex.core.exceptions import InvalidKernelSpecializationError
-from numba_dpex.core.kernel_interface.indexers import Range
+from numba_dpex.kernel_api import Range
 
 i64arrty = DpnpNdArray(ndim=1, dtype=int64, layout="C")
 f32arrty = DpnpNdArray(ndim=1, dtype=float32, layout="C")
diff --git a/numba_dpex/tests/kernel_api/__init__.py b/numba_dpex/tests/kernel_api/__init__.py
new file mode 100644
index 0000000000..ee20629a35
--- /dev/null
+++ b/numba_dpex/tests/kernel_api/__init__.py
@@ -0,0 +1,3 @@
+# SPDX-FileCopyrightText: 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
diff --git a/numba_dpex/tests/kernel_api/test_ndrange_kernel_launch.py b/numba_dpex/tests/kernel_api/test_ndrange_kernel_launch.py
new file mode 100644
index 0000000000..d2bf75f8dd
--- /dev/null
+++ b/numba_dpex/tests/kernel_api/test_ndrange_kernel_launch.py
@@ -0,0 +1,52 @@
+# SPDX-FileCopyrightText: 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import numpy
+
+from numba_dpex import kernel_api as kapi
+
+
+def test_range_kernel_call1D():
+    def vecadd(item: kapi.NdItem, a, b, c):
+        idx = item.get_global_id(0)
+        c[idx] = a[idx] + b[idx]
+
+    a = numpy.ones(100)
+    b = numpy.ones(100)
+    c = numpy.empty(100)
+
+    kapi.call_kernel(vecadd, kapi.NdRange((100,), (20,)), a, b, c)
+
+    assert numpy.allclose(c, a + b)
+
+
+def test_range_kernel_call2D():
+    def vecadd(item: kapi.NdItem, a, b, c):
+        idx = item.get_global_id(0)
+        jdx = item.get_global_id(1)
+        c[idx, jdx] = a[idx, jdx] + b[idx, jdx]
+
+    a = numpy.ones((10, 10))
+    b = numpy.ones((10, 10))
+    c = numpy.empty((10, 10))
+
+    kapi.call_kernel(vecadd, kapi.NdRange((10, 10), (2, 2)), a, b, c)
+
+    assert numpy.allclose(c, a + b)
+
+
+def test_range_kernel_call3D():
+    def vecadd(item: kapi.Item, a, b, c):
+        idx = item.get_global_id(0)
+        jdx = item.get_global_id(1)
+        kdx = item.get_global_id(2)
+        c[idx, jdx, kdx] = a[idx, jdx, kdx] + b[idx, jdx, kdx]
+
+    a = numpy.ones((8, 8, 8))
+    b = numpy.ones((8, 8, 8))
+    c = numpy.empty((8, 8, 8))
+
+    kapi.call_kernel(vecadd, kapi.NdRange((8, 8, 8), (2, 2, 2)), a, b, c)
+
+    assert numpy.allclose(c, a + b)
diff --git a/numba_dpex/tests/kernel_api/test_range_kernel_launch.py b/numba_dpex/tests/kernel_api/test_range_kernel_launch.py
new file mode 100644
index 0000000000..04c41ea4b6
--- /dev/null
+++ b/numba_dpex/tests/kernel_api/test_range_kernel_launch.py
@@ -0,0 +1,52 @@
+# SPDX-FileCopyrightText: 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import numpy
+
+from numba_dpex import kernel_api as kapi
+
+
+def test_range_kernel_call1D():
+    def vecadd(item: kapi.Item, a, b, c):
+        idx = item.get_id(0)
+        c[idx] = a[idx] + b[idx]
+
+    a = numpy.ones(100)
+    b = numpy.ones(100)
+    c = numpy.empty(100)
+
+    kapi.call_kernel(vecadd, kapi.Range(100), a, b, c)
+
+    assert numpy.allclose(c, a + b)
+
+
+def test_range_kernel_call2D():
+    def vecadd(item: kapi.Item, a, b, c):
+        idx = item.get_id(0)
+        jdx = item.get_id(1)
+        c[idx, jdx] = a[idx, jdx] + b[idx, jdx]
+
+    a = numpy.ones((10, 10))
+    b = numpy.ones((10, 10))
+    c = numpy.empty((10, 10))
+
+    kapi.call_kernel(vecadd, kapi.Range(10, 10), a, b, c)
+
+    assert numpy.allclose(c, a + b)
+
+
+def test_range_kernel_call3D():
+    def vecadd(item: kapi.Item, a, b, c):
+        idx = item.get_id(0)
+        jdx = item.get_id(1)
+        kdx = item.get_id(2)
+        c[idx, jdx, kdx] = a[idx, jdx, kdx] + b[idx, jdx, kdx]
+
+    a = numpy.ones((5, 5, 5))
+    b = numpy.ones((5, 5, 5))
+    c = numpy.empty((5, 5, 5))
+
+    kapi.call_kernel(vecadd, kapi.Range(5, 5, 5), a, b, c)
+
+    assert numpy.allclose(c, a + b)
diff --git a/numba_dpex/tests/kernel_tests/test_kernel_launch_params.py b/numba_dpex/tests/kernel_tests/test_kernel_launch_params.py
index 6505401030..2d2c564f41 100644
--- a/numba_dpex/tests/kernel_tests/test_kernel_launch_params.py
+++ b/numba_dpex/tests/kernel_tests/test_kernel_launch_params.py
@@ -7,12 +7,8 @@
 import pytest
 
 import numba_dpex as dpex
-from numba_dpex.core.exceptions import (
-    IllegalRangeValueError,
-    InvalidKernelLaunchArgsError,
-    UnknownGlobalRangeError,
-)
-from numba_dpex.core.kernel_interface.indexers import Range
+from numba_dpex.core.exceptions import UnknownGlobalRangeError
+from numba_dpex.kernel_api import NdRange, Range
 
 
 @dpex.kernel
@@ -28,7 +24,7 @@ def test_1D_global_range_as_one_tuple():
 
 
 def test_2D_global_range_and_2D_local_range4():
-    k = vecadd[dpex.NdRange((10, 10), (10, 10))]
+    k = vecadd[NdRange((10, 10), (10, 10))]
     assert k._global_range == [10, 10]
     assert k._local_range == [10, 10]
 
diff --git a/numba_dpex/tests/kernel_tests/test_kernel_specialization.py b/numba_dpex/tests/kernel_tests/test_kernel_specialization.py
index 276c8fec0f..99076461f6 100644
--- a/numba_dpex/tests/kernel_tests/test_kernel_specialization.py
+++ b/numba_dpex/tests/kernel_tests/test_kernel_specialization.py
@@ -11,7 +11,7 @@
     InvalidKernelSpecializationError,
     MissingSpecializationError,
 )
-from numba_dpex.core.kernel_interface.indexers import Range
+from numba_dpex.kernel_api import Range
 
 i64arrty = usm_ndarray(ndim=1, dtype=int64, layout="C")
 f32arrty = usm_ndarray(ndim=1, dtype=float32, layout="C")
diff --git a/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py b/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py
index 079cc551b3..a73f18767a 100644
--- a/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py
+++ b/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py
@@ -6,7 +6,7 @@
 import pytest
 
 import numba_dpex as ndpx
-from numba_dpex.core.kernel_interface.indexers import NdRange
+from numba_dpex.kernel_api import NdRange
 
 
 # Data parallel kernel implementing vector sum
diff --git a/numba_dpex/tests/test_vectorize.py b/numba_dpex/tests/test_vectorize.py
deleted file mode 100644
index ec25ce9864..0000000000
--- a/numba_dpex/tests/test_vectorize.py
+++ /dev/null
@@ -1,65 +0,0 @@
-#! /usr/bin/env python
-
-# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numpy as np
-import pytest
-from numba import float32, float64, int32, int64, vectorize
-
-list_of_shape = [
-    (100, 100),
-    (100, (10, 10)),
-    (100, (2, 5, 10)),
-]
-
-
-@pytest.fixture(params=list_of_shape)
-def shape(request):
-    return request.param
-
-
-list_of_dtype = [
-    (np.int32, int32),
-    (np.float32, float32),
-    (np.int64, int64),
-    (np.float64, float64),
-]
-
-
-@pytest.fixture(params=list_of_dtype)
-def dtypes(request):
-    return request.param
-
-
-list_of_input_type = ["array", "scalar"]
-
-
-@pytest.fixture(params=list_of_input_type)
-def input_type(request):
-    return request.param
-
-
-@pytest.mark.xfail
-def test_vectorize(shape, dtypes, input_type):
-    def vector_add(a, b):
-        return a + b
-
-    dtype, sig_dtype = dtypes
-    sig = [sig_dtype(sig_dtype, sig_dtype)]
-    size, shape = shape
-
-    if input_type == "array":
-        A = np.arange(size, dtype=dtype).reshape(shape)
-        B = np.arange(size, dtype=dtype).reshape(shape)
-    elif input_type == "scalar":
-        A = dtype(1.2)
-        B = dtype(2.3)
-
-    f = vectorize(sig, target="dpex")(vector_add)
-    expected = f(A, B)
-    actual = vector_add(A, B)
-
-    max_abs_err = np.sum(expected) - np.sum(actual)
-    assert max_abs_err < 1e-5