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

Issue:
Updated CHANGELOG for 0.21.2
Updated change log for 0.21.2


Discussion / hints:

Fix patch:
diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
index ffec74c911..4733d9f49c 100644
--- a/.github/CODEOWNERS
+++ b/.github/CODEOWNERS
@@ -1,5 +1,5 @@
-*.yml   @mingjie-intel @diptorupd
-conda-recipe/*  @mingjie-intel @diptorupd
+*.yml  @diptorupd
+conda-recipe/*  @diptorupd
 numba_dpex/dpnp_iface @diptorupd @chudur-budur
 numba_dpex/core @diptorupd
-numba_dpex/core/passes/parfor_lowering_pass.py @mingjie-intel
+numba_dpex/core/passes/parfor_lowering_pass.py @diptorupd
diff --git a/.github/actions/setup-conda-build/action.yml b/.github/actions/setup-conda-build/action.yml
deleted file mode 100644
index 7bf8755173..0000000000
--- a/.github/actions/setup-conda-build/action.yml
+++ /dev/null
@@ -1,9 +0,0 @@
-name: Setup conda-build
-description: 'Activate conda and install conda-build'
-runs:
-  using: "composite"
-  steps:
-    - run: echo $CONDA/bin >> $GITHUB_PATH
-      shell: bash
-    - run: conda install conda-build
-      shell: bash
diff --git a/.github/actions/setup-numba-dpex/action.yml b/.github/actions/setup-numba-dpex/action.yml
deleted file mode 100644
index 6db253c703..0000000000
--- a/.github/actions/setup-numba-dpex/action.yml
+++ /dev/null
@@ -1,31 +0,0 @@
-name: Setup numba-dpex
-description: 'Setup conda environment, build numba-dpex and setup devices'
-inputs:
-  environment:
-    description: Environment file
-    required: true
-    default: environment/dev.yml
-runs:
-  using: "composite"
-  steps:
-    - name: Setup conda
-      run: echo $CONDA/bin >> $GITHUB_PATH
-      shell: bash
-
-    - name: Setup conda environment
-      run: conda env update -n base -f ${{ inputs.environment }} --prune
-      shell: bash
-
-    - name: Test conda environment
-      run: conda list
-      shell: bash
-
-    - name: Build numba-dpex
-      run: |
-        export PATH=$CONDA/bin-llvm:$PATH
-        python setup.py develop
-      shell: bash
-
-    - name: Setup OpenCL CPU device
-      run: echo "OCL_ICD_FILENAMES=libintelocl.so" >> $GITHUB_ENV
-      shell: bash
diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml
index d6932d684b..05be76ddc2 100644
--- a/.github/workflows/conda-package.yml
+++ b/.github/workflows/conda-package.yml
@@ -13,55 +13,37 @@ on:
 env:
   PACKAGE_NAME: numba-dpex
   MODULE_NAME: numba_dpex
+  # There is a separate action that removes defaults.
+  CHANNELS: 'dppy/label/dev,conda-forge,intel'
+  VER_JSON_NAME: 'version.json'
+  VER_SCRIPT1: "import json; f = open('version.json', 'r'); j = json.load(f); f.close(); "
+  VER_SCRIPT2: "d = j['numba-dpex'][0]; print('='.join((d[s] for s in ('version', 'build'))))"
 
 jobs:
-  build_linux:
-    runs-on: ubuntu-20.04
+  build:
+    name: Build ['${{ matrix.os }}', python='${{ matrix.python }}']
 
     strategy:
+      fail-fast: false
       matrix:
-        python: ["3.8", "3.9", "3.10"]
+        python: ['3.8', '3.9', '3.10']
+        os: [ubuntu-latest, windows-latest]
 
-    steps:
-      - name: Checkout repo
-        uses: actions/checkout@v3
-        with:
-          fetch-depth: 0
+    runs-on: ${{ matrix.os }}
 
-      - name: Setup conda-build
-        uses: ./.github/actions/setup-conda-build
+    defaults:
+      run:
+        shell: bash -l {0}
 
-      - name: Build conda package
-        env:
-          CHANNELS: -c dppy/label/dev -c intel -c defaults -c numba -c numba/label/dev --override-channels
-        run: ./scripts/build_conda_package.sh ${{ matrix.python }}
-
-      - name: Dump build linux conda environment
-        run: |
-          conda env export > /tmp/env-build-linux.yml
-          cat /tmp/env-build-linux.yml
+    continue-on-error: false
 
-      - name: Upload artifact
-        uses: actions/upload-artifact@v3
+    steps:
+      - name: Cancel Previous Runs
+        uses: styfle/cancel-workflow-action@0.11.0
         with:
-          name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Python ${{ matrix.python }}
-          path: /usr/share/miniconda/conda-bld/linux-64/${{ env.PACKAGE_NAME }}-*.tar.bz2
+          access_token: ${{ github.token }}
 
-  build_windows:
-    runs-on: windows-latest
-
-    strategy:
-      matrix:
-        python: ["3.8", "3.9", "3.10"]
-        experimental: [false]
-        artifact_name: [""]
-    continue-on-error: ${{ matrix.experimental }}
-    env:
-      CHANNELS: -c dppy/label/dev -c intel -c defaults -c numba -c numba/label/dev --override-channels
-      conda-bld: C:\Miniconda\conda-bld\win-64\
-
-    steps:
-      - name: Checkout repo
+      - name: Checkout ${{ env.PACKAGE_NAME }} repo
         uses: actions/checkout@v3
         with:
           fetch-depth: 0
@@ -69,253 +51,165 @@ jobs:
       - name: Setup miniconda
         uses: conda-incubator/setup-miniconda@v2
         with:
-          auto-activate-base: true
-          activate-environment: ""
+          python-version: ${{ inputs.python-version }}
+          miniforge-variant: Mambaforge
+          miniforge-version: latest
+          activate-environment: "build"
+          channels: ${{ env.CHANNELS }}
+          channel-priority: "disabled"
+          run-post: false
 
-      - name: Cache conda packages
-        uses: actions/cache@v3
-        env:
-          CACHE_NUMBER: 0  # Increase to reset cache
-        with:
-          path: /home/runner/conda_pkgs_dir
-          key:
-            ${{ runner.os }}-conda-${{ env.CACHE_NUMBER }}-python-${{ matrix.python }}-${{hashFiles('**/meta.yaml') }}
-          restore-keys: |
-            ${{ runner.os }}-conda-${{ env.CACHE_NUMBER }}-python-${{ matrix.python }}-
-            ${{ runner.os }}-conda-${{ env.CACHE_NUMBER }}-
+      - name: Disable defautls
+        run: conda config --remove channels defaults
+
+      - name: Store conda paths as envs
+        run: echo "CONDA_BLD=$CONDA_PREFIX/conda-bld/${{ runner.os == 'Linux' && 'linux' || 'win' }}-64/" | tr "\\" '/' >> $GITHUB_ENV
 
+      # boa is an extention to conda so we can use mamba resolver in conda build
       - name: Install conda-build
-        run: conda install conda-build
+        run: mamba install boa
 
       - name: Build conda package
-        run: conda build --no-test --python ${{ matrix.python }} ${{ env.CHANNELS }} conda-recipe
-
-      - name: Dump build windows conda environment
-        run: |
-          conda env export > env-build-win.yml
-          type env-build-win.yml
+        run: conda mambabuild --no-test --python ${{ matrix.python }} conda-recipe
 
       - name: Upload artifact
         uses: actions/upload-artifact@v3
-        with:
-          name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Python ${{ matrix.python }} ${{ matrix.artifact_name }}
-          path: ${{ env.conda-bld }}${{ env.PACKAGE_NAME }}-*.tar.bz2
-
-  test_linux:
-    needs: build_linux
-    runs-on: ubuntu-20.04
-
-    strategy:
-      matrix:
-        python: ["3.8", "3.9", "3.10"]
-        numba: ["0.57"]
-        dpnp: ["0.11"]
-
-    steps:
-      - name: Download artifact
-        uses: actions/download-artifact@v3
         with:
           name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Python ${{ matrix.python }}
+          path: ${{ env.CONDA_BLD }}${{ env.PACKAGE_NAME }}-*.tar.bz2
 
-      - name: Checkout repo
-        uses: actions/checkout@v3
-        with:
-          path: numba-dpex
-
-      - name: Setup conda-build
-        uses: ./numba-dpex/.github/actions/setup-conda-build
 
-      - name: Create conda channel
-        run: |
-          mkdir -p $GITHUB_WORKSPACE/channel/linux-64
-          mv ${PACKAGE_NAME}-*.tar.bz2 $GITHUB_WORKSPACE/channel/linux-64
-          conda index $GITHUB_WORKSPACE/channel
-          # Test channel
-          conda search $PACKAGE_NAME -c $GITHUB_WORKSPACE/channel --override-channels --info
-
-      - name: Install numba-dpex
-        env:
-          CHANNELS: -c dppy/label/dev -c intel -c defaults -c numba -c numba/label/dev --override-channels
-        run: |
-          CHANNELS="-c $GITHUB_WORKSPACE/channel $CHANNELS"
-          conda list
-          conda create -n numba_dpex_env $PACKAGE_NAME pytest dpcpp_linux-64 python=${{ matrix.python }} numba=${{ matrix.numba }} dpctl dpnp=${{ matrix.dpnp }} dpcpp-llvm-spirv $CHANNELS
-          # Test installed packages
-          conda list
+  test:
+    name: Test ['${{ matrix.os }}', python='${{ matrix.python }}']
 
-      - name: Check DPNP
-        run: |
-          source $CONDA/etc/profile.d/conda.sh
-          conda activate numba_dpex_env
-          python -c "import dpnp"
+    needs: build
 
-      - name: Check dpcpp-llvm-spirv
-        run: |
-          source $CONDA/etc/profile.d/conda.sh
-          conda activate numba_dpex_env
-          python -c "import dpcpp_llvm_spirv as p; print(p.get_llvm_spirv_path())"
-
-      - name: Dump test linux conda environment
-        run: |
-          conda env export > /tmp/env-test-linux.yml
-          cat /tmp/env-test-linux.yml
-
-      - name: Run tests
-        run: |
-          source $CONDA/etc/profile.d/conda.sh
-          conda activate numba_dpex_env
-          # echo "libintelocl.so" | tee /etc/OpenCL/vendors/intel-cpu.icd
-          pytest -q -ra --disable-warnings --pyargs $MODULE_NAME -vv
-
-      - name: Run examples
-        run: |
-          ls
-          pwd
-          cd ./numba-dpex/numba_dpex/examples/
-          source $CONDA/etc/profile.d/conda.sh
-          conda activate numba_dpex_env
-          # echo "libintelocl.so" | tee /etc/OpenCL/vendors/intel-cpu.icd
-          for script in $(find . \( -not -name "_*" -not -name "side-by-side*" -not -name "vectorize.py" -not -name "scan.py" -and -name "*.py" \))
-          do
-            echo "Executing ${script}"
-            python ${script} || exit 1
-          done
+    runs-on: ${{ matrix.os }}
 
-  test_windows:
-    needs: build_windows
-    runs-on: windows-latest
+    defaults:
+      run:
+        shell: ${{ matrix.os == 'windows-latest' && 'cmd /C CALL {0}' || 'bash -l {0}' }}
 
     strategy:
+      fail-fast: false
       matrix:
-        python: ["3.8", "3.9", "3.10"]
-        integration_channels: [""]
-        experimental: [false]  # packages are not available on -c intel yet
+        python: ['3.8', '3.9', '3.10']
+        os: [ubuntu-20.04, ubuntu-latest, windows-latest]
+        experimental: [false]
+
     continue-on-error: ${{ matrix.experimental }}
-    env:
-      # conda-forge: llvm-spirv 11 not on intel channel yet
-      CHANNELS: -c dppy/label/dev -c intel -c defaults -c numba -c numba/label/dev --override-channels
 
     steps:
-      - name: Create dir for numba-dpex repo
-        run: |
-          mkdir -p ${{ github.workspace }}/dpex-repo
-
-      - name: Checout repo
-        uses: actions/checkout@v3
+      - name: Setup miniconda
+        uses: conda-incubator/setup-miniconda@v2
         with:
-          path: dpex-repo
-          fetch-depth: 0
+          python-version: ${{ matrix.python }}
+          miniforge-variant: Mambaforge
+          miniforge-version: latest
+          activate-environment: "test"
+          channels: ${{ env.CHANNELS }}
+          channel-priority: "disabled"
+          run-post: false
+
+      - name: Disable defautls
+        run: conda config --remove channels defaults
+
+      - name: Store conda paths as envs
+        shell: bash -l {0}
+        run: |
+          echo "CHANNEL_PATH=${{ github.workspace }}/channel/" | tr "\\" "/" >> $GITHUB_ENV
+          echo "EXTRACTED_PKG_PATH=${{ github.workspace }}/pkg/" | tr "\\" "/" >> $GITHUB_ENV
+          echo "VER_JSON_PATH=${{ github.workspace }}/version.json" | tr "\\" "/" >> $GITHUB_ENV
+          echo "PKG_PATH_IN_CHANNEL=${{ github.workspace }}/channel/${{ runner.os == 'Linux' && 'linux' || 'win' }}-64/" | tr "\\" "/" >> $GITHUB_ENV
+          echo "EXAMPLES_PATH=$CONDA_PREFIX/${{ runner.os == 'Linux' && format('lib/python{0}',matrix.python) || 'Lib' }}/site-packages/numba_dpex/examples/" | tr "\\" "/" >> $GITHUB_ENV
 
       - name: Download artifact
         uses: actions/download-artifact@v3
         with:
           name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Python ${{ matrix.python }}
+          path: ${{ env.PKG_PATH_IN_CHANNEL }}
 
-      - name: Setup miniconda
-        uses: conda-incubator/setup-miniconda@v2
-        with:
-          use-only-tar-bz2: true
-          auto-update-conda: true
-          conda-build-version: '*'
-          miniconda-version: 'latest'
-          activate-environment: "numba_dpex_wintest"
+      # We need --force-local because on windows path looks like C:/foo/bar
+      # and tar interprits semicolons as remote device.
+      - name: Extract package archive
+        shell: bash -l {0}
+        run: |
+          mkdir -p ${EXTRACTED_PKG_PATH}
+          tar -xvf ${PKG_PATH_IN_CHANNEL}${PACKAGE_NAME}-*.tar.bz2 -C ${EXTRACTED_PKG_PATH} --force-local
+
+      - name: Install conda build
+        run: |
+          mamba install conda-build
 
       - name: Create conda channel
         run: |
-          mkdir -p $env:GITHUB_WORKSPACE/channel/win-64
-          mv ${{ env.PACKAGE_NAME }}-*.tar.bz2 $env:GITHUB_WORKSPACE/channel/win-64
-          conda index $env:GITHUB_WORKSPACE/channel
-          # Test channel
-          conda search ${{ env.PACKAGE_NAME }} -c $env:GITHUB_WORKSPACE/channel --override-channels
+          conda index ${{ env.CHANNEL_PATH }}
 
-      - name: Collect dependencies
-        run: conda install ${{ env.PACKAGE_NAME }} python=${{ matrix.bpython }} -c $env:GITHUB_WORKSPACE/channel ${{ env.CHANNELS }} --only-deps --dry-run > lockfile
+      - name: Test conda channel
+        run: |
+          mamba search ${{ env.PACKAGE_NAME }} -c ${{ env.CHANNEL_PATH }} --override-channels --info --json > ${{ env.VER_JSON_PATH }}
+          cat ${{ env.VER_JSON_PATH }}
 
-      - name: Cache conda packages
-        uses: actions/cache@v3
-        env:
-          CACHE_NUMBER: 1  # Increase to reset cache
-        with:
-          path: /home/runner/conda_pkgs_dir
-          key:
-            ${{ runner.os }}-conda-${{ env.CACHE_NUMBER }}-python-${{ matrix.python }}-${{hashFiles('lockfile') }}
-          restore-keys: |
-            ${{ runner.os }}-conda-${{ env.CACHE_NUMBER }}-python-${{ matrix.python }}-
-            ${{ runner.os }}-conda-${{ env.CACHE_NUMBER }}-
-
-      - name: Install numba-dpex
+      - name: Collect dependencies
+        shell: bash -l {0}
         run: |
-          conda install ${{ env.PACKAGE_NAME }} pytest dpcpp_win-64 dpcpp-llvm-spirv python=${{ matrix.python }} dpctl -c $env:GITHUB_WORKSPACE/channel ${{ env.CHANNELS }}
-          # Test installed packages
-          conda list
+          export PACKAGE_VERSION=$(python -c "${{ env.VER_SCRIPT1 }} ${{ env.VER_SCRIPT2 }}")
+          echo "PACKAGE_VERSION=$PACKAGE_VERSION" >> $GITHUB_ENV
 
-      - name: Install opencl_rt
-        run: conda install opencl_rt -c intel --override-channels
+      # We want to make sure that all dependecies install automatically.
+      - name: Install builded package
+        run: mamba install ${{ env.PACKAGE_NAME }}=${{ env.PACKAGE_VERSION }} opencl_rt pytest -c ${{ env.CHANNEL_PATH }}
 
-      - name: Add library
+      - name: Setup OpenCL CPU device
+        if: runner.os == 'Windows'
         shell: pwsh
         run: |
           $script_path="$env:CONDA_PREFIX\Scripts\set-intel-ocl-icd-registry.ps1"
           &$script_path
-          # Check the variable assisting OpenCL CPU driver to find TBB DLLs which are not located where it expects them by default
-          $cl_cfg="$env:CONDA_PREFIX\Library\lib\cl.cfg"
-          Get-Content -Tail 5 -Path $cl_cfg
 
-      - name: Add dpnp skip variable
-        run: echo "NUMBA_DPEX_TESTING_SKIP_NO_DPNP=1" | Out-File -FilePath $env:GITHUB_ENV -Encoding utf8 -Append
+      - name: List installed packages
+        run: conda list
 
-      - name: Smoke test
-        run: python -c "import numba_dpex.core.runtime._dpexrt_python"
-
-      - name: Smoke test for dpcpp-llvm-spirv
-        run: python -c "import dpcpp_llvm_spirv as p; print(p.get_llvm_spirv_path())"
-
-      - name: Dump test windows conda environment
+      - name: Check dpcpp-llvm-spirv
         run: |
-          conda env export > env-test-win.yml
-          type env-test-win.yml
+          python -c "import dpcpp_llvm_spirv as p; print(p.get_llvm_spirv_path())"
+
+      - name: Smoke test
+        run: python -c "import dpnp, dpctl, numba_dpex; dpctl.lsplatform()"
 
       - name: Run tests
         run: |
           pytest -q -ra --disable-warnings --pyargs ${{ env.MODULE_NAME }} -vv
 
-  upload_linux:
-    needs: test_linux
+      - name: Run 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" \))
+          do
+            echo "Executing ${script}"
+            python ${script} || exit 1
+          done
+
+  upload_anaconda:
+    name: Upload dppy/label/dev ['${{ matrix.os }}', python='${{ matrix.python }}']
     if: ${{github.ref == 'refs/heads/main' || (startsWith(github.ref, 'refs/heads/release') == true)}}
-    runs-on: ubuntu-20.04
+
+    needs: [test]
+
     strategy:
       matrix:
-        python: ["3.8", "3.9", "3.10"]
-    steps:
-      - name: Download artifact
-        uses: actions/download-artifact@v3
-        with:
-          name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Python ${{ matrix.python }}
+        python: ['3.8', '3.9', '3.10']
+        os: [ubuntu-latest, windows-latest]
 
-      - name: Setup miniconda
-        uses: conda-incubator/setup-miniconda@v2
-        with:
-          auto-activate-base: true
-          activate-environment: ""
+    runs-on: ${{ matrix.os }}
 
-      - name: Install anaconda-client
-        run: conda install anaconda-client
+    defaults:
+      run:
+        shell: bash -l {0}
 
-      - name: Upload
-        env:
-          ANACONDA_TOKEN: ${{ secrets.ANACONDA_TOKEN }}
-        run: |
-          source /usr/share/miniconda/etc/profile.d/conda.sh
-          conda activate
-          anaconda --token $ANACONDA_TOKEN upload --user dppy --label dev ${PACKAGE_NAME}-*.tar.bz2
+    continue-on-error: false
 
-  upload_windows:
-    needs: test_windows
-    if: ${{github.ref == 'refs/heads/main' || (startsWith(github.ref, 'refs/heads/release') == true)}}
-    runs-on: windows-latest
-    strategy:
-      matrix:
-        python: ["3.8", "3.9", "3.10"]
     steps:
       - name: Download artifact
         uses: actions/download-artifact@v3
@@ -325,14 +219,16 @@ jobs:
       - name: Setup miniconda
         uses: conda-incubator/setup-miniconda@v2
         with:
-          auto-activate-base: true
-          activate-environment: ""
+          python-version: ${{ matrix.python }}
+          miniforge-variant: Mambaforge
+          miniforge-version: latest
+          activate-environment: "anaconda"
+          channels: ${{ env.CHANNELS }}
+          channel-priority: "disabled"
+          run-post: false
 
       - name: Install anaconda-client
-        run: conda install anaconda-client
+        run: mamba install anaconda-client
 
       - name: Upload
-        env:
-          ANACONDA_TOKEN: ${{ secrets.ANACONDA_TOKEN }}
-        run: |
-          anaconda --token ${{ env.ANACONDA_TOKEN }} upload --user dppy --label dev ${{ env.PACKAGE_NAME }}-*.tar.bz2
+        run: anaconda --token ${{ secrets.ANACONDA_TOKEN }} upload --user dppy --label dev ${{ env.PACKAGE_NAME }}-*.tar.bz2
diff --git a/.github/workflows/coverage.yml b/.github/workflows/coverage.yml
index ee1fbb0dd6..d615760451 100644
--- a/.github/workflows/coverage.yml
+++ b/.github/workflows/coverage.yml
@@ -12,7 +12,9 @@ jobs:
   main:
     name: Generate coverage and push to Coveralls.io
     runs-on: ubuntu-20.04
-
+    defaults:
+      run:
+        shell: bash -l {0}
     steps:
       - name: Cancel Previous Runs
         uses: styfle/cancel-workflow-action@0.9.1
@@ -24,10 +26,19 @@ jobs:
         with:
           fetch-depth: 0
 
-      - name: Setup numba-dpex
-        uses: ./.github/actions/setup-numba-dpex
+      - uses: conda-incubator/setup-miniconda@v2
         with:
-          environment: environment/coverage.yml
+          python-version: '3.10'
+          miniforge-variant: Mambaforge
+          miniforge-version: latest
+          activate-environment: "coverage"
+          channel-priority: "disabled"
+          environment-file: environment/coverage.yml
+
+      - name: Build numba-dpex
+        run: |
+          export PATH=$CONDA/bin-llvm:$PATH
+          python setup.py develop
 
       - name: Test installation
         run: |
diff --git a/.github/workflows/gh-pages-tags.yml b/.github/workflows/gh-pages-tags.yml
deleted file mode 100644
index efd250d0ad..0000000000
--- a/.github/workflows/gh-pages-tags.yml
+++ /dev/null
@@ -1,54 +0,0 @@
-name: GitHub Pages [Release, Tag]
-on:
-  push:
-    branches:
-      - release*
-    tags:
-      - '[0-9]+.[0-9]+*'
-
-jobs:
-  main:
-    runs-on: ubuntu-latest
-
-    steps:
-      - uses: actions/checkout@v3
-        with:
-          fetch-depth: 0
-
-      - name: Setup numba-dpex
-        uses: ./.github/actions/setup-numba-dpex
-        with:
-          environment: environment/docs.yml
-
-      - name: Make documentation
-        working-directory: docs
-        run: make html
-
-      - name: Publish release
-        if: startsWith(github.ref, 'refs/heads/release')
-        uses: peaceiris/actions-gh-pages@v3
-        with:
-          github_token: ${{ secrets.GITHUB_TOKEN }}
-          destination_dir : next_release
-          publish_dir: docs/_build/html/
-          allow_empty_commit : true
-          commit_message: ${{ github.event.head_commit.message }}
-
-
-      - name: Capture tag
-        id: capture_tag
-        if: startsWith(github.ref, 'refs/tags/')
-        run: |
-          echo "${GITHUB_REF#refs/heads/}"
-          echo "${GITHUB_REF#refs/tags/}"
-          echo ::set-output name=tag_number::${GITHUB_REF#refs/tags/}
-
-      - name: Publish tag
-        if: startsWith(github.ref, 'refs/tags/')
-        uses: peaceiris/actions-gh-pages@v3
-        with:
-          github_token: ${{ secrets.GITHUB_TOKEN }}
-          destination_dir : ${{ steps.capture_tag.outputs.tag_number }}
-          publish_dir: docs/_build/html/
-          allow_empty_commit : true
-          commit_message: ${{ github.event.head_commit.message }}
diff --git a/.github/workflows/gh-pages.yml b/.github/workflows/gh-pages.yml
index 9dba355135..920881a652 100644
--- a/.github/workflows/gh-pages.yml
+++ b/.github/workflows/gh-pages.yml
@@ -3,6 +3,9 @@ on:
   push:
     branches:
       - main
+      - release*
+    tags:
+      - '[0-9]+.[0-9]+*'
   pull_request:
     types: [opened, synchronize, reopened, closed]
     paths:
@@ -13,16 +16,29 @@ jobs:
   main:
     if: ${{ !(github.event.pull_request && github.event.action == 'closed') }}
     runs-on: ubuntu-latest
+    defaults:
+      run:
+        shell: bash -l {0}
 
     steps:
       - uses: actions/checkout@v3
         with:
           fetch-depth: 0
 
-      - name: Setup numba-dpex
-        uses: ./.github/actions/setup-numba-dpex
+      - uses: conda-incubator/setup-miniconda@v2
         with:
-          environment: environment/docs.yml
+          python-version: '3.10'
+          miniforge-variant: Mambaforge
+          miniforge-version: latest
+          activate-environment: "docs"
+          channels: conda-forge,nodefaults
+          channel-priority: "disabled"
+          environment-file: "environment/docs.yml"
+
+      - name: Build numba-dpex
+        run: |
+          export PATH=$CONDA/bin-llvm:$PATH
+          python setup.py develop
 
       - name: Make documentation
         working-directory: docs
@@ -58,13 +74,37 @@ jobs:
         if: ${{ github.event.pull_request && github.event.action != 'closed' }}
         env:
           PR_NUM: ${{ github.event.number }}
-        uses: mshick/add-pr-comment@v1
+        uses: mshick/add-pr-comment@v2
         with:
           message: |
             Documentation preview: [show](https://intelpython.github.io/numba-dpex/pull/${{ env.PR_NUM }}).
-          repo-token: ${{ secrets.GITHUB_TOKEN }}
-          repo-token-user-login: 'github-actions[bot]'
-          allow-repeats: true
+          # repo-token: ${{ secrets.GITHUB_TOKEN }}
+
+      - name: Publish release
+        if: startsWith(github.ref, 'refs/heads/release')
+        uses: peaceiris/actions-gh-pages@v3
+        with:
+          github_token: ${{ secrets.GITHUB_TOKEN }}
+          destination_dir : next_release
+          publish_dir: docs/_build/html/
+          allow_empty_commit : true
+          commit_message: ${{ github.event.head_commit.message }}
+
+      - name: Capture tag
+        id: capture_tag
+        if: startsWith(github.ref, 'refs/tags/')
+        run: |
+          echo ::set-output name=tag_number::${GITHUB_REF#refs/tags/}
+
+      - name: Publish tag
+        if: startsWith(github.ref, 'refs/tags/')
+        uses: peaceiris/actions-gh-pages@v3
+        with:
+          github_token: ${{ secrets.GITHUB_TOKEN }}
+          destination_dir : ${{ steps.capture_tag.outputs.tag_number }}
+          publish_dir: docs/_build/html/
+          allow_empty_commit : true
+          commit_message: ${{ github.event.head_commit.message }}
 
   clean:
     if: ${{ github.event.pull_request && github.event.action == 'closed' }}
diff --git a/CHANGELOG.md b/CHANGELOG.md
index 7b1f2d996e..0689a6e7ed 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -4,6 +4,19 @@ All notable changes to this project will be documented in this file.
 The format is based on [Keep a Changelog](https://keepachangelog.com/en/1.0.0/),
 and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html).
 
+## [0.21.2] - 2023-08-07
+
+### Fixed
+* Bugs (#1068, #774) in atomic addition caused due to imporper floating point atomic emulation. (#1103)
+
+### Changed
+* Updated documentation and user guides (#1097, #879)
+
+### Removed
+* Dependency on `spirv-tools` (#1103, #1108)
+* floating point atomic add emulation using `atomic_ops.cl` (#1103)
+* `NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE` configuration option (#1103)
+
 ## [0.21.1] - 2023-07-17
 
 ### Changed
diff --git a/Dockerfile b/Dockerfile
index 0e90599329..1bfa3f78ed 100644
--- a/Dockerfile
+++ b/Dockerfile
@@ -199,7 +199,7 @@ RUN \
     --mount=type=bind,target=/opt/toolkit,source=/opt/toolkit,from=toolkit-dist \
     export http_proxy=$http_proxy https_proxy=$https_proxy \
     && apt-get update && apt-get install -y \
-    spirv-tools spirv-headers \
+    spirv-headers \
     rsync \
     && rm -rf /var/lib/apt/lists/* \
     && rsync -a /opt/toolkit/bin/ /usr/local/bin/ \
diff --git a/README.md b/README.md
index 20ec88adc6..0a5af67a80 100644
--- a/README.md
+++ b/README.md
@@ -39,7 +39,7 @@ to learn more.
 Numba-dpex is part of the Intel&reg; Distribution of Python (IDP) and Intel&reg;
 oneAPI AIKit, and can be installed along with oneAPI. Additionally, we support
 installing it from Anaconda cloud. Please refer the instructions
-on our [documentation page](https://intelpython.github.io/numba-dpex/latest/user_guides/getting_started.html)
+on our [documentation page](https://intelpython.github.io/numba-dpex/latest/getting_started.html)
 for more details.
 
 Once the package is installed, a good starting point is to run the examples in
diff --git a/conda-recipe/meta.yaml b/conda-recipe/meta.yaml
index 317557ae91..1a597f3cf8 100644
--- a/conda-recipe/meta.yaml
+++ b/conda-recipe/meta.yaml
@@ -26,7 +26,6 @@ requirements:
         - python
         - numba >=0.57*
         - dpctl >=0.14*
-        - spirv-tools
         - dpcpp-llvm-spirv
         - dpnp >=0.11*
         - packaging
diff --git a/docs/.gitignore b/docs/.gitignore
index 73f85073bb..5c490088af 100644
--- a/docs/.gitignore
+++ b/docs/.gitignore
@@ -1,1 +1,2 @@
 apidoc
+sources/_build
diff --git a/docs/CoreFeatures.rst b/docs/CoreFeatures.rst
deleted file mode 100644
index 96fcbec334..0000000000
--- a/docs/CoreFeatures.rst
+++ /dev/null
@@ -1,70 +0,0 @@
-.. _core_features:
-
-Code-generation based on a device
-=================================
-
-In numba-dpex, kernels are written in a device-agnostic fashion making it
-easy to write portable code. A kernel is compiled for the device on which the
-kernel is enqueued to be executed. The device is specified using a
-``dpctl.device_context`` context manager. In the following example, two versions
-of the ``sum`` kernel are compiled, one for a GPU and another for a CPU based on
-which context the function was invoked. Numba-dpex supports
-OpenCL CPU and GPU devices and Level Zero GPU devices. In future, compilation
-support may be extended to other type of SYCL devices that are supported by
-DPC++'s runtime.
-
-    .. code-block:: python
-
-        import numpy as np
-        import numba_dpex as dpex
-        import dpctl
-
-
-        @dpex.kernel
-        def sum(a, b, c):
-            i = dpex.get_global_id(0)
-            c[i] = a[i] + b[i]
-
-
-        a = np.array(np.random.random(20), dtype=np.float32)
-        b = np.array(np.random.random(20), dtype=np.float32)
-        c = np.ones_like(a)
-
-        with dpctl.device_context("level_zero:gpu"):
-            sum[20, dpex.DEFAULT_LOCAL_SIZE](a, b, c)
-
-        with dpctl.device_context("opencl:cpu"):
-            sum[20, dpex.DEFAULT_LOCAL_SIZE](a, b, c)
-
-Automatic offload of NumPy expressions
-======================================
-
-A key distinction between numba-dpex and other the GPU backends in Numba is
-the ability to automatically offload specific data-parallel sections of a
-Numba ``jit`` function.
-
-.. todo::
-
-    Details and examples to be added.
-
-Controllable Fallback
----------------------
-
-By default, if a section of code cannot be offloaded to the GPU, it is
-automatically executed on the CPU and warning is printed. This behavior is only
-applicable to ``jit`` functions, auto-offloading of NumPy calls, array
-expressions and ``prange`` loops. To disable this functionality and force code
-running on GPU set the environment variable ``NUMBA_DPEX_FALLBACK_ON_CPU`` to
-false (e.g. ``export NUMBA_DPEX_FALLBACK_ON_CPU=0``). In this case the code is
-not automatically offloaded to the CPU and errors occur if any.
-
-Offload Diagnostics
--------------------
-
-Setting the debug environment variable ``NUMBA_DPEX_OFFLOAD_DIAGNOSTICS`` (e.g.
-``export NUMBA_DPEX_OFFLOAD_DIAGNOSTICS=1``) provides emission of the parallel
-and offload diagnostics information based on produced parallel transforms. The
-level of detail depends on the integer value between 1 and 4 that is set to the
-environment variable (higher is more detailed). In the "Auto-offloading" section
-there is the information on which device (device name) this parfor or kernel was
-offloaded.
diff --git a/docs/LICENSE_THIRD_PARTY b/docs/LICENSE_THIRD_PARTY
new file mode 100644
index 0000000000..e69de29bb2
diff --git a/docs/Makefile b/docs/Makefile
index d4bb2cbb9e..1ed4090504 100644
--- a/docs/Makefile
+++ b/docs/Makefile
@@ -5,7 +5,7 @@
 # from the environment for the first two.
 SPHINXOPTS    ?=
 SPHINXBUILD   ?= sphinx-build
-SOURCEDIR     = .
+SOURCEDIR     = ./source
 BUILDDIR      = _build
 
 # Put it first so that "make" without argument is like "make help".
@@ -18,3 +18,6 @@ help:
 # "make mode" option.  $(O) is meant as a shortcut for $(SPHINXOPTS).
 %: Makefile
 	@$(SPHINXBUILD) -M $@ "$(SOURCEDIR)" "$(BUILDDIR)" $(SPHINXOPTS) $(O)
+
+clean:
+	rm -rf "$(BUILDDIR)"
diff --git a/docs/asset/images/DPEP-large.png b/docs/asset/images/DPEP-large.png
new file mode 100644
index 0000000000..0ccbb3d38f
Binary files /dev/null and b/docs/asset/images/DPEP-large.png differ
diff --git a/docs/asset/images/DPEP.png b/docs/asset/images/DPEP.png
new file mode 100644
index 0000000000..a6e2610b94
Binary files /dev/null and b/docs/asset/images/DPEP.png differ
diff --git a/docs/asset/images/dpep-all.png b/docs/asset/images/dpep-all.png
new file mode 100644
index 0000000000..ffa6a96dac
Binary files /dev/null and b/docs/asset/images/dpep-all.png differ
diff --git a/docs/asset/images/kernel-queue-device.png b/docs/asset/images/kernel-queue-device.png
new file mode 100644
index 0000000000..15482f58d4
Binary files /dev/null and b/docs/asset/images/kernel-queue-device.png differ
diff --git a/docs/asset/images/kernel_prog_model.png b/docs/asset/images/kernel_prog_model.png
new file mode 100644
index 0000000000..d9fda7f48a
Binary files /dev/null and b/docs/asset/images/kernel_prog_model.png differ
diff --git a/docs/asset/images/numba-dpex-logo.png b/docs/asset/images/numba-dpex-logo.png
new file mode 100644
index 0000000000..a0f224e883
Binary files /dev/null and b/docs/asset/images/numba-dpex-logo.png differ
diff --git a/docs/developer_guides/dpnp_integration.rst b/docs/backups/developer_guides/dpnp_integration.rst
similarity index 100%
rename from docs/developer_guides/dpnp_integration.rst
rename to docs/backups/developer_guides/dpnp_integration.rst
diff --git a/docs/developer_guides/tools.rst b/docs/backups/developer_guides/tools.rst
similarity index 100%
rename from docs/developer_guides/tools.rst
rename to docs/backups/developer_guides/tools.rst
diff --git a/docs/backups/docker.rst b/docs/backups/docker.rst
new file mode 100644
index 0000000000..b37f4428be
--- /dev/null
+++ b/docs/backups/docker.rst
@@ -0,0 +1,146 @@
+.. 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
new file mode 100644
index 0000000000..d96aefdd9a
--- /dev/null
+++ b/docs/backups/sdc-sphinx-theme/globaltoc.html
@@ -0,0 +1,2 @@
+<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
new file mode 100644
index 0000000000..9a6badaf37
--- /dev/null
+++ b/docs/backups/sdc-sphinx-theme/layout.html
@@ -0,0 +1,99 @@
+{% 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
new file mode 100644
index 0000000000..00752f4e12
--- /dev/null
+++ b/docs/backups/sdc-sphinx-theme/localtoc.html
@@ -0,0 +1,2 @@
+<h3>Page Contents</h3>
+{{ toc }}
diff --git a/docs/backups/sdc-sphinx-theme/searchbox.html b/docs/backups/sdc-sphinx-theme/searchbox.html
new file mode 100644
index 0000000000..0d17cc1491
--- /dev/null
+++ b/docs/backups/sdc-sphinx-theme/searchbox.html
@@ -0,0 +1,7 @@
+{%- 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
new file mode 100644
index 0000000000..8f12e23d6c
--- /dev/null
+++ b/docs/backups/sdc-sphinx-theme/static/copybutton.js
@@ -0,0 +1,62 @@
+$(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
new file mode 100644
index 0000000000..a0f224e883
Binary files /dev/null and b/docs/backups/sdc-sphinx-theme/static/numba-dpex-logo.png differ
diff --git a/docs/backups/sdc-sphinx-theme/static/sdc.css b/docs/backups/sdc-sphinx-theme/static/sdc.css
new file mode 100644
index 0000000000..b9e6e37205
--- /dev/null
+++ b/docs/backups/sdc-sphinx-theme/static/sdc.css
@@ -0,0 +1,694 @@
+/*!
+ * 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
new file mode 100644
index 0000000000..9fde61751d
Binary files /dev/null and b/docs/backups/sdc-sphinx-theme/static/sdc_email.png differ
diff --git a/docs/backups/sdc-sphinx-theme/static/sdc_examples.png b/docs/backups/sdc-sphinx-theme/static/sdc_examples.png
new file mode 100644
index 0000000000..dd7095b80e
Binary files /dev/null and b/docs/backups/sdc-sphinx-theme/static/sdc_examples.png differ
diff --git a/docs/backups/sdc-sphinx-theme/static/sdc_github.png b/docs/backups/sdc-sphinx-theme/static/sdc_github.png
new file mode 100644
index 0000000000..89467c51c4
Binary files /dev/null and b/docs/backups/sdc-sphinx-theme/static/sdc_github.png differ
diff --git a/docs/backups/sdc-sphinx-theme/static/sdc_issues.png b/docs/backups/sdc-sphinx-theme/static/sdc_issues.png
new file mode 100644
index 0000000000..13d597429e
Binary files /dev/null and b/docs/backups/sdc-sphinx-theme/static/sdc_issues.png differ
diff --git a/docs/backups/sdc-sphinx-theme/static/sidebar.js b/docs/backups/sdc-sphinx-theme/static/sidebar.js
new file mode 100644
index 0000000000..8922cda0ca
--- /dev/null
+++ b/docs/backups/sdc-sphinx-theme/static/sidebar.js
@@ -0,0 +1,160 @@
+﻿/*
+ * 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
new file mode 100644
index 0000000000..4184857de0
--- /dev/null
+++ b/docs/backups/sdc-sphinx-theme/theme.conf
@@ -0,0 +1,11 @@
+# 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/user_guides/getting_started.rst b/docs/backups/user_guides/getting_started.rst
similarity index 96%
rename from docs/user_guides/getting_started.rst
rename to docs/backups/user_guides/getting_started.rst
index 795f1c0396..50b4b69534 100644
--- a/docs/user_guides/getting_started.rst
+++ b/docs/backups/user_guides/getting_started.rst
@@ -11,7 +11,6 @@ Numba-dpex depends on following components:
 * dpnp 0.10.1 (`Intel Python DPNP`_)
 * `dpcpp-llvm-spirv`_ (SPIRV generation from LLVM IR)
 * `llvmdev`_ (LLVM IR generation)
-* `spirv-tools`_
 * `packaging`_
 * `scipy`_ (for testing)
 * `pytest`_ (for testing)
@@ -59,7 +58,7 @@ installed in conda environment:
 .. code-block:: bash
 
     export ONEAPI_ROOT=/opt/intel/oneapi
-    conda create -n numba-dpex-env -c ${ONEAPI_ROOT}/conda_channel python=3.7 dpctl dpnp numba spirv-tools dpcpp-llvm-spirv llvmdev pytest
+    conda create -n numba-dpex-env -c ${ONEAPI_ROOT}/conda_channel python=3.7 dpctl dpnp numba dpcpp-llvm-spirv llvmdev pytest
     conda activate numba-dpex-env
 
 Activate DPC++ compiler:
@@ -150,7 +149,6 @@ Refer to :ref:`Docker <docker>` section for more options.
 .. _`Intel Python dpnp`: https://github.com/IntelPython/dpnp
 .. _`dpcpp-llvm-spirv`: https://github.com/IntelPython/dpcpp-llvm-spirv
 .. _`llvmdev`: https://anaconda.org/intel/llvmdev
-.. _`spirv-tools`: https://anaconda.org/intel/spirv-tools
 .. _`packaging`: https://packaging.pypa.io/
 .. _`scipy`: https://anaconda.org/intel/scipy
 .. _`pytest`: https://docs.pytest.org
diff --git a/docs/user_guides/kernel_programming_guide/atomic-operations.rst b/docs/backups/user_guides/kernel_programming_guide/atomic-operations.rst
similarity index 60%
rename from docs/user_guides/kernel_programming_guide/atomic-operations.rst
rename to docs/backups/user_guides/kernel_programming_guide/atomic-operations.rst
index 95b5a060b4..3c9626e52c 100644
--- a/docs/user_guides/kernel_programming_guide/atomic-operations.rst
+++ b/docs/backups/user_guides/kernel_programming_guide/atomic-operations.rst
@@ -21,21 +21,6 @@ Example usage of atomic operations
     The ``numba_dpex.atomic.add`` function is analogous to The
     ``numba.cuda.atomic.add`` provided by the ``numba.cuda`` backend.
 
-Generating Native FP Atomics
-----------------------------
-Numba-dpex supports generating native floating-point atomics.
-This feature is experimental. Users will need to provide
-the following environment variables to activate it.
-
-    NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1
-
-Example command:
-
-.. code-block:: bash
-
-    NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1 \
-    python program.py
-
 Full examples
 -------------
 
diff --git a/docs/user_guides/kernel_programming_guide/device-functions.rst b/docs/backups/user_guides/kernel_programming_guide/device-functions.rst
similarity index 100%
rename from docs/user_guides/kernel_programming_guide/device-functions.rst
rename to docs/backups/user_guides/kernel_programming_guide/device-functions.rst
diff --git a/docs/user_guides/kernel_programming_guide/index.rst b/docs/backups/user_guides/kernel_programming_guide/index.rst
similarity index 100%
rename from docs/user_guides/kernel_programming_guide/index.rst
rename to docs/backups/user_guides/kernel_programming_guide/index.rst
diff --git a/docs/user_guides/kernel_programming_guide/memory-management.rst b/docs/backups/user_guides/kernel_programming_guide/memory-management.rst
similarity index 100%
rename from docs/user_guides/kernel_programming_guide/memory-management.rst
rename to docs/backups/user_guides/kernel_programming_guide/memory-management.rst
diff --git a/docs/user_guides/kernel_programming_guide/memory_allocation_address_space.rst b/docs/backups/user_guides/kernel_programming_guide/memory_allocation_address_space.rst
similarity index 100%
rename from docs/user_guides/kernel_programming_guide/memory_allocation_address_space.rst
rename to docs/backups/user_guides/kernel_programming_guide/memory_allocation_address_space.rst
diff --git a/docs/user_guides/kernel_programming_guide/random.rst b/docs/backups/user_guides/kernel_programming_guide/random.rst
similarity index 100%
rename from docs/user_guides/kernel_programming_guide/random.rst
rename to docs/backups/user_guides/kernel_programming_guide/random.rst
diff --git a/docs/user_guides/kernel_programming_guide/reduction.rst b/docs/backups/user_guides/kernel_programming_guide/reduction.rst
similarity index 100%
rename from docs/user_guides/kernel_programming_guide/reduction.rst
rename to docs/backups/user_guides/kernel_programming_guide/reduction.rst
diff --git a/docs/user_guides/kernel_programming_guide/selecting_device.rst b/docs/backups/user_guides/kernel_programming_guide/selecting_device.rst
similarity index 100%
rename from docs/user_guides/kernel_programming_guide/selecting_device.rst
rename to docs/backups/user_guides/kernel_programming_guide/selecting_device.rst
diff --git a/docs/user_guides/kernel_programming_guide/supported-python-features.rst b/docs/backups/user_guides/kernel_programming_guide/supported-python-features.rst
similarity index 100%
rename from docs/user_guides/kernel_programming_guide/supported-python-features.rst
rename to docs/backups/user_guides/kernel_programming_guide/supported-python-features.rst
diff --git a/docs/user_guides/kernel_programming_guide/synchronization.rst b/docs/backups/user_guides/kernel_programming_guide/synchronization.rst
similarity index 100%
rename from docs/user_guides/kernel_programming_guide/synchronization.rst
rename to docs/backups/user_guides/kernel_programming_guide/synchronization.rst
diff --git a/docs/user_guides/kernel_programming_guide/ufunc.rst b/docs/backups/user_guides/kernel_programming_guide/ufunc.rst
similarity index 100%
rename from docs/user_guides/kernel_programming_guide/ufunc.rst
rename to docs/backups/user_guides/kernel_programming_guide/ufunc.rst
diff --git a/docs/user_guides/kernel_programming_guide/writing_kernels.rst b/docs/backups/user_guides/kernel_programming_guide/writing_kernels.rst
similarity index 100%
rename from docs/user_guides/kernel_programming_guide/writing_kernels.rst
rename to docs/backups/user_guides/kernel_programming_guide/writing_kernels.rst
diff --git a/docs/backups/user_guides/memory-management.rst b/docs/backups/user_guides/memory-management.rst
new file mode 100644
index 0000000000..4a867ea6e5
--- /dev/null
+++ b/docs/backups/user_guides/memory-management.rst
@@ -0,0 +1,106 @@
+.. _memory_management:
+.. include:: ./../../ext_links.txt
+
+Memory Management
+=================
+
+Numba-dpex uses DPC++'s USM shared memory allocator (``memory_alloc``) to
+enable host to device and *vice versa* data transfer. By using USM shared
+memory allocator, numba-dpex allows seamless interoperability between
+numba-dpex and other SYCL-based Python extensions and across multiple
+kernels written using ``numba_dpex.kernel`` decorator.
+
+Numba-dpex uses the USM memory manager provided by dpctl and supports
+the **SYCL USM Array Interface** protocol to enable zero-copy data
+exchange across USM memory-backed Python objects.
+
+.. note::
+
+    USM pointers make sense within a SYCL context and can be of four allocation
+    types: ``host``, ``device``, ``shared``, or ``unknown``. Host applications,
+    including CPython interpreter, can work with USM pointers of type ``host``
+    and ``shared`` as if they were ordinary host pointers. Accessing ``device``
+    USM pointers by host applications is not allowed.
+
+SYCL USM Array Interface
+------------------------
+
+A SYCL library may allocate USM memory for the result that needs to be passed to
+Python. A native Python extension that makes use of such a library may expose
+this memory as an instance of Python class that will implement memory management
+logic (ensures that memory is freed when the instance is no longer needed).
+The need to manage memory arises whenever a library uses a custom allocator.
+For example, |dp4p|_ uses Python capsule to ensure that a native
+library-allocated memory is freed using the appropriate deallocator.
+
+To enable native extensions to pass the memory allocated by a native SYCL
+library to Numba, or another SYCL-aware Python extension without making a copy,
+the class must provide ``__sycl_usm_array_interface__`` attribute which
+returns a Python dictionary with the following fields:
+
+- ``shape``: tuple of ``int``
+- ``typestr``: ``string``
+- ``typedescr``: a list of tuples
+- ``data``: (``int``, ``bool``)
+- ``strides``: tuple of ``int``
+- ``offset``: ``int``
+- ``version``: ``int``
+- ``syclobj``: ``dpctl.SyclQueue`` or ``dpctl.SyclContext`` object
+
+The dictionary keys align with those of |npai|_ and |cai|_. For host accessible
+USM pointers, the object may also implement CPython
+`PEP-3118 <https://www.python.org/dev/peps/pep-3118/>`_
+compliant buffer interface which will be used if a ``data`` key is not present
+in the dictionary. Use of a buffer interface extends the interoperability to
+other Python objects, such as ``bytes``, ``bytearray``, ``array.array``, or
+``memoryview``. The type of the USM pointer stored in the object can be queried
+using methods of the dpctl.
+
+.. |npai| replace:: ``numpy.ndarray.__array_interface__``
+.. _npai: https://numpy.org/doc/stable/reference/arrays.interface.html
+
+.. |cai| replace:: ``__cuda_array_interface__``
+.. _cai: http://numba.pydata.org/numba-doc/latest/cuda/cuda_array_interface.html
+
+.. |dp4p| replace:: ``daal4py``
+.. _dp4p: https://intelpython.github.io/daal4py/
+
+
+Device-only memory and explicit data transfer
+---------------------------------------------
+
+At the moment, there is no mechanism for the explicit transfer of arrays to
+the device and back. Please use usm arrays.
+
+
+Local memory
+------------
+
+In SYCL's memory model, local memory is a contiguous region of memory allocated
+per work group and is visible to all the work items in that group. Local memory
+is device-only and cannot be accessed from the host. From the perspective offers
+the device, the local memory is exposed as a contiguous array of a specific
+types. The maximum available local memory is hardware-specific. The SYCL local
+memory concept is analogous to CUDA's shared memory concept.
+
+Numba-dpex provides a special function ``numba_dpex.local.array`` to
+allocate local memory for a kernel.
+
+.. literalinclude:: ./../../../../numba_dpex/examples/barrier.py
+   :pyobject: local_memory
+
+.. note::
+
+  To go convert from ``numba.cuda`` to ``numba-dpex``, replace
+  ``numba.cuda.shared.array`` with
+  ``numba_dpex.local.array(shape=blocksize, dtype=float32)``.
+
+.. todo::
+
+  Add details about current limitations for local memory allocation in
+  ``numba-dpex``.
+
+Private and Constant memory
+---------------------------
+
+``numba-dpex`` does not yet support SYCL private and constant memory.
diff --git a/docs/user_guides/migrating_from_numba_cuda.rst b/docs/backups/user_guides/migrating_from_numba_cuda.rst
similarity index 100%
rename from docs/user_guides/migrating_from_numba_cuda.rst
rename to docs/backups/user_guides/migrating_from_numba_cuda.rst
diff --git a/docs/backups/user_guides/performance_tips.rst b/docs/backups/user_guides/performance_tips.rst
new file mode 100644
index 0000000000..2e3132ee33
--- /dev/null
+++ b/docs/backups/user_guides/performance_tips.rst
@@ -0,0 +1,5 @@
+.. _performance_tips:
+.. include:: ./../ext_links.txt
+
+Performance Tips
+================
diff --git a/docs/backups/user_guides/selecting_device.rst b/docs/backups/user_guides/selecting_device.rst
new file mode 100644
index 0000000000..d187cc7d9a
--- /dev/null
+++ b/docs/backups/user_guides/selecting_device.rst
@@ -0,0 +1,62 @@
+Defining the execution queue for a kernel function
+==================================================
+There are two ways to specify the queue where a kernel is executed. The first
+way follows the notion of "compute follows data" (CFD). The second way is for a
+programmer to specify the execution queue using a dpctl.device_context context
+manager.
+
+
+In the CFD style of programming kernels, the execution queue is determined based
+on the input arguments passed to a kernel function. Currently, numba-dpex's
+kernel API only supports array arguments that provide the
+:code:`__sycl_usm_array_interface__` (SUAI) attribute for CFD style programming.
+The SUAI attribute encodes the queue where the array was defined.
+
+
+We also allow passing arrays and data types that do not provide SUAI. For such
+cases, programmers need to specify the queue using the
+:code:`dpctl.device_context` context manager. Do note that the use of
+:code:`dpctl.device_context` is deprecated and slotted for removal in some
+future release.
+
+
+
+**Users are not allowed to pass mixed type of arrays to a ``numba_dpex.kernel``
+.** For example, if the first array argument to a ``numba_dpex.kernel`` is of
+type :code:`numpy.ndarray`, the rest of the array argument will also have to be
+of type :code:`numpy.ndarray`.
+
+The following are how users can specify in which device they want to offload their computation.
+
+- :code:`numpy.ndarray`
+    Using context manager, :code:`with dpctl.device_context(SYCL_device)`. Please look at method :code:`select_device_ndarray()` in the example below.
+
+- Array with :code:`__sycl_usm_array_interface__` attribute
+     Numba-dpex supports the Compute Follows Data semantics in this case.
+     Compute Follows Data stipulates that computation must be off-loaded to
+     device where data is resident.
+
+     Expected behavior in different cases:
+            - Users are allowed to mix arrays created using equivalent SYCL
+              queues. Where equivalent queues are defined as:
+
+                Two SYCL queues are equivalent if they have the same:
+                    1. SYCL context
+                    2. SYCL device
+                    3. Same queue properties
+
+
+            - All usm-types are accessible from device. Users can mix arrays
+              with different usm-type as long as they were allocated using the
+              equivalent SYCL queue.
+
+
+            - Using the context_manager to specify a queue when passing SUAI
+              args will have no effect on queue selection and Numba-dpex will
+              print out a warning to inform the user.
+
+
+Example
+~~~~~~~
+
+.. literalinclude:: ./../../../../numba_dpex/tests/kernel_tests/test_sycl_usm_array_iface_interop.py
diff --git a/docs/backups/user_guides/troubleshooting.rst b/docs/backups/user_guides/troubleshooting.rst
new file mode 100644
index 0000000000..4b813015e4
--- /dev/null
+++ b/docs/backups/user_guides/troubleshooting.rst
@@ -0,0 +1,5 @@
+.. _troubleshooting:
+.. include:: ./../ext_links.txt
+
+Troubleshooting
+===============
diff --git a/docs/conf.py b/docs/conf.py
deleted file mode 100644
index 6ee87b4ee5..0000000000
--- a/docs/conf.py
+++ /dev/null
@@ -1,80 +0,0 @@
-# Configuration file for the Sphinx documentation builder.
-#
-# This file only contains a selection of the most common options. For a full
-# list see the documentation:
-# https://www.sphinx-doc.org/en/master/usage/configuration.html
-
-# -- Path setup --------------------------------------------------------------
-
-# If extensions (or modules to document with autodoc) are in another directory,
-# add these directories to sys.path here. If the directory is relative to the
-# documentation root, use os.path.abspath to make it absolute, like shown here.
-#
-# import os
-# import sys
-# sys.path.insert(0, os.path.abspath('.'))
-
-import sphinx_rtd_theme
-
-import numba_dpex
-
-# -- Project information -----------------------------------------------------
-
-project = "numba-dpex"
-copyright = "2021-2023 Intel Corporation"
-author = "Intel"
-
-# The full version, including alpha/beta/rc tags
-release = numba_dpex.__version__
-
-
-# -- General configuration ---------------------------------------------------
-
-# Add any Sphinx extension module names here, as strings. They can be
-# extensions coming with Sphinx (named 'sphinx.ext.*') or your custom
-# ones.
-extensions = [
-    "recommonmark",
-    "sphinx_rtd_theme",
-    "sphinx.ext.autodoc",
-    "sphinx.ext.todo",
-    "sphinx.ext.intersphinx",
-    "sphinxcontrib.apidoc",
-    "sphinxcontrib.googleanalytics",
-]
-
-googleanalytics_id = "G-LGGL0NJK6P"
-googleanalytics_enabled = True
-
-todo_include_todos = True
-
-source_parsers = {".md": "recommonmark.parser.CommonMarkParser"}
-
-# Add any paths that contain templates here, relative to this directory.
-templates_path = ["_templates"]
-
-# List of patterns, relative to source directory, that match files and
-# directories to ignore when looking for source files.
-# This pattern also affects html_static_path and html_extra_path.
-exclude_patterns = ["_build", "Thumbs.db", ".DS_Store"]
-
-
-# -- Options for HTML output -------------------------------------------------
-
-# The theme to use for HTML and HTML Help pages.  See the documentation for
-# a list of builtin themes.
-#
-html_theme = "sphinx_rtd_theme"
-
-# Add any paths that contain custom static files (such as style sheets) here,
-# relative to this directory. They are copied after the builtin static files,
-# so a file named "default.css" will overwrite the builtin "default.css".
-# html_static_path = ["_static"]
-
-modindex_common_prefix = ["numba_dpex."]
-
-# See https://github.com/sphinx-contrib/apidoc
-apidoc_module_dir = "../numba_dpex"
-apidoc_output_dir = "apidoc"
-apidoc_excluded_paths = ["tests"]
-apidoc_separate_modules = True
diff --git a/docs/developer_guides/caching.rst b/docs/developer_guides/caching.rst
deleted file mode 100644
index 49186ca472..0000000000
--- a/docs/developer_guides/caching.rst
+++ /dev/null
@@ -1,24 +0,0 @@
-.. _caching:
-
-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, we do not perform file-based caching, instead we use an Least Recently Used (LRU) caching mechanism. However when a kernel needs to be evicted, we utilize numba's file-based caching mechanism 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 whenever 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.
-
-Settings
-========
-
-Therefore, we employ similar environment variables as used in Numba, i.e. ``NUMBA_CACHE_DIR`` etc. However we add three more environment variables to control the caching mechanism.
-
-- In order to specify cache capacity, one can use ``NUMBA_DPEX_CACHE_SIZE``. 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, one can set ``NUMBA_DPEX_DEBUG_CACHE`` to 1. All environment variables are defined in :file:`numba_dpex/config.py`.
diff --git a/docs/index.rst b/docs/index.rst
deleted file mode 100644
index 9c3c4329aa..0000000000
--- a/docs/index.rst
+++ /dev/null
@@ -1,133 +0,0 @@
-Welcome to numba-dpex's documentation!
-======================================
-
-Numba data-parallel extension (`numba-dpex
-<https://github.com/IntelPython/numba-dpex>`_) is an Intel |reg|-developed
-extension to the `Numba <https://numba.pydata.org/>`_ JIT compiler. The
-extension adds kernel programming and automatic offload capabilities to the
-Numba compiler. Numba-dpex is part of `Intel oneAPI Base Toolkit
-<https://www.intel.com/content/www/us/en/developer/tools/oneapi/base-toolkit.html>`_
-and distributed with the `Intel Distribution for Python*
-<https://www.intel.com/content/www/us/en/developer/tools/oneapi/distribution-for-python.html>`_.
-The goal of the extension is to make it easy for Python programmers to
-write efficient and portable code for a mix of architectures across CPUs, GPUs,
-FPGAs and other accelerators.
-
-Numba-dpex provides an API to write data-parallel kernels directly in Python and
-compiles the kernels to a lower-level kernels that are executed using a `SYCL
-<https://www.khronos.org/sycl/>`_ runtime library. Presently, only Intel's
-`DPC++ <https://github.com/intel/llvm/blob/sycl/sycl/doc/GetStartedGuide.md>`_
-SYCL runtime is supported via the `dpctl
-<https://github.com/IntelPython/dpctl>`_ package, and only OpenCL and Level Zero
-devices are supported. Support for other SYCL runtime libraries and hardwares
-may be added in the future.
-
-Along with the kernel programming API an auto-offload feature is also provided.
-The feature enables automatic generation of kernels from data-parallel NumPy
-library calls and array expressions, Numba ``prange`` loops, and `other
-"data-parallel by construction" expressions
-<https://numba.pydata.org/numba-doc/latest/user/parallel.html>`_ that Numba is
-able to parallelize. Following two examples demonstrate the two ways in
-which kernels may be written using numba-dpex.
-
-  - Defining a data-parallel kernel explicitly.
-
-    .. code-block:: python
-
-        import numpy as np
-        import numba_dpex as dpex
-        import dpctl
-
-
-        @dpex.kernel
-        def sum(a, b, c):
-            i = dpex.get_global_id(0)
-            c[i] = a[i] + b[i]
-
-
-        a = np.array(np.random.random(20), dtype=np.float32)
-        b = np.array(np.random.random(20), dtype=np.float32)
-        c = np.ones_like(a)
-
-        with dpctl.device_context("opencl:gpu"):
-            sum[20, dpex.DEFAULT_LOCAL_SIZE](a, b, c)
-
-  - Writing implicitly data-parallel expressions in the fashion of `Numba
-    parallel loops
-    <https://numba.pydata.org/numba-doc/dev/user/parallel.html#explicit-parallel-loops>`_.
-
-    .. code-block:: python
-
-        from numba import njit
-        import numpy as np
-        import dpctl
-
-
-        @njit
-        def f1(a, b):
-            c = a + b
-            return c
-
-
-        global_size = 64
-        local_size = 32
-        N = global_size * local_size
-        a = np.ones(N, dtype=np.float32)
-        b = np.ones(N, dtype=np.float32)
-        with dpctl.device_context("opencl:gpu:0"):
-            c = f1(a, b)
-
-.. toctree::
-   :maxdepth: 1
-   :caption: Core Features
-
-   CoreFeatures
-
-.. toctree::
-   :maxdepth: 1
-   :caption: User Guides
-
-    Getting Started <user_guides/getting_started>
-    Direct kernel programming <user_guides/kernel_programming_guide/index>
-    Debugging with GDB <user_guides/debugging/index>
-    Docker <user_guides/docker>
-    numba-dpex for numba.cuda Programmers <user_guides/migrating_from_numba_cuda>
-
-.. toctree::
-    :maxdepth: 1
-    :caption: Developer Guides
-
-    developer_guides/dpnp_integration
-    developer_guides/tools
-
-
-Contributing
-============
-
-Refer the `contributing guide
-<https://github.com/IntelPython/numba-dpex/blob/main/CONTRIBUTING.md>`_ for
-information on coding style and standards used in numba-dpex.
-
-License
-=======
-
-Numba-dpex is Licensed under Apache License 2.0 that can be found in `LICENSE
-<https://github.com/IntelPython/numba-dpex/blob/main/LICENSE>`_. All usage and
-contributions to the project are subject to the terms and conditions of this
-license.
-
-
-Indices and tables
-==================
-
-.. only:: builder_html
-
-   * :ref:`genindex`
-   * :ref:`modindex`
-   * :ref:`search`
-
-.. only:: not builder_html
-
-   * :ref:`modindex`
-
-.. |reg|    unicode:: U+000AE .. REGISTERED SIGN
diff --git a/docs/make.bat b/docs/make.bat
index 2119f51099..6f7cf0f0d4 100644
--- a/docs/make.bat
+++ b/docs/make.bat
@@ -7,7 +7,7 @@ REM Command file for Sphinx documentation
 if "%SPHINXBUILD%" == "" (
 	set SPHINXBUILD=sphinx-build
 )
-set SOURCEDIR=.
+set SOURCEDIR=./source
 set BUILDDIR=_build
 
 if "%1" == "" goto help
@@ -31,5 +31,8 @@ goto end
 :help
 %SPHINXBUILD% -M help %SOURCEDIR% %BUILDDIR% %SPHINXOPTS% %O%
 
+:clean
+rd /s %BUILDDIR%
+
 :end
 popd
diff --git a/docs/source/api_reference/index.rst b/docs/source/api_reference/index.rst
new file mode 100644
index 0000000000..728caaabcf
--- /dev/null
+++ b/docs/source/api_reference/index.rst
@@ -0,0 +1,8 @@
+.. _index:
+.. include:: ./../ext_links.txt
+
+
+API Reference
+=============
+
+Coming soon
diff --git a/docs/source/conf.py b/docs/source/conf.py
new file mode 100644
index 0000000000..60d42b9b98
--- /dev/null
+++ b/docs/source/conf.py
@@ -0,0 +1,139 @@
+# *****************************************************************************
+# Copyright (c) 2022, Intel Corporation All rights reserved.
+#
+# 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.
+# *****************************************************************************
+
+# coding: utf-8
+# Configuration file for the Sphinx documentation builder.
+
+import numba_dpex
+
+# -- Project information -----------------------------------------------------
+
+project = "numba-dpex"
+copyright = "2020-2023, Intel Corporation"
+author = "Intel Corporation"
+
+# The full version, including alpha/beta/rc tags
+# release = "main"
+
+# -- General configuration ----------------------------------------------------
+
+# Add any Sphinx extension module names here, as strings. They can be
+# extensions coming with Sphinx (named 'sphinx.ext.*') or your custom
+# ones.
+extensions = [
+    "sphinx.ext.todo",
+    "sphinx.ext.intersphinx",
+    "sphinx.ext.extlinks",
+    "sphinx.ext.githubpages",
+    "sphinx.ext.napoleon",
+    "sphinx.ext.autosectionlabel",
+    "sphinxcontrib.programoutput",
+    "sphinxcontrib.googleanalytics",
+    "myst_parser",
+]
+
+# Add any paths that contain templates here, relative to this directory.
+# templates_path = ['_templates']
+templates_path = []
+
+# List of patterns, relative to source directory, that match files and
+# directories to ignore when looking for source files.
+# This pattern also affects html_static_path and html_extra_path.
+exclude_patterns = []
+
+
+# -- Options for HTML output -------------------------------------------------
+
+# 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_options = {
+    "icon_links": [
+        {
+            "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"
+googleanalytics_enabled = True
+
+# Add any paths that contain custom static files (such as style sheets) here,
+# relative to this directory. They are copied after the builtin static files,
+# so a file named "default.css" will overwrite the builtin "default.css".
+html_static_path = []
+
+html_sidebars = {
+    # "**": [
+    #     "globaltoc.html",
+    #     "sourcelink.html",
+    #     "searchbox.html",
+    #     "relations.html",
+    # ],
+}
+
+html_show_sourcelink = False
+
+# -- Todo extension configuration  ----------------------------------------------
+todo_include_todos = True
+todo_link_only = True
+
+# -- InterSphinx configuration: looks for objects in external projects -----
+# Add here external classes you want to link from Intel SDC documentation
+# Each entry of the dictionary has the following format:
+#      'class name': ('link to object.inv file for that class', None)
+# intersphinx_mapping = {
+#    'pandas': ('https://pandas.pydata.org/pandas-docs/stable/', None),
+#    'python': ('http://docs.python.org/2', None),
+#    'numpy': ('http://docs.scipy.org/doc/numpy', None)
+# }
+intersphinx_mapping = {}
+
+# -- Napoleon extension configuration (Numpy and Google docstring options) -------
+# napoleon_google_docstring = True
+# napoleon_numpy_docstring = True
+# napoleon_include_init_with_doc = True
+# napoleon_include_private_with_doc = True
+# napoleon_include_special_with_doc = True
+# napoleon_use_admonition_for_examples = False
+# napoleon_use_admonition_for_notes = False
+# napoleon_use_admonition_for_references = False
+# napoleon_use_ivar = False
+# napoleon_use_param = True
+# napoleon_use_rtype = True
+
+# -- Prepend module name to an object name or not -----------------------------------
+add_module_names = False
diff --git a/docs/source/contribution_guide.rst b/docs/source/contribution_guide.rst
new file mode 100644
index 0000000000..971a7fd411
--- /dev/null
+++ b/docs/source/contribution_guide.rst
@@ -0,0 +1,26 @@
+.. _contributions:
+.. include:: ./ext_links.txt
+
+Contribution Guide
+==================
+
+Classification of Contributions
+-------------------------------
+
+Development Cycle
+-----------------
+
+Issues and Pull Requests
+------------------------
+
+Coding Guidelines
+-----------------
+
+Unit Testing
+-------------
+
+Documentation
+-------------
+
+Tips for Developers
+-------------------
diff --git a/docs/source/examples.rst b/docs/source/examples.rst
new file mode 100644
index 0000000000..f05801f7ab
--- /dev/null
+++ b/docs/source/examples.rst
@@ -0,0 +1,20 @@
+.. _examples:
+.. include:: ./ext_links.txt
+
+List of examples
+================
+
+.. todo::
+   Provide list of examples for numba-dpex
+
+Benchmarks
+**********
+
+.. todo::
+   Provide instructions for dpbench
+
+Jupyter* Notebooks
+******************
+
+.. todo::
+   Provide instructions for Jupyter Notebook samples
diff --git a/docs/source/ext_links.txt b/docs/source/ext_links.txt
new file mode 100644
index 0000000000..13a02c5f0d
--- /dev/null
+++ b/docs/source/ext_links.txt
@@ -0,0 +1,27 @@
+..
+    **********************************************************
+    THESE ARE EXTERNAL PROJECT LINKS USED IN THE DOCUMENTATION
+    **********************************************************
+
+.. _NumPy*: https://numpy.org/
+.. _Numba*: https://numba.pydata.org/
+.. _numba-dpex: https://github.com/IntelPython/numba-dpex
+.. _Python* Array API Standard: https://data-apis.org/array-api/
+.. _OpenCl*: https://www.khronos.org/opencl/
+.. _oneAPI Level Zero: https://spec.oneapi.io/level-zero/latest/index.html
+.. _DPC++: https://www.apress.com/gp/book/9781484255735
+.. _Data Parallel Extension for Numba*: https://intelpython.github.io/numba-dpex/latest/index.html
+.. _SYCL*: https://www.khronos.org/sycl/
+.. _dpctl: https://intelpython.github.io/dpctl/latest/index.html
+.. _Data Parallel Control: https://intelpython.github.io/dpctl/latest/index.html
+.. _Dpnp: https://intelpython.github.io/dpnp/
+.. _dpnp: https://intelpython.github.io/dpnp/
+.. _Data Parallel Extension for Numpy*: https://intelpython.github.io/dpnp/
+.. _IEEE 754-2019 Standard for Floating-Point Arithmetic: https://standards.ieee.org/ieee/754/6210/
+.. _Intel oneAPI Base Toolkit: https://www.intel.com/content/www/us/en/developer/tools/oneapi/base-toolkit.html
+.. _Intel Distribution for Python*: https://www.intel.com/content/www/us/en/developer/tools/oneapi/distribution-for-python.html
+.. _Intel AI Analytics Toolkit: https://www.intel.com/content/www/us/en/developer/tools/oneapi/ai-analytics-toolkit.html
+.. _Data Parallel Extensions for Python*: https://intelpython.github.io/DPEP/main/
+.. _Intel VTune Profiler: https://www.intel.com/content/www/us/en/developer/tools/oneapi/vtune-profiler.html
+.. _Intel Advisor: https://www.intel.com/content/www/us/en/developer/tools/oneapi/advisor.html
+.. _oneMKL: https://www.intel.com/content/www/us/en/docs/oneapi/programming-guide/2023-2/intel-oneapi-math-kernel-library-onemkl.html
diff --git a/docs/source/getting_started.rst b/docs/source/getting_started.rst
new file mode 100644
index 0000000000..da928df207
--- /dev/null
+++ b/docs/source/getting_started.rst
@@ -0,0 +1,247 @@
+.. _getting_started:
+.. include:: ./ext_links.txt
+
+.. |copy| unicode:: U+000A9
+
+.. |trade| unicode:: U+2122
+
+Getting Started
+===============
+
+
+Installing pre-built conda packages
+-----------------------------------
+
+``numba-dpex`` along with its dependencies can be installed using ``conda``.
+It is recommended to use conda packages from the ``anaconda.org/intel`` channel
+to get the latest production releases.
+
+.. code-block:: bash
+
+    conda create -n numba-dpex-env                                             \
+        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
+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                     \
+        -c dppy/label/dev -c intel -c conda-forge
+
+
+
+Building from source
+--------------------
+
+``numba-dpex`` can be built from source using either ``conda-build`` or
+``setuptools``.
+
+Steps to build using ``conda-build``:
+
+1. Ensure ``conda-build`` is installed in the ``base`` conda environment or
+   create a new conda environment with ``conda-build`` installed.
+
+.. code-block:: bash
+
+    conda create -n build-env conda-build
+    conda activate build-env
+
+2. Build using the vendored conda recipe
+
+.. code-block:: bash
+
+    conda build conda-recipe -c intel -c conda-forge
+
+3. Install the conda package
+
+.. code-block:: bash
+
+    conda install -c local numba-dpex
+
+Steps to build using ``setup.py``:
+
+As before, a conda environment with all necessary dependencies is the suggested
+first step.
+
+.. code-block:: bash
+
+    # Create a conda environment that hass needed dependencies installed
+    conda create -n numba-dpex-env                                             \
+        dpctl dpnp numba dpcpp-llvm-spirv llvmdev pytest           \
+        -c intel -c conda-forge
+    # Activate the environment
+    conda activate numba-dpex-env
+    # Clone the numba-dpex repository
+    git clone https://github.com/IntelPython/numba-dpex.git
+    cd numba-dpex
+    python setup.py develop
+
+Building inside Docker
+----------------------
+
+A Dockerfile is provided on the GitHub repository to build ``numba-dpex``
+as well as its direct dependencies: ``dpctl`` and ``dpnp``. Users can either use
+one of the pre-built images on the ``numba-dpex`` GitHub page or use the
+bundled Dockerfile to build ``numba-dpex`` from source. Using the Dockerfile
+also ensures that all device drivers and runtime libraries are pre-installed.
+
+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::
+
+    When trying to build a docker image with Intel GPU support, the Dockerfile
+    will attempt to use the GitHub API to get the latest Intel GPU drivers.
+    Users may run into an issue related to  Github API call limits. The issue
+    can be bypassed by providing valid GitHub credentials using the
+    ``GITHUB_USER`` and ``GITHUB_PASSWORD``
+    `build args <https://docs.docker.com/engine/reference/commandline/build/#build-arg>`_
+    to increase the call limit. A GitHub
+    `access token <https://docs.github.com/en/authentication/keeping-your-account-and-data-secure/creating-a-personal-access-token>`_
+    can also be used instead of the password.
+
+.. note::
+
+    When building the docker image behind a firewall the proxy server settings
+    should be provided using the ``http_proxy`` and ``https_proxy`` build args.
+    These build args must be specified in lowercase.
+
+The bundled Dockerfile supports different python versions that can be specified
+via the ``PYTHON_VERSION`` build arg. By default, the docker image is based on
+the official python image based on slim debian. The requested python version
+must be from the available python docker images.
+
+The ``BASE_IMAGE`` build arg can be used to build the docker image from a
+custom image. Note that as the Dockerfile is based on debian any custom base
+image should be debian-based, like debian or ubuntu.
+
+The list of other build args are as follows. Please refer the Dockerfile to
+see currently all available build args.
+
+- ``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``
+
+
+Using the pre-built images
+~~~~~~~~~~~~~~~~~~~~~~~~~~
+
+There are several pre-built docker images available:
+
+- ``runtime`` package that provides a pre-built environment with ``numba-dpex``
+              already installed. It is ideal to quickly setup and try
+              ``numba-dpex``.
+
+.. code-block:: text
+
+    ghcr.io/intelpython/numba-dpex/runtime:<numba_dpex_version>-py<python_version>[-gpu]
+
+- ``builder`` package that has all required dependencies pre-installed and is
+              ideal for building ``numba-dpex`` from source.
+
+.. code-block:: text
+
+    ghcr.io/intelpython/numba-dpex/builder:<numba_dpex_version>-py<python_version>[-gpu]
+
+- ``stages`` package primarily meant for creating a new docker image that is
+             built on top of one of the pre-built images.
+
+After setting up the docker image, to run ``numba-dpex`` the following snippet
+can be used.
+
+.. code-block:: bash
+
+    docker run --rm -it ghcr.io/intelpython/numba-dpex/runtime:0.20.0-py3.10 bash
+
+It is advisable to verify the SYCL runtime and driver installation within the
+container by either running,
+
+.. code-block:: bash
+
+    sycl-ls
+
+or,
+
+.. code-block:: bash
+
+    python -m dpctl -f
+
+.. note::
+
+    To enable GPU device, the ``device`` argument should be used and one of the
+    ``*-gpu`` images should be used.
+
+    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.
+
+For example, to run ``numba-dpex`` 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
+
+
+
+Testing
+-------
+
+``numba-dpex`` uses pytest for unit testing and the following example
+shows a way to run the unit tests.
+
+.. code-block:: bash
+
+    python -m pytest --pyargs numba_dpex.tests
+
+Examples
+--------
+
+A set of examples on how to use ``numba-dpex`` can be found in
+``numba_dpex/examples``.
diff --git a/docs/source/glossary.rst b/docs/source/glossary.rst
new file mode 100644
index 0000000000..5c3414c7f8
--- /dev/null
+++ b/docs/source/glossary.rst
@@ -0,0 +1,5 @@
+.. _glossary:
+.. include:: ./ext_links.txt
+
+Glossary
+========
diff --git a/docs/source/index.rst b/docs/source/index.rst
new file mode 100644
index 0000000000..c8ce20c0c9
--- /dev/null
+++ b/docs/source/index.rst
@@ -0,0 +1,64 @@
+.. _index:
+.. include:: ./ext_links.txt
+
+.. Feedbacks:
+.. 	Use proper names (as provided by marketing)
+.. 	Do not use "we/I"
+.. 	Use correct acronym, expand acronym at the first instance
+
+
+.. Follow cupy docs.
+
+.. Overview:
+.. 	Main Features
+.. 	Project Goal
+
+.. Getting Started:
+.. 	Yevheni
+.. 	Docker page
+
+.. User Guide:
+.. 	https://intelpython.github.io/oneAPI-for-SciPy/details/programming_model/
+.. 	Kernel Programming API --> Kernel Programming Basics
+.. 	Compiling and Offloading DPNP
+.. 		- Introduce @dpjit
+.. 		- Array constructor in dpjit
+.. 		- ufunc (refer to numba)
+.. 		- prange
+.. 		-
+.. 	Debugging using GDB
+.. 	# Performance Tips
+.. 	# Troubleshooting
+.. 	Useful Links
+.. 	# To-Do
+
+.. API Reference:
+
+
+Data Parallel Extension for Numba*
+==================================
+
+.. module:: numba_dpex
+
+.. toctree::
+   :maxdepth: 2
+
+   overview
+   getting_started
+   user_guide/index
+   api_reference/index
+
+.. toctree::
+   :maxdepth: 2
+   :caption: Development
+
+   contribution_guide
+
+.. toctree::
+   :maxdepth: 2
+   :caption: Misc Notes
+
+   examples
+   glossary
+   license
+   release-notes
diff --git a/docs/source/license.rst b/docs/source/license.rst
new file mode 100644
index 0000000000..8752d7d639
--- /dev/null
+++ b/docs/source/license.rst
@@ -0,0 +1,10 @@
+.. _license:
+.. include:: ./ext_links.txt
+
+License
+=======
+
+Numba-dpex is Licensed under Apache License 2.0 that can be found in `LICENSE
+<https://github.com/IntelPython/numba-dpex/blob/main/LICENSE>`_. All usage and
+contributions to the project are subject to the terms and conditions of this
+license.
diff --git a/docs/source/overview.rst b/docs/source/overview.rst
new file mode 100644
index 0000000000..25ad9c139b
--- /dev/null
+++ b/docs/source/overview.rst
@@ -0,0 +1,150 @@
+.. _overview:
+.. include:: ./ext_links.txt
+
+Overview
+========
+
+Data Parallel Extension for Numba* (`numba-dpex`_) is an extension to
+the `Numba*`_ Python JIT compiler adding an architecture-agnostic kernel
+programming API, and a new front-end to compile the Data Parallel Extension
+for Numpy* (`dpnp`_) library. The ``dpnp`` Python library is a data-parallel
+implementation of `NumPy*`_'s API using the `SYCL*`_ language.
+
+.. ``numba-dpex``'s support for ``dpnp`` compilation is a new way for Numba* users
+.. to write code in a NumPy-like API that is already supported by Numba*, while at
+.. the same time automatically running such code parallelly on various types of
+.. architecture.
+
+``numba-dpex`` is an open-source project and can be installed as part of `Intel
+AI Analytics Toolkit`_ or the `Intel Distribution for Python*`_. The package is
+also available on Anaconda cloud and as a Docker image on GitHub. Please refer
+the :doc:`getting_started` page to learn more.
+
+Main Features
+-------------
+
+Portable Kernel Programming
+~~~~~~~~~~~~~~~~~~~~~~~~~~~
+
+The ``numba-dpex`` kernel programming API has a design similar to Numba's
+``cuda.jit`` sub-module. The API is modeled after the `SYCL*`_ language and uses
+the `DPC++`_ SYCL runtime. Currently, compilation of kernels is supported for
+SPIR-V-based OpenCL and `oneAPI Level Zero`_ devices CPU and GPU devices. In the
+future, compilation support for other types of hardware that are supported by
+DPC++ will be added.
+
+The following example illustrates a vector addition kernel written with
+``numba-dpex`` kernel API.
+
+.. code-block:: python
+
+    import dpnp
+    import numba_dpex as dpex
+
+
+    @dpex.kernel
+    def vecadd_kernel(a, b, c):
+        i = dpex.get_global_id(0)
+        c[i] = a[i] + b[i]
+
+
+    a = dpnp.ones(1024, device="gpu")
+    b = dpnp.ones(1024, device="gpu")
+    c = dpnp.empty_like(a)
+
+    vecadd_kernel[dpex.Range(1024)](a, b, c)
+    print(c)
+
+In the above example, three arrays are allocated on a default ``gpu`` device
+using the ``dpnp`` library. The arrays are then passed as input arguments to the
+kernel function. The compilation target and the subsequent execution of the
+kernel is determined by the input arguments and follow the
+"compute-follows-data" programming model as specified in the `Python* Array API
+Standard`_. To change the execution target to a CPU, the device keyword needs to
+be changed to ``cpu`` when allocating the ``dpnp`` arrays. It is also possible
+to leave the ``device`` keyword undefined and let the ``dpnp`` library select a
+default device based on environment flag settings. Refer the
+:doc:`user_guide/kernel_programming/index` for further details.
+
+``dpjit`` decorator
+~~~~~~~~~~~~~~~~~~~
+
+The ``numba-dpex`` package provides a new decorator ``dpjit`` that extends
+Numba's ``njit`` decorator. The new decorator is equivalent to
+``numba.njit(parallel=True)``, but additionally supports compiling ``dpnp``
+functions, ``prange`` loops, and array expressions that use ``dpnp.ndarray``
+objects.
+
+Unlike Numba's NumPy parallelization that only supports CPUs, ``dpnp``
+expressions are first converted to data-parallel kernels and can then be
+`offloaded` to different types of devices. As ``dpnp`` implements the same API
+as NumPy*, an existing ``numba.njit`` decorated function that uses
+``numpy.ndarray`` may be refactored to use ``dpnp.ndarray`` and decorated with
+``dpjit``. Such a refactoring can allow the parallel regions to be offloaded
+to a supported GPU device, providing users an additional option to execute their
+code parallelly.
+
+The vector addition example depicted using the kernel API can also be
+expressed in several different ways using ``dpjit``.
+
+.. code-block:: python
+
+    import dpnp
+    import numba_dpex as dpex
+
+
+    @dpex.dpjit
+    def vecadd_v1(a, b):
+        return a + b
+
+
+    @dpex.dpjit
+    def vecadd_v2(a, b):
+        return dpnp.add(a, b)
+
+
+    @dpex.dpjit
+    def vecadd_v3(a, b):
+        c = dpnp.empty_like(a)
+        for i in prange(a.shape[0]):
+            c[i] = a[i] + b[i]
+        return c
+
+As with the kernel API example, a ``dpjit`` function if invoked with ``dpnp``
+input arguments follows the compute-follows-data programming model. Refer
+:doc:`user_manual/dpnp_offload/index` for further details.
+
+
+.. Project Goal
+.. ------------
+
+.. If C++ is not your language, you can skip writing data-parallel kernels in SYCL
+.. and directly write them in Python.
+
+.. Our package ``numba-dpex`` extends the Numba compiler to allow kernel creation
+.. directly in Python via a custom compute API
+
+
+.. Contributing
+.. ------------
+
+.. Refer the `contributing guide
+.. <https://github.com/IntelPython/numba-dpex/blob/main/CONTRIBUTING>`_ for
+.. information on coding style and standards used in ``numba-dpex``.
+
+.. License
+.. -------
+
+.. ``numba-dpex`` is Licensed under Apache License 2.0 that can be found in `LICENSE
+.. <https://github.com/IntelPython/numba-dpex/blob/main/LICENSE>`_. All usage and
+.. contributions to the project are subject to the terms and conditions of this
+.. license.
+
+
+.. Along with the kernel programming API an auto-offload feature is also provided.
+.. The feature enables automatic generation of kernels from data-parallel NumPy
+.. library calls and array expressions, Numba ``prange`` loops, and `other
+.. "data-parallel by construction" expressions
+.. <https://numba.pydata.org/numba-doc/latest/user/parallel.html>`_ that Numba is
+.. able to parallelize. Following two examples demonstrate the two ways in which
+.. kernels may be written using numba-dpex.
diff --git a/docs/source/release-notes.rst b/docs/source/release-notes.rst
new file mode 100644
index 0000000000..ed091cce62
--- /dev/null
+++ b/docs/source/release-notes.rst
@@ -0,0 +1,8 @@
+.. _release-notes:
+.. include:: ./ext_links.txt
+
+Release Notes
+=============
+
+.. include:: ../../CHANGELOG.md
+   :parser: myst_parser.sphinx_
diff --git a/docs/source/user_guide/caching.rst b/docs/source/user_guide/caching.rst
new file mode 100644
index 0000000000..59c9e48385
--- /dev/null
+++ b/docs/source/user_guide/caching.rst
@@ -0,0 +1,49 @@
+.. 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/config.rst b/docs/source/user_guide/config.rst
new file mode 100644
index 0000000000..a9c188d660
--- /dev/null
+++ b/docs/source/user_guide/config.rst
@@ -0,0 +1,62 @@
+.. include:: ./../ext_links.txt
+
+Configuration Options for ``numba-dpex``
+========================================
+
+``numba-dpex`` provides a set of environment variables and flags for configuring different aspects of the compilation, debugging and execution of programs. The configuration flags of ``numba-dpex`` are mostly inherited from those of Numba*. They are defined in :file:`numba_dpex/core/config.py`.
+
+.. note::
+    In order to enable/disable each of the configuration flags, a ``NUMBA_DPEX``
+    prefix needs to be appended before each variable. For example, in order to
+    turn ``SAVE_IR_FILES`` flag on, it needs to be passed as ``NUMBA_DPEX_SAVE_IR_FILES=1``
+
+    For example:
+
+    .. code-block:: bash
+
+        user@host:~/NUMBA_DPEX_SAVE_IR_FILES=1 python numba_dpex_program.py
+
+
+The list of available configuration flags are listed as follows:
+
+``SAVE_IR_FILES``:
+    A flag to save the Numba* intermediate representation (IR) files generated by the compiler. Set to ``0`` by default.
+
+``SPIRV_VAL``:
+    A flag to turn Numba*'s ``SPIRV-VALIDATION`` switch. Set to ``0`` by default.
+
+``OFFLOAD_DIAGNOSTICS``:
+    A flag to dump the offload diagnostics. Set to ``0`` by default.
+
+``NATIVE_FP_ATOMICS``:
+    A flag to activate the native floating point (FP) atomcis support for supported devices. Requires ``llvm-spirv`` supporting the FP atomics extension. Set to ``0`` by default.
+
+``DEBUG``:
+    A flag to emit the debug info, inherited from |numba.core.config.DEBUG|_.
+
+``DEBUGINFO_DEFAULT``:
+    The default value for the `debug` flag. Inherited from |numba.core.config.DEBUGINFO_DEFAULT|_.
+
+``DUMP_KERNEL_LLVM``:
+    A flag to emit LLVM assembly language format (``.ll``). Inherited from |numba.core.config.DUMP_OPTIMIZED|_.
+
+``ENABLE_CACHE``:
+    A flag to enable caching, set ``NUMBA_DPEX_ENABLE_CACHE=0`` to turn off. Set to ``1`` by default.
+
+``CACHE_SIZE``:
+    A flag to specify the default cache size. Set to ``20`` by default.
+
+``DEBUG_CACHE``:
+    A flag to enable debugging of cahcing mechanism, set ``1`` to turn it on.
+
+``STATIC_LOCAL_MEM_PASS``:
+    A flag to turn on the ``ConstantSizeStaticLocalMemoryPass`` in the kernel pipeline. The pass is turned off by default.
+
+
+.. |numba.core.config.DEBUG| replace:: ``numba.core.config.DEBUG``
+.. |numba.core.config.DEBUGINFO_DEFAULT| replace:: ``numba.core.config.DEBUGINFO_DEFAULT``
+.. |numba.core.config.DUMP_OPTIMIZED| replace:: ``numba.core.config.DUMP_OPTIMIZED``
+
+.. _`numba.core.config.DEBUG`: https://github.com/numba/numba/blob/main/numba/core/config.py#L202
+.. _`numba.core.config.DEBUGINFO_DEFAULT`: https://github.com/numba/numba/blob/main/numba/core/config.py#L488
+.. _`numba.core.config.DUMP_OPTIMIZED`: https://github.com/numba/numba/blob/main/numba/core/config.py#L301
diff --git a/docs/user_guides/debugging/altering.rst b/docs/source/user_guide/debugging/altering.rst
similarity index 91%
rename from docs/user_guides/debugging/altering.rst
rename to docs/source/user_guide/debugging/altering.rst
index 4b5789f4cb..c7813a3b49 100644
--- a/docs/user_guides/debugging/altering.rst
+++ b/docs/source/user_guide/debugging/altering.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Altering Execution
 ==================
 
@@ -22,7 +24,7 @@ Example
 
 Source code :file:`numba_dpex/examples/debug/side-by-side-2.py`:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/side-by-side-2.py
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/side-by-side-2.py
    :pyobject: common_loop_body
    :linenos:
    :lineno-match:
diff --git a/docs/user_guides/debugging/backtrace.rst b/docs/source/user_guide/debugging/backtrace.rst
similarity index 72%
rename from docs/user_guides/debugging/backtrace.rst
rename to docs/source/user_guide/debugging/backtrace.rst
index ed18cd722b..6935df4f51 100644
--- a/docs/user_guides/debugging/backtrace.rst
+++ b/docs/source/user_guide/debugging/backtrace.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Backtrace
 ==========
 
@@ -5,8 +7,8 @@ The ``backtrace`` command displays a summary of how your program got where it
 is. Consider the following example
 ``numba_dpex/examples/debug/simple_dpex_func.py``:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/simple_dpex_func.py
-    :lines: 15-
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/simple_dpex_func.py
+    :lines: 5-
     :linenos:
     :lineno-match:
 
@@ -19,13 +21,13 @@ numba_dpex.func.
 
 Example 1:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/backtrace_kernel
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/backtrace_kernel
     :language: shell-session
     :emphasize-lines: 8,9
 
 Example 2:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/backtrace
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/backtrace
     :language: shell-session
     :emphasize-lines: 8-10
 
diff --git a/docs/user_guides/debugging/breakpoints.rst b/docs/source/user_guide/debugging/breakpoints.rst
similarity index 63%
rename from docs/user_guides/debugging/breakpoints.rst
rename to docs/source/user_guide/debugging/breakpoints.rst
index 386fb57466..bb6b2504cf 100644
--- a/docs/user_guides/debugging/breakpoints.rst
+++ b/docs/source/user_guide/debugging/breakpoints.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Breakpoints
 ===========
 
@@ -9,10 +11,10 @@ your program should stop in the kernel. Define breakpoints by line numbers or
 function names.
 
 You have several ways to set breakpoints:
-  - break <function>
-  - break <filename>:<linenumber>
-  - break <filename>:<function>
-  - break … if <condition>
+  - ``break <function>``
+  - ``break <filename>:<linenumber>``
+  - ``break <filename>:<function>``
+  - ``break … if <condition>``
 
 See also:
   - `Breakpoints in GDB*`_.
@@ -22,8 +24,8 @@ See also:
 Consider the following numba-dpex kernel code (refer
 ``numba_dpex/examples/debug/simple_sum.py`` for full example):
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/simple_sum.py
-    :lines: 15-
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/simple_sum.py
+    :lines: 5-
     :linenos:
     :lineno-match:
 
@@ -32,7 +34,7 @@ Consider the following numba-dpex kernel code (refer
 
 The debugger output:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/break_func
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/break_func
     :language: shell-session
     :emphasize-lines: 3
 
@@ -41,7 +43,7 @@ The debugger output:
 
 The debugger output:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/break_line_number
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/break_line_number
     :language: shell-session
     :emphasize-lines: 3
 
@@ -50,7 +52,7 @@ The debugger output:
 
 The debugger output:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/break_file_func
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/break_file_func
     :language: shell-session
     :emphasize-lines: 3
 
@@ -59,7 +61,7 @@ The debugger output:
 
 The debugger output:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/break_conditional
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/break_conditional
     :language: shell-session
     :emphasize-lines: 3
 
@@ -69,13 +71,13 @@ Breakpoints with nested functions
 Consider numba-dpex kernel code. See the source file
 ``numba_dpex/examples/debug/simple_dpex_func.py``:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/simple_dpex_func.py
-    :lines: 15-
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/simple_dpex_func.py
+    :lines: 5-
     :linenos:
     :lineno-match:
 
 The debugger output:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/break_nested_func
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/break_nested_func
     :language: shell-session
     :emphasize-lines: 3
diff --git a/docs/user_guides/debugging/common_issues.rst b/docs/source/user_guide/debugging/common_issues.rst
similarity index 98%
rename from docs/user_guides/debugging/common_issues.rst
rename to docs/source/user_guide/debugging/common_issues.rst
index 0579554371..a214a9883d 100644
--- a/docs/user_guides/debugging/common_issues.rst
+++ b/docs/source/user_guide/debugging/common_issues.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Common issues and tips
 ======================
 
diff --git a/docs/user_guides/debugging/data.rst b/docs/source/user_guide/debugging/data.rst
similarity index 87%
rename from docs/user_guides/debugging/data.rst
rename to docs/source/user_guide/debugging/data.rst
index fcb27b6f05..0926a8020e 100644
--- a/docs/user_guides/debugging/data.rst
+++ b/docs/source/user_guide/debugging/data.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Examining Data
 ==============
 
@@ -10,7 +12,7 @@ See `GDB* documentation <https://www.sourceware.org/gdb/onlinedocs/gdb/Data.html
 
 To print the value of a variable, run the ``print <variable>`` command.
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/local_variables_0
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/local_variables_0
     :language: shell-session
     :lines: 67-72
     :emphasize-lines: 1-6
@@ -24,7 +26,7 @@ Example - Complex Data Types
 
 Source code :file:`numba_dpex/examples/debug/side-by-side-2.py`:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/side-by-side-2.py
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/side-by-side-2.py
    :pyobject: common_loop_body
    :linenos:
    :lineno-match:
diff --git a/docs/user_guides/debugging/debugging_environment.rst b/docs/source/user_guide/debugging/debugging_environment.rst
similarity index 91%
rename from docs/user_guides/debugging/debugging_environment.rst
rename to docs/source/user_guide/debugging/debugging_environment.rst
index 747cbec1cc..abedd0840d 100644
--- a/docs/user_guides/debugging/debugging_environment.rst
+++ b/docs/source/user_guide/debugging/debugging_environment.rst
@@ -1,4 +1,4 @@
-.. _debugging-environment:
+.. include:: ./../../ext_links.txt
 
 Configure debugging environment
 =================================
@@ -27,8 +27,8 @@ Configure debugging environment
 
     You can check the correctness of the work with the following example:
 
-    .. literalinclude:: ../../../numba_dpex/examples/debug/simple_sum.py
-        :lines: 15-
+    .. literalinclude:: ./../../../../numba_dpex/examples/debug/simple_sum.py
+        :lines: 5-
         :linenos:
         :lineno-match:
 
diff --git a/docs/user_guides/debugging/features.rst b/docs/source/user_guide/debugging/features.rst
similarity index 90%
rename from docs/user_guides/debugging/features.rst
rename to docs/source/user_guide/debugging/features.rst
index 80ac81c4f6..e6d95e6983 100644
--- a/docs/user_guides/debugging/features.rst
+++ b/docs/source/user_guide/debugging/features.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Supported Features
 ==================
 
diff --git a/docs/user_guides/debugging/frame_info.rst b/docs/source/user_guide/debugging/frame_info.rst
similarity index 79%
rename from docs/user_guides/debugging/frame_info.rst
rename to docs/source/user_guide/debugging/frame_info.rst
index 68f7cccdaa..bd2b9e1f76 100644
--- a/docs/user_guides/debugging/frame_info.rst
+++ b/docs/source/user_guide/debugging/frame_info.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Information About a Frame
 =========================
 
@@ -20,7 +22,7 @@ Example
 
 Source code :file:`numba_dpex/examples/debug/side-by-side.py`:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/side-by-side.py
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/side-by-side.py
    :pyobject: common_loop_body
    :linenos:
    :lineno-match:
@@ -59,21 +61,21 @@ Example
 
 Source code :file:`sum_local_vars.py`:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/sum_local_vars.py
-    :lines: 15-
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/sum_local_vars.py
+    :lines: 5-
     :linenos:
     :lineno-match:
 
 Run the debugger with ``NUMBA_OPT=0``:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/local_variables_0
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/local_variables_0
     :language: shell-session
     :lines: 1-6
 
 Use ``info locals``.
 Note that uninitialized variables are zeros:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/local_variables_0
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/local_variables_0
     :language: shell-session
     :lines: 8-48
     :emphasize-lines: 1-16, 24-39
diff --git a/docs/user_guides/debugging/index.rst b/docs/source/user_guide/debugging/index.rst
similarity index 88%
rename from docs/user_guides/debugging/index.rst
rename to docs/source/user_guide/debugging/index.rst
index 015a9b4903..04916265d6 100644
--- a/docs/user_guides/debugging/index.rst
+++ b/docs/source/user_guide/debugging/index.rst
@@ -1,3 +1,6 @@
+.. _index:
+.. include:: ./../../ext_links.txt
+
 Debugging with Intel® Distribution for GDB*
 ===========================================
 
@@ -39,15 +42,15 @@ Example of Intel® Distribution for GDB* usage
 
 You can use a sample kernel code, :file:`simple_sum.py`, for basic debugging:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/simple_sum.py
-    :lines: 15-
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/simple_sum.py
+    :lines: 5-
     :linenos:
     :lineno-match:
 
 Use the following commands to create a breakpoint inside the kernel and run the
 debugger:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/simple_sum
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/simple_sum
     :language: shell-session
 
 .. _debugging-features-and-limitations:
diff --git a/docs/user_guides/debugging/limitations.rst b/docs/source/user_guide/debugging/limitations.rst
similarity index 95%
rename from docs/user_guides/debugging/limitations.rst
rename to docs/source/user_guide/debugging/limitations.rst
index 5e5ba34a1b..e061148c9f 100644
--- a/docs/user_guides/debugging/limitations.rst
+++ b/docs/source/user_guide/debugging/limitations.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Limitations
 ===========
 
diff --git a/docs/user_guides/debugging/local_variables.rst b/docs/source/user_guide/debugging/local_variables.rst
similarity index 94%
rename from docs/user_guides/debugging/local_variables.rst
rename to docs/source/user_guide/debugging/local_variables.rst
index 77976c6fdb..6bb61d6637 100644
--- a/docs/user_guides/debugging/local_variables.rst
+++ b/docs/source/user_guide/debugging/local_variables.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Debugging Local Variables
 =========================
 
@@ -25,7 +27,7 @@ Example
 
 Source code :file:`numba_dpex/examples/debug/sum_local_vars.py`:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/sum_local_vars.py
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/sum_local_vars.py
     :pyobject: data_parallel_sum
     :linenos:
     :lineno-match:
@@ -97,7 +99,7 @@ Example 1 - Using ``NUMBA_EXTEND_VARIABLE_LIFETIMES``
 
 Source code :file:`numba_dpex/tests/debugging/test_info.py`:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/side-by-side.py
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/side-by-side.py
    :pyobject: common_loop_body
    :linenos:
    :lineno-match:
@@ -150,7 +152,7 @@ Example 2 - Using ``NUMBA_DUMP_ANNOTATION``
 
 Source code :file:`numba_dpex/examples/debug/sum_local_vars.py`:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/sum_local_vars.py
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/sum_local_vars.py
     :pyobject: data_parallel_sum
     :linenos:
     :lineno-match:
@@ -230,8 +232,8 @@ As a workaround you can expand lifetime of the variable by using it (i.e.
 passing to dummy function `revive()`) at the end of the function. So numba will
 not insert `del a` until the end of the function.
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/sum_local_vars_revive.py
-    :lines: 20-31
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/sum_local_vars_revive.py
+    :lines: 5-
     :linenos:
     :lineno-match:
 
@@ -350,27 +352,27 @@ Example 3 - Using ``info locals``
 
 Source code :file:`sum_local_vars.py`:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/sum_local_vars.py
-    :lines: 15-
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/sum_local_vars.py
+    :lines: 5-
     :linenos:
     :lineno-match:
 
 Run the debugger with ``NUMBA_OPT=0``:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/local_variables_0
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/local_variables_0
     :language: shell-session
     :lines: 1-6
 
 Run the ``info locals`` command. The sample output on "no optimization" level ``NUMBA_OPT=0`` is as follows:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/local_variables_0
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/local_variables_0
     :language: shell-session
     :lines: 8-48
     :emphasize-lines: 1-16, 24-39
 
 Since the debugger does not hit a line with the target variable ``l1``, the value equals 0. The true value of the variable ``l1`` is shown after stepping to line 22.
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/local_variables_0
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/local_variables_0
     :language: shell-session
     :lines: 49-66
     :emphasize-lines: 1-16
diff --git a/docs/user_guides/debugging/numba-0.55.rst b/docs/source/user_guide/debugging/numba-0.55.rst
similarity index 94%
rename from docs/user_guides/debugging/numba-0.55.rst
rename to docs/source/user_guide/debugging/numba-0.55.rst
index ae8f3aeb4e..846a37472f 100644
--- a/docs/user_guides/debugging/numba-0.55.rst
+++ b/docs/source/user_guide/debugging/numba-0.55.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Debugging Features in Numba 0.55
 ================================
 
@@ -28,7 +30,7 @@ Example
 
 Source code :file:`numba_dpex/examples/debug/side-by-side.py`:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/side-by-side.py
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/side-by-side.py
     :pyobject: common_loop_body
     :linenos:
     :lineno-match:
diff --git a/docs/user_guides/debugging/set_up_machine.rst b/docs/source/user_guide/debugging/set_up_machine.rst
similarity index 99%
rename from docs/user_guides/debugging/set_up_machine.rst
rename to docs/source/user_guide/debugging/set_up_machine.rst
index 015ea9ad9a..a42666859c 100644
--- a/docs/user_guides/debugging/set_up_machine.rst
+++ b/docs/source/user_guide/debugging/set_up_machine.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Set up the machine for debugging
 ====================================
 
diff --git a/docs/user_guides/debugging/stepping.rst b/docs/source/user_guide/debugging/stepping.rst
similarity index 72%
rename from docs/user_guides/debugging/stepping.rst
rename to docs/source/user_guide/debugging/stepping.rst
index b57748429e..a2b9127776 100644
--- a/docs/user_guides/debugging/stepping.rst
+++ b/docs/source/user_guide/debugging/stepping.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Stepping
 ========
 
@@ -8,15 +10,15 @@ Consider the following examples.
 
 ``numba_dpex/examples/debug/simple_sum.py``:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/simple_sum.py
-    :lines: 15-
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/simple_sum.py
+    :lines: 5-
     :linenos:
     :lineno-match:
 
 Example with a nested function ``numba_dpex/examples/debug/simple_dpex_func.py``:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/simple_dpex_func.py
-    :lines: 15-
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/simple_dpex_func.py
+    :lines: 5-
     :linenos:
     :lineno-match:
 
@@ -26,13 +28,13 @@ Example with a nested function ``numba_dpex/examples/debug/simple_dpex_func.py``
 
 Run the debugger and use the following commands:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/step_sum
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/step_sum
     :language: shell-session
     :emphasize-lines: 8-13
 
 You can use stepping to switch to a nested function. See the example below:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/step_dpex_func
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/step_dpex_func
     :language: shell-session
     :emphasize-lines: 8-14
 
@@ -41,7 +43,7 @@ You can use stepping to switch to a nested function. See the example below:
 
 The command allows you to move forward by machine instructions. The example uses an additional command ``x/i $pc``, which prints the instruction to be executed.
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/stepi
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/stepi
     :language: shell-session
     :emphasize-lines: 8-13
 
@@ -50,7 +52,7 @@ The command allows you to move forward by machine instructions. The example uses
 
 The command has stepping-like behavior, but it skips nested functions.
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/next
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/next
     :language: shell-session
     :emphasize-lines: 8-14
 
@@ -66,7 +68,7 @@ configure it for more efficient debugging. To ensure the current thread executes
 a single line without interference, set the scheduler-locking setting to `on` or
 `step`:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/sheduler_locking
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/sheduler_locking
     :language: shell-session
     :emphasize-lines: 8-13
 
diff --git a/docs/user_guides/debugging/symbols.rst b/docs/source/user_guide/debugging/symbols.rst
similarity index 84%
rename from docs/user_guides/debugging/symbols.rst
rename to docs/source/user_guide/debugging/symbols.rst
index 88f776db0a..0586ec9c7b 100644
--- a/docs/user_guides/debugging/symbols.rst
+++ b/docs/source/user_guide/debugging/symbols.rst
@@ -1,3 +1,5 @@
+.. include:: ./../../ext_links.txt
+
 Examining the Symbol Table
 ==========================
 
@@ -23,14 +25,14 @@ Example
 
 Source file ``numba_dpex/examples/debug/simple_sum.py``:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/simple_sum.py
-    :lines: 15-
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/simple_sum.py
+    :lines: 5-
     :linenos:
     :lineno-match:
 
 Output of the debug session:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/info_func
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/info_func
     :language: shell-session
     :emphasize-lines: 5-9
 
@@ -41,7 +43,7 @@ Output of the debug session:
 
 To print the type of a variable, run the ``ptype <variable>`` or ``whatis <variable>`` commands:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/commands/docs/local_variables_0
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/commands/docs/local_variables_0
     :language: shell-session
     :lines: 73-81
     :emphasize-lines: 1-6
@@ -51,7 +53,7 @@ Example - Complex Data Types
 
 Source code :file:`numba_dpex/examples/debug/side-by-side-2.py`:
 
-.. literalinclude:: ../../../numba_dpex/examples/debug/side-by-side-2.py
+.. literalinclude:: ./../../../../numba_dpex/examples/debug/side-by-side-2.py
    :pyobject: common_loop_body
    :linenos:
    :lineno-match:
diff --git a/docs/source/user_guide/dpnp_offload.rst b/docs/source/user_guide/dpnp_offload.rst
new file mode 100644
index 0000000000..d04183df20
--- /dev/null
+++ b/docs/source/user_guide/dpnp_offload.rst
@@ -0,0 +1,219 @@
+.. include:: ./../ext_links.txt
+
+Compiling and Offloading ``dpnp`` statements
+============================================
+
+Data Parallel Extension for NumPy* (``dpnp``) is a drop-in ``NumPy*``
+replacement library built on top of oneMKL and SYCL. ``numba-dpex`` allows
+various ``dpnp`` library function calls to be JIT-compiled using the
+``numba_dpex.dpjit`` decorator. Presently, ``numba-dpex`` can compile several
+``dpnp`` array constructors (``ones``, ``zeros``, ``full``, ``empty``), most
+universal functions, ``prange`` loops, and vector expressions using
+``dpnp.ndarray`` objects.
+
+An example of a supported usage of ``dpnp`` statements in ``numba-dpex`` is
+provided in the following code snippet:
+
+
+.. ``numba-dpex`` implements its own runtime library to support offloading ``dpnp``
+.. library functions to SYCL devices. For each ``dpnp`` function signature to be
+.. offloaded, ``numba-dpex`` implements the corresponding direct SYCL function call
+.. in the runtime and the function call is inlined in the generated LLVM IR.
+
+.. code-block:: python
+
+    import dpnp
+    from numba_dpex import dpjit
+
+
+    @dpjit
+    def foo():
+        a = dpnp.ones(1024, device="gpu")
+        return dpnp.sqrt(a)
+
+
+    a = foo()
+    print(a)
+    print(type(a))
+
+.. :samp:`dpnp.ones(10)` will be called through |ol_dpnp_ones(...)|_.
+
+
+.. Design
+.. -------
+
+.. ``numba_dpex`` uses the |numba.extending.overload| decorator to create a Numba*
+.. implementation of a function that can be used in `nopython mode`_ functions.
+.. This is done through translation of ``dpnp`` function signature so that they can
+.. be called in ``numba_dpex.dpjit`` decorated code.
+
+.. The specific SYCL operation for a certain ``dpnp`` function is performed by the
+.. runtime interface. During compiling a function decorated with the ``@dpjit``
+.. decorator, ``numba-dpex`` generates the corresponding SYCL function call through
+.. its runtime library and injects it into the LLVM IR through
+.. |numba.extending.intrinsic|_. The ``@intrinsic`` decorator is used for marking a
+.. ``dpnp`` function as typing and implementing the function in nopython mode using
+.. the `llvmlite IRBuilder API`_. This is an escape hatch to build custom LLVM IR
+.. that will be inlined into the caller.
+
+.. The code injection logic to enable ``dpnp`` functions calls in the Numba IR is
+.. implemented by :mod:`numba_dpex.core.dpnp_iface.arrayobj` module which replaces
+.. Numba*'s :mod:`numba.np.arrayobj`. Each ``dpnp`` function signature is provided
+.. with a concrete implementation to generates the actual code using Numba's
+.. ``overload`` function API. e.g.:
+
+.. .. code-block:: python
+
+..     @overload(dpnp.ones, prefer_literal=True)
+..     def ol_dpnp_ones(
+..         shape, dtype=None, order="C", device=None, usm_type="device", sycl_queue=None
+..     ):
+..         ...
+
+.. The corresponding intrinsic implementation is in :file:`numba_dpex/dpnp_iface/_intrinsic.py`.
+
+.. .. code-block:: python
+
+..    @intrinsic
+..    def impl_dpnp_ones(
+..        ty_context,
+..        ty_shape,
+..        ty_dtype,
+..        ty_order,
+..        ty_device,
+..        ty_usm_type,
+..        ty_sycl_queue,
+..        ty_retty_ref,
+..    ):
+..        ...
+
+Parallel Range
+---------------
+
+``numba-dpex`` supports using the ``numba.prange`` statements with
+``dpnp.ndarray`` objects. All such ``prange`` loops are offloaded as kernels and
+executed on a device inferred using the compute follows data programming model.
+The next examples shows using a ``prange`` loop.
+
+.. implements the ability to run loops in parallel, the language
+.. construct is adapted from Numba*'s ``prange`` concept that was initially
+.. designed to run OpenMP parallel for loops. In Numba*, the loop-body is scheduled
+.. in seperate threads, and they execute in a ``nopython`` Numba* context.
+.. ``prange`` automatically takes care of data privatization. ``numba-dpex``
+.. employs the ``prange`` compilation mechanism to offload parallel loop like
+.. programming constructs onto SYCL enabled devices.
+
+.. The ``prange`` compilation pass is delegated through Numba's
+.. :file:`numba/parfor/parfor_lowering.py` module where ``numba-dpex`` provides
+.. :file:`numba_dpex/core/parfors/parfor_lowerer.py` module to be used as the
+.. *lowering* mechanism through
+.. :py:class:`numba_dpex.core.parfors.parfor_lowerer.ParforLowerImpl` class. This
+.. provides a custom lowerer for ``prange`` nodes that generates a SYCL kernel for
+.. a ``prange`` node and submits it to a queue. Here is an example of a ``prange``
+.. use case in ``@dpjit`` context:
+
+.. code-block:: python
+
+    import dpnp
+    from numba_dpex import dpjit, prange
+
+
+    @dpjit
+    def foo():
+        x = dpnp.ones(1024, device="gpu")
+        o = dpnp.empty_like(a)
+        for i in prange(x.shape[0]):
+            o[i] = x[i] * x[i]
+        return o
+
+
+    c = foo()
+    print(c)
+    print(type(c))
+
+.. Each ``prange`` instruction in Numba* has an optional *lowerer* attribute. The
+.. lowerer attribute determines how the parfor instruction should be lowered to
+.. LLVM IR. In addition, the lower attribute decides which ``prange`` instructions
+.. can be fused together. At this point ``numba-dpex`` does not generate
+.. device-specific code and the lowerer used is same for all device types. However,
+.. a different :py:class:`numba_dpex.core.parfors.parfor_lowerer.ParforLowerImpl`
+.. instance is returned for every ``prange`` instruction for each corresponding CFD
+.. (Compute Follows Data) inferred device to prevent illegal ``prange`` fusion.
+
+``prange`` loop statements can also be used to write reduction loops as
+demonstrated by the following naive pairwise distance computation.
+
+.. code-block:: python
+
+  from numba_dpex import dpjit, prange
+  import dpnp
+  import dpctl
+
+
+  @dpjit
+  def pairwise_distance(X1, X2, D):
+      """Naïve pairwise distance impl - take an array representing M points in N
+      dimensions, and return the M x M matrix of Euclidean distances
+
+      Args:
+          X1 : Set of points
+          X2 : Set of points
+          D  : Outputted distance matrix
+      """
+      # Size of inputs
+      X1_rows = X1.shape[0]
+      X2_rows = X2.shape[0]
+      X1_cols = X1.shape[1]
+
+      float0 = X1.dtype.type(0.0)
+
+      # Outermost parallel loop over the matrix X1
+      for i in prange(X1_rows):
+          # Loop over the matrix X2
+          for j in range(X2_rows):
+              d = float0
+              # Compute exclidean distance
+              for k in range(X1_cols):
+                  tmp = X1[i, k] - X2[j, k]
+                  d += tmp * tmp
+              # Write computed distance to distance matrix
+              D[i, j] = dpnp.sqrt(d)
+
+
+  q = dpctl.SyclQueue()
+  X1 = dpnp.ones((10, 2), sycl_queue=q)
+  X2 = dpnp.zeros((10, 2), sycl_queue=q)
+  D = dpnp.empty((10, 2), sycl_queue=q)
+
+  pairwise_distance(X1, X2, D)
+  print(D)
+
+
+.. Fusion of Kernels
+.. ------------------
+
+.. ``numba-dpex`` can identify each NumPy* (or ``dpnp``) array expression as a
+.. data-parallel kernel and fuse them together to generate a single SYCL kernel.
+.. The kernel is automatically offloaded to the specified device where the fusion
+.. operation is invoked. Here is a simple example of a Black-Scholes formula
+.. computation where kernel fusion occurs at different ``dpnp`` math functions:
+
+.. .. literalinclude:: ./../../../numba_dpex/examples/blacksholes_njit.py
+..    :language: python
+..    :pyobject: blackscholes
+..    :caption: **EXAMPLE:** Data parallel kernel implementing the vector sum a+b
+..    :name: blackscholes_dpjit
+
+
+.. .. |numba.extending.overload| replace:: ``numba.extending.overload``
+.. .. |numba.extending.intrinsic| replace:: ``numba.extending.intrinsic``
+.. .. |ol_dpnp_ones(...)| replace:: ``ol_dpnp_ones(...)``
+.. .. |numba.np.arrayobj| replace:: ``numba.np.arrayobj``
+
+.. .. _low-level API: https://github.com/IntelPython/dpnp/tree/master/dpnp/backend
+.. .. _`ol_dpnp_ones(...)`: https://github.com/IntelPython/numba-dpex/blob/main/numba_dpex/dpnp_iface/arrayobj.py#L358
+.. .. _`numba.extending.overload`: https://numba.pydata.org/numba-doc/latest/extending/high-level.html#implementing-functions
+.. .. _`numba.extending.intrinsic`: https://numba.pydata.org/numba-doc/latest/extending/high-level.html#implementing-intrinsics
+.. .. _nopython mode: https://numba.pydata.org/numba-doc/latest/glossary.html#term-nopython-mode
+.. .. _`numba.np.arrayobj`: https://github.com/numba/numba/blob/main/numba/np/arrayobj.py
+.. .. _`llvmlite IRBuilder API`: http://llvmlite.pydata.org/en/latest/user-guide/ir/ir-builder.html
diff --git a/docs/source/user_guide/index.rst b/docs/source/user_guide/index.rst
new file mode 100644
index 0000000000..eb61b268b3
--- /dev/null
+++ b/docs/source/user_guide/index.rst
@@ -0,0 +1,21 @@
+.. _index:
+.. include:: ./../ext_links.txt
+
+.. .. image:: ./../../asset/images/DPEP-large.png
+..     :width: 400px
+..     :align: center
+..     :alt: Data Parallel Extensions for Python
+
+User Guide
+==========
+
+.. toctree::
+    :maxdepth: 2
+
+    programming_model
+    kernel_programming/index
+    dpnp_offload
+    debugging/index
+    caching
+    config
+    useful_links
diff --git a/docs/source/user_guide/kernel_programming/atomic-operations.rst b/docs/source/user_guide/kernel_programming/atomic-operations.rst
new file mode 100644
index 0000000000..07c2be4795
--- /dev/null
+++ b/docs/source/user_guide/kernel_programming/atomic-operations.rst
@@ -0,0 +1,27 @@
+Supported Atomic Operations
+===========================
+
+Numba-dpex supports some of the atomic operations supported in DPC++.
+Those that are presently implemented are as follows:
+
+.. automodule:: numba_dpex.ocl.stubs
+   :members: atomic
+   :noindex:
+
+Example
+-------
+
+Example usage of atomic operations
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/atomic_op.py
+   :pyobject: main
+
+.. note::
+
+    The ``numba_dpex.atomic.add`` function is analogous to The
+    ``numba.cuda.atomic.add`` provided by the ``numba.cuda`` backend.
+
+Full examples
+-------------
+
+- :file:`numba_dpex/examples/atomic_op.py`
diff --git a/docs/source/user_guide/kernel_programming/device-functions.rst b/docs/source/user_guide/kernel_programming/device-functions.rst
new file mode 100644
index 0000000000..0ad5a10441
--- /dev/null
+++ b/docs/source/user_guide/kernel_programming/device-functions.rst
@@ -0,0 +1,27 @@
+Writing Device Functions
+========================
+
+The user-level API of SYCL does not have a notion for device-only functions,
+*i.e.* functions that can be only invoked from a kernel and not from a host
+function. However, numba-dpex provides a special decorator
+``numba_dpex.func`` specifically to implement device functions.
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/device_func.py
+   :pyobject: a_device_function
+
+To use a device function from an another device function:
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/device_func.py
+   :pyobject: another_device_function
+
+To use a device function from a kernel function ``numba_dpex.kernel``:
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/device_func.py
+   :pyobject: a_kernel_function
+
+Unlike a kernel function, a device function can return a value like normal
+functions.
+
+.. todo::
+
+   Specific capabilities and limitations for device functions need to be added.
diff --git a/docs/source/user_guide/kernel_programming/index.rst b/docs/source/user_guide/kernel_programming/index.rst
new file mode 100644
index 0000000000..0008eccc6c
--- /dev/null
+++ b/docs/source/user_guide/kernel_programming/index.rst
@@ -0,0 +1,53 @@
+.. _index:
+.. include:: ./../../ext_links.txt
+
+Kernel Programming Basics
+=========================
+
+`Data Parallel Extensions for Python*`_ introduce a concept of an *offload kernel*, defined as
+a part of a Python program being submitted for execution to the device queue.
+
+.. image:: ./../../../asset/images/kernel-queue-device.png
+    :scale: 50%
+    :align: center
+    :alt: Offload Kernel
+
+There are multiple ways how to write offload kernels. CUDA*, OpenCl*, and SYCL* offer similar programming model
+known as the *data parallel kernel programming*. In this model you express the work in terms of *work items*.
+You split data into small pieces, and each piece will be a unit of work, or a *work item*. The total number of
+work items is called *global size*. You can also group work items into bigger chunks called *work groups*.
+The number of work items in the work group is called the *local size*.
+
+.. image:: ./../../../asset/images/kernel_prog_model.png
+    :scale: 50%
+    :align: center
+    :alt: Offload Kernel
+
+In this example there are 48 *work items* (8 in dimension 0, and 6 in dimension 1), that is the *global size* is 48.
+Work items are grouped in *work groups* with the *local size* 8 (4 in dimension 0, and 2 in dimension 1). There are
+total 48/8 = 6 work groups.
+
+In the *data parallel kernel programming* model you write a function that processes a given work item.
+Such a function is called the *data parallel kernel*.
+
+**Data Parallel Extension for Numba** offers a way to write data parallel kernels directly using Python using
+``numba_dpex.kernel``. It bears similarities with ``numba.cuda`` and ``numba.roc``, but unlike these proprietary
+programming models ``numba_dpex`` is built on top of `SYCL*`_ , which is hardware agnostic, meaning
+that with ``numba_dpex.kernel`` programming model you will be able to write a portable code targeting different
+hardware vendors.
+
+.. note::
+   The current version of ``numba-dpex`` supports Intel SYCL devices only
+
+.. toctree::
+   :caption: This document will cover the following chapters:
+   :maxdepth: 2
+
+   writing_kernels
+   synchronization
+   device-functions
+   atomic-operations
+   memory_allocation_address_space
+   reduction
+   ufunc
+   supported-python-features
diff --git a/docs/source/user_guide/kernel_programming/memory_allocation_address_space.rst b/docs/source/user_guide/kernel_programming/memory_allocation_address_space.rst
new file mode 100644
index 0000000000..7fc02e6a58
--- /dev/null
+++ b/docs/source/user_guide/kernel_programming/memory_allocation_address_space.rst
@@ -0,0 +1,36 @@
+Supported Address Space Qualifiers
+==================================
+
+The address space qualifier may be used to specify the region of memory that is
+used to allocate the object.
+
+Numba-dpex supports three disjoint named address spaces:
+
+1. Global Address Space
+    Global Address Space refers to memory objects allocated from the global
+    memory pool and will be shared among all work-items. Arguments passed to any
+    kernel are allocated in the global address space. In the below example,
+    arguments `a`, `b` and `c` will be allocated in the global address space:
+
+    .. literalinclude:: ./../../../../numba_dpex/examples/kernel/vector_sum.py
+
+
+2. Local Address Space
+    Local Address Space refers to memory objects that need to be allocated in
+    local memory pool and are shared by all work-items of a work-group.
+    Numba-dpex does not support passing arguments that are allocated in the
+    local address space to `@numba_dpex.kernel`. Users are allowed to allocate
+    static arrays in the local address space inside the `@numba_dpex.kernel`. In
+    the example below `numba_dpex.local.array(shape, dtype)` is the API used to
+    allocate a static array in the local address space:
+
+    .. literalinclude:: ./../../../../numba_dpex/examples/barrier.py
+      :lines: 54-87
+
+3. Private Address Space
+    Private Address Space refers to memory objects that are local to each
+    work-item and is not shared with any other work-item. In the example below
+    `numba_dpex.private.array(shape, dtype)` is the API used to allocate a
+    static array in the private address space:
+
+    .. literalinclude:: ./../../../../numba_dpex/examples/kernel/kernel_private_memory.py
diff --git a/docs/source/user_guide/kernel_programming/random.rst b/docs/source/user_guide/kernel_programming/random.rst
new file mode 100644
index 0000000000..e2aef1d14b
--- /dev/null
+++ b/docs/source/user_guide/kernel_programming/random.rst
@@ -0,0 +1,58 @@
+Random Number Generation
+========================
+
+Numba-dpex does not provide a random number generation algorithm that can be
+executed on the SYCL-supported Device.
+
+Numba-dpex provides access to NumPy random algorithms that can be executed on the
+SYCL-supported device via integration with `dpnp Random`_.
+
+.. _`dpnp Random`: https://intelpython.github.io/dpnp/reference/comparison.html#random-sampling
+
+
+Supported functions
+-------------------
+
+Simple random
+`````````````
+
+- `random <https://numpy.org/doc/1.16/reference/generated/numpy.random.random.html#numpy.random.random>`_
+- `sample <https://numpy.org/doc/1.16/reference/generated/numpy.random.sample.html#numpy.random.sample>`_
+- `ranf <https://numpy.org/doc/1.16/reference/generated/numpy.random.ranf.html#numpy.random.ranf>`_
+- `random_sample <https://numpy.org/doc/1.16/reference/generated/numpy.random.random_sample.html#numpy.random.random_sample>`_
+- `rand <https://numpy.org/doc/1.16/reference/generated/numpy.random.rand.html#numpy.random.rand>`_
+- `randint <https://numpy.org/doc/1.16/reference/generated/numpy.random.randint.html#numpy.random.randint>`_
+- `random_integers <https://numpy.org/doc/1.16/reference/generated/numpy.random.random_integers.html#numpy.random.random_integers>`_
+
+Distribution
+````````````
+
+- `beta <https://numpy.org/doc/1.16/reference/generated/numpy.random.beta.html#numpy.random.beta>`_
+- `binomial <https://numpy.org/doc/1.16/reference/generated/numpy.random.binomial.html#numpy.random.binomial>`_
+- `chisquare <https://numpy.org/doc/1.16/reference/generated/numpy.random.chisquare.html#numpy.random.chisquare>`_
+- `exponential <https://numpy.org/doc/1.16/reference/generated/numpy.random.exponential.html#numpy.random.exponential>`_
+- `gamma <https://numpy.org/doc/1.16/reference/generated/numpy.random.gamma.html#numpy.random.gamma>`_
+- `geometric <https://numpy.org/doc/1.16/reference/generated/numpy.random.geometric.html#numpy.random.geometric>`_
+- `gumbel <https://numpy.org/doc/1.16/reference/generated/numpy.random.gumbel.html#numpy.random.gumbel>`_
+- `hypergeometric <https://numpy.org/doc/1.16/reference/generated/numpy.random.hypergeometric.html#numpy.random.hypergeometric>`_
+- `laplace <https://numpy.org/doc/1.16/reference/generated/numpy.random.laplace.html#numpy.random.laplace>`_
+- `lognormal <https://numpy.org/doc/1.16/reference/generated/numpy.random.lognormal.html#numpy.random.lognormal>`_
+- `multinomial <https://numpy.org/doc/1.16/reference/generated/numpy.random.multinomial.html#numpy.random.multinomial>`_
+- `multivariate_normal <https://numpy.org/doc/1.16/reference/generated/numpy.random.multivariate_normal.html#numpy.random.multivariate_normal>`_
+- `negative_binomial <https://numpy.org/doc/1.16/reference/generated/numpy.random.negative_binomial.html#numpy.random.negative_binomial>`_
+- `normal <https://numpy.org/doc/1.16/reference/generated/numpy.random.normal.html#numpy.random.normal>`_
+- `poisson <https://numpy.org/doc/1.16/reference/generated/numpy.random.poisson.html#numpy.random.poisson>`_
+- `rayleigh <https://numpy.org/doc/1.16/reference/generated/numpy.random.rayleigh.html#numpy.random.rayleigh>`_
+- `standard_cauchy <https://numpy.org/doc/1.16/reference/generated/numpy.random.standard_cauchy.html#numpy.random.standard_cauchy>`_
+- `standard_exponential <https://numpy.org/doc/1.16/reference/generated/numpy.random.standard_exponential.html#numpy.random.standard_exponential>`_
+- `standard_gamma <https://numpy.org/doc/1.16/reference/generated/numpy.random.standard_gamma.html#numpy.random.standard_gamma>`_
+- `standard_normal <https://numpy.org/doc/1.16/reference/generated/numpy.random.standard_normal.html#numpy.random.standard_normal>`_
+- `uniform <https://numpy.org/doc/1.16/reference/generated/numpy.random.uniform.html#numpy.random.uniform>`_
+- `weibull <https://numpy.org/doc/1.16/reference/generated/numpy.random.weibull.html#numpy.random.weibull>`_
+
+Example:
+
+.. note::
+    To ensure the code is executed on GPU set ``DEBUG=1`` (or ``NUMBA_DPEX_DEBUG=1``) and look to stdout
+
+.. literalinclude:: ./../../../../numba_dpex/examples/rand.py
diff --git a/docs/source/user_guide/kernel_programming/reduction.rst b/docs/source/user_guide/kernel_programming/reduction.rst
new file mode 100644
index 0000000000..0f57b4b40e
--- /dev/null
+++ b/docs/source/user_guide/kernel_programming/reduction.rst
@@ -0,0 +1,51 @@
+Reduction on SYCL-supported Devices
+===================================
+
+Numba-dpex does not yet provide any specific decorator to implement
+reduction kernels. However, a kernel reduction can be written explicitly. This
+section provides two approaches for writing a reduction kernel as a
+``numba_dpex.kernel`` function.
+
+
+Example 1
+---------
+
+This example demonstrates a summation reduction on a one-dimensional array.
+
+Full example can be found at ``numba_dpex/examples/sum_reduction.py``.
+
+In this example, to reduce the array we invoke the kernel multiple times.
+
+.. literalinclude:: ./../../../../numba_dpex/examples/sum_reduction.py
+   :pyobject: sum_reduction_kernel
+
+.. literalinclude:: ./../../../../numba_dpex/examples/sum_reduction.py
+   :pyobject: sum_reduce
+
+Example 2
+---------
+
+Full example can be found at
+``numba_dpex/examples/sum_reduction_recursive_ocl.py``.
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/sum_reduction_recursive_ocl.py
+   :pyobject: sum_reduction_kernel
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/sum_reduction_recursive_ocl.py
+   :pyobject: sum_recursive_reduction
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/sum_reduction_recursive_ocl.py
+   :pyobject: sum_reduce
+
+.. note::
+
+    Numba-dpex does not yet provide any analogue to the ``numba.cuda.reduce``
+    decorator for writing reductions kernel. Such a decorator will be added in
+    future releases.
+
+Full examples
+-------------
+
+- :file:`numba_dpex/examples/sum_reduction_recursive_ocl.py`
+- :file:`numba_dpex/examples/sum_reduction_ocl.py`
+- :file:`numba_dpex/examples/sum_reduction.py`
diff --git a/docs/source/user_guide/kernel_programming/supported-python-features.rst b/docs/source/user_guide/kernel_programming/supported-python-features.rst
new file mode 100644
index 0000000000..fcfc41b904
--- /dev/null
+++ b/docs/source/user_guide/kernel_programming/supported-python-features.rst
@@ -0,0 +1,110 @@
+Supported Python Features inside ``numba_dpex.kernel``
+======================================================
+
+This page lists the Python features supported inside a ``numba_dpex.kernel``
+function.
+
+Built-in types
+--------------
+
+**Supported Types**
+
+- ``int``
+- ``float``
+
+**Unsupported Types**
+
+- ``complex``
+- ``bool``
+- ``None``
+- ``tuple``
+
+Built-in functions
+------------------
+
+The following built-in functions are supported:
+
+- ``abs()``
+- ``float``
+- ``int``
+- ``len()``
+- ``range()``
+- ``round()``
+
+Standard library modules
+------------------------
+
+The following functions from the math module are supported:
+
+-  ``math.acos()``
+-  ``math.asin()``
+-  ``math.atan()``
+-  ``math.acosh()``
+-  ``math.asinh()``
+-  ``math.atanh()``
+-  ``math.cos()``
+-  ``math.sin()``
+-  ``math.tan()``
+-  ``math.cosh()``
+-  ``math.sinh()``
+-  ``math.tanh()``
+-  ``math.erf()``
+-  ``math.erfc()``
+-  ``math.exp()``
+-  ``math.expm1()``
+-  ``math.fabs()``
+-  ``math.gamma()``
+-  ``math.lgamma()``
+-  ``math.log()``
+-  ``math.log10()``
+-  ``math.log1p()``
+-  ``math.sqrt()``
+-  ``math.ceil()``
+-  ``math.floor()``
+
+The following functions from the operator module are supported:
+
+-  ``operator.add()``
+-  ``operator.eq()``
+-  ``operator.floordiv()``
+-  ``operator.ge()``
+-  ``operator.gt()``
+-  ``operator.iadd()``
+-  ``operator.ifloordiv()``
+-  ``operator.imod()``
+-  ``operator.imul()``
+-  ``operator.ipow()``
+-  ``operator.isub()``
+-  ``operator.itruediv()``
+-  ``operator.le()``
+-  ``operator.lshift()``
+-  ``operator.lt()``
+-  ``operator.mod()``
+-  ``operator.mul()``
+-  ``operator.ne()``
+-  ``operator.neg()``
+-  ``operator.not_()``
+-  ``operator.or_()``
+-  ``operator.pos()``
+-  ``operator.pow()``
+-  ``operator.sub()``
+-  ``operator.truediv()``
+
+Unsupported Constructs
+----------------------
+
+The following Python constructs are **not supported**:
+
+- Exception handling (``try .. except``, ``try .. finally``)
+- Context management (the ``with`` statement)
+- Comprehensions (either list, dict, set or generator comprehensions)
+- Generator (any ``yield`` statements)
+- The ``raise`` statement
+- The ``assert`` statement
+
+
+NumPy support
+-------------
+
+NumPy functions are whole array operations and are not supported within a
+``numba_dpex.kernel``.
diff --git a/docs/source/user_guide/kernel_programming/synchronization.rst b/docs/source/user_guide/kernel_programming/synchronization.rst
new file mode 100644
index 0000000000..d590b0801b
--- /dev/null
+++ b/docs/source/user_guide/kernel_programming/synchronization.rst
@@ -0,0 +1,33 @@
+Synchronization Functions
+=========================
+
+``numba-dpex`` only supports some of the SYCL synchronization operations. For
+synchronization of all threads in the same thread block, numba-dpex provides
+a helper function called ``numba_dpex.barrier()``. This function implements the
+same pattern as barriers in traditional multi-threaded programming: invoking the
+function forces a thread to wait until all threads in the block reach the
+barrier, at which point it returns control to all its callers.
+
+``numba_dpex.barrier()`` supports two memory fence options:
+
+- ``numba_dpex.GLOBAL_MEM_FENCE``: The barrier function will queue a memory
+  fence to ensure correct ordering of memory operations to global memory. Using
+  the option can be useful when work-items, for example, write to buffer or
+  image objects and then want to read the updated data. Passing no arguments to
+  ``numba_dpex.barrier()`` is equivalent to setting the global memory fence
+  option.
+
+  .. .. literalinclude:: ./../../../../numba_dpex/examples/barrier.py
+  ..  :pyobject: no_arg_barrier_support
+
+- ``numba_dpex.LOCAL_MEM_FENCE``: The barrier function will either flush
+  any variables stored in local memory or queue a memory fence to ensure
+  correct ordering of memory operations to local memory.
+
+.. .. literalinclude:: ./../../../../numba_dpex/examples/barrier.py
+..    :pyobject: local_memory
+
+.. note::
+
+    The ``numba_dpex.barrier()`` function is semantically equivalent to
+    ``numba.cuda.syncthreads``.
diff --git a/docs/source/user_guide/kernel_programming/ufunc.rst b/docs/source/user_guide/kernel_programming/ufunc.rst
new file mode 100644
index 0000000000..eb1b47d99b
--- /dev/null
+++ b/docs/source/user_guide/kernel_programming/ufunc.rst
@@ -0,0 +1,52 @@
+Universal Functions
+===================
+
+Numba provides a `set of decorators
+<https://numba.pydata.org/numba-doc/latest/user/vectorize.html>`_ to create
+`NumPy universal functions
+<http://docs.scipy.org/doc/numpy/reference/ufuncs.html>`_-like routines that are
+JIT compiled. Although, a close analog to NumPy universal functions Numba's
+``@vectorize`` are not fully compatible with a regular NumPy ufunc.. Refer
+`Creating NumPy universal functions`_ for details.
+
+
+Numba-dpex only supports ``numba.vectorize`` decorator and not yet the
+``numba.guvectorize`` decorator. Another present limitation is that
+numba-dpex ufunc kernels cannot invoke a ``numba_dpex.kernel`` function.
+.. Ongoing work is in progress to address these limitations.
+
+Example 1: Basic Usage
+----------------------
+
+Full example can be found at :file:`numba_dpex/examples/vectorize.py`.
+
+.. literalinclude:: ./../../../../numba_dpex/examples/vectorize.py
+   :pyobject: ufunc_kernel
+
+.. literalinclude:: ./../../../../numba_dpex/examples/vectorize.py
+   :pyobject: test_njit
+
+Example 2: Calling ``numba.vectorize`` inside a ``numba_dpex.kernel``
+---------------------------------------------------------------------
+
+Full example can be found at :file:`numba_dpex/examples/blacksholes_njit.py`.
+
+.. literalinclude:: ./../../../../numba_dpex/examples/blacksholes_njit.py
+   :pyobject: cndf2
+
+.. note::
+
+    ``numba.cuda`` requires ``target='cuda'`` parameter for ``numba.vectorize``
+    and ``numba.guvectorize`` functions. Numba-dpex eschews the ``target``
+    parameter for ``@vectorize`` and infers the target from the
+    ``dpctl.device_context`` in which the ``numba.vectorize`` function is
+    called.
+
+Full Examples
+-------------
+
+- :file:`numba_dpex/examples/vectorize.py`
+- :file:`numba_dpex/examples/blacksholes_njit.py`
+
+.. _`Universal functions (ufunc)`: http://docs.scipy.org/doc/numpy/reference/ufuncs.html
+.. _`Creating NumPy universal functions`: https://numba.pydata.org/numba-doc/latest/user/vectorize.html
diff --git a/docs/source/user_guide/kernel_programming/writing_kernels.rst b/docs/source/user_guide/kernel_programming/writing_kernels.rst
new file mode 100644
index 0000000000..1e0deae459
--- /dev/null
+++ b/docs/source/user_guide/kernel_programming/writing_kernels.rst
@@ -0,0 +1,124 @@
+.. include:: ./../../ext_links.txt
+
+Writing Data Parallel Kernels
+=============================
+
+Kernel Declaration
+------------------
+A kernel function is a device function that is meant to be called from host
+code, where a device can be any SYCL supported device such as a GPU, CPU, or an
+FPGA. The main characteristics of a kernel function are:
+
+- **Scalars must be passed as an array**. Kernels operate with ``dpnp`` array
+  arguments only. If your want a scalar argument, then represent it as
+  0-dimensional ``dpnp`` array.
+
+.. note::
+   Please refer to `Data Parallel Extension for Numpy*`_ to learn more about ``dpnp``.
+
+- **Kernels cannot explicitly return a value**. All result data must be written
+  to ``dpnp`` array passed as a function's argument.
+
+Here is an example of a kernel that computes sum of two vectors ``a`` and ``b``.
+Arguments are two input vectors ``a`` and ``b`` and one output vector ``c`` for
+storing the result of vector summation:
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/vector_sum.py
+   :language: python
+   :lines: 8-9, 11-15
+   :caption: **EXAMPLE:** Data parallel kernel implementing the vector sum a+b
+   :name: ex_kernel_declaration_vector_sum
+
+
+.. Kernel Invocation
+.. ------------------
+
+.. When a kernel is launched you must specify the *global size* and the *local size*,
+.. which determine the hierarchy of threads, that is the order in which kernels
+.. will be invoked.
+
+.. The following syntax is used in ``numba-dpex`` for kernel invocation with
+.. specified global and local sizes:
+
+.. ``kernel_function_name[global_size, local_size](kernel arguments)``
+
+.. In the following example we invoke kernel ``kernel_vector_sum`` with global size
+.. specified via variable ``global_size``, and use ``numba_dpex.DEFAULT_LOCAL_SIZE``
+.. constant for setting local size to some default value:
+
+.. .. code-block:: python
+
+..    import numba_dpex as ndpx
+
+..    global_size = 10
+..    kernel_vector_sum[global_size, ndpx.DEFAULT_LOCAL_SIZE](a, b, c)
+
+.. .. note::
+..   Each kernel is compiled once, but it can be called multiple times with different global and local sizes settings.
+
+
+Kernel Invocation
+------------------
+
+The kernel launch parameter syntax for specifying global and local sizes are
+similar to ``SYCL``'s ``range`` and ``ndrange`` classes. The global and local
+sizes need to be specified with ``numba_dpex``'s ``Range`` and ``NdRange``
+classes.
+
+For example, below is a following kernel that computes a sum of two vectors:
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/vector_sum.py
+   :language: python
+   :lines: 8-9, 11-15
+   :caption: **EXAMPLE:** A vector sum kernel
+   :name: vector_sum_kernel
+
+If the ``global size`` parameter is needed to run, it could be like this (where
+``global_size`` is an ``int``):
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/vector_sum.py
+   :language: python
+   :lines: 8-9, 18-24
+   :emphasize-lines: 5
+   :caption: **EXAMPLE:** A vector sum kernel with a global size/range
+   :name: vector_sum_kernel_with_launch_param
+
+If both local and global ranges are needed, they can be specified using two
+instances of ``Range`` inside an ``NdRange`` object. For example, below is a
+kernel to compute pair-wise Euclidean distances of n-dimensional data points:
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/pairwise_distance.py
+   :language: python
+   :lines: 14-15, 36-51
+   :caption: **EXAMPLE:** A kernel to compute pair-wise Euclidean distances
+   :name: pairwise_distance_kernel
+
+Now the local and global sizes can be specified as follows (here both ``args.n``
+and ``args.l`` are ``int``):
+
+.. literalinclude:: ./../../../../numba_dpex/examples/kernel/pairwise_distance.py
+   :language: python
+   :lines: 14-15, 27-31, 54-67
+   :emphasize-lines: 4,6,13
+   :caption: **EXAMPLE:** A kernel to compute pair-wise Euclidean distances with
+               a global and a local size/range
+   :name: pairwise_distance_kernel_with_launch_param
+
+
+Kernel Indexing Functions
+-------------------------
+
+In *data parallel kernel programming* all work items are enumerated and accessed
+by their index. You will use ``numba_dpex.get_global_id()`` function to get the
+index of a current work item from the kernel. The total number of work items can
+be determined by calling ``numba_dpex.get_global_size()`` function.
+
+The work group size can be determined by calling ``numba_dpex.get_local_size()``
+function. Work items in the current work group are accessed by calling
+``numba_dpex.get_local_id()``.
+
+The total number of work groups are determined by calling
+``numba_dpex.get_num_groups()`` function. The current work group index is
+obtained by calling ``numba_dpex.get_group_id()`` function.
+
+.. _Black Scholes: https://en.wikipedia.org/wiki/Black%E2%80%93Scholes_model
diff --git a/docs/source/user_guide/programming_model.rst b/docs/source/user_guide/programming_model.rst
new file mode 100644
index 0000000000..dac9475bd1
--- /dev/null
+++ b/docs/source/user_guide/programming_model.rst
@@ -0,0 +1,279 @@
+.. _programming_model:
+.. include:: ./../ext_links.txt
+
+Programming Model
+=================
+
+In a heterogeneous system there may be **multiple** devices a Python user may
+want to engage. For example, it is common for a consumer-grade laptop to feature
+an integrated or a discrete GPU alongside a CPU.
+
+To harness their power one needs to know how to answer the following 3 key
+questions:
+
+1. How does a Python program recognize available computational devices?
+2. How does a Python workload specify computations to be offloaded to selected
+   devices?
+3. How does a Python application manage data sharing?
+
+Recognizing available devices
+~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
+
+Python package ``dpctl`` answers these questions. All the computational devices
+known to the underlying DPC++ runtime can be accessed using
+``dpctl.get_devices()``. A specific device of interest `can be selected
+<https://intelpython.github.io/dpctl/latest/docfiles/user_guides/manual/dpctl/device_selection.html>`__
+either using a helper function, e.g. ``dpctl.select_gpu_device()``, or by
+passing a filter selector string to ``dpctl.SyclDevice`` constructor.
+
+.. code:: python
+
+   import dpctl
+
+   # select a GPU device. If multiple devices present,
+   # let the underlying runtime select from GPUs
+   dev_gpu = dpctl.SyclDevice("gpu")
+   # select a CPU device
+   dev_cpu = dpctl.SyclDevice("cpu")
+
+   # stand-alone function, equivalent to C++
+   #   `auto dev = sycl::gpu_selector().select_device();`
+   dev_gpu_alt = dpctl.select_gpu_device()
+   # stand-alone function, equivalent to C++
+   #   `auto dev = sycl::cpu_selector().select_device();`
+   dev_cpu_alt = dpctl.select_cpu_device()
+
+A `device object
+<https://intelpython.github.io/dpctl/latest/docfiles/user_guides/manual/dpctl/devices.html>`__
+can be used to query properies of the device, such as its name, vendor, maximal
+number of computational units, memory size, etc.
+
+Specifying offload target
+~~~~~~~~~~~~~~~~~~~~~~~~~
+
+To answer the second question on the list we need a digression to explain
+offloading in oneAPI DPC++ first.
+
+.. note::
+   In DPC++, a computation kernel can be specified using generic C++
+   programming and then the kernel can be offloaded to any device that is
+   supported by an underlying SYCL runtime. The device to which the kernel
+   is offloaded is specified using an **execution queue** when *launching
+   the kernel*.
+
+   The oneAPI unified programming model brings portability across heterogeneous
+   architectures. Another important aspect of the programming model is its
+   inherent flexibility that makes it possible to go beyond portability and even
+   strive for performance portability. An oneAPI library may be implemented
+   using C++ techniques such as template metaprogramming or dynamic polymorphism
+   to implement specializations for a generic kernel. If a kernel is implemented
+   polymorphically, the specialized implementation will be dispatched based on
+   the execution queue specified during kernel launch. The oneMKL library is an
+   example of a performance portable oneAPI library.
+
+A computational task is offloaded for execution on a device by submitting it to
+DPC++ runtime which inserts the task in a computational graph. Once the device
+becomes available the runtime selects a task whose dependencies are met for
+execution. The computational graph as well as the device targeted by its tasks
+are stored in a `SYCL queue
+<https://intelpython.github.io/dpctl/latest/docfiles/user_guides/manual/dpctl/queues.html>`__
+object. The task submission is therefore always associated with a queue.
+
+Queues can be constructed directly from a device object, or by using a filter
+selector string to indicate the device to construct:
+
+.. code:: python
+
+   # construct queue from device object
+   q1 = dpctl.SyclQueue(dev_gpu)
+   # construct queue using filter selector
+   q2 = dpctl.SyclQueue("gpu")
+
+The computational tasks can be stored in an oneAPI native extension in which
+case their submission is orchestrated during Python API calls. Let’s consider a
+function that offloads an evaluation of a polynomial for every point of a NumPy
+array ``X``. Such a function needs to receive a queue object to indicate which
+device the computation must be offloaded to:
+
+.. code:: python
+
+   # allocate space for the result
+   Y = np.empty_like(X)
+   # evaluate polynomial on the device targeted by the queue, Y[i] = p(X[i])
+   onapi_ext.offloaded_poly_evaluate(exec_q, X, Y)
+
+Python call to ``onapi_ext.offloaded_poly_evaluate`` applied to NumPy arrays of
+double precision floating pointer numbers gets translated to the following
+sample C++ code:
+
+.. code:: cpp
+
+   void
+   cpp_offloaded_poly_evaluate(
+     sycl::queue q, const double *X, double *Y, size_t n) {
+       // create buffers from malloc allocations to make data accessible from device
+       sycl::buffer<1, double> buf_X(X, n);
+       sycl::buffer<1, double> buf_Y(Y, n);
+
+       q.submit([&](sycl::handler &cgh) {
+           // create buffer accessors indicating kernel data-flow pattern
+           sycl::accessor acc_X(buf_X, cgh, sycl::read_only);
+           sycl::accessor acc_Y(buf_Y, cgh, sycl::write_only, sycl::no_init);
+
+           cgh.parallel_for(n,
+              // lambda function that gets executed by different work-items with
+              // different arguments in parallel
+              [=](sycl::id<1> id) {
+                 auto x = accX[id];
+                 accY[id] = 3.0 + x * (1.0 + x * (-0.5 + 0.3 * x));
+              });
+       }).wait();
+
+       return;
+   }
+
+We refer an interested reader to an excellent and freely available “`Data
+Parallel C++ <https://link.springer.com/book/10.1007%2F978-1-4842-5574-2>`__”
+book for details of this data parallel C++.
+
+Our package ``numba_dpex`` allows one to write kernels directly in Python.
+
+.. code:: python
+
+   import numba_dpex
+
+
+   @numba_dpex.kernel
+   def numba_dpex_poly(X, Y):
+       i = numba_dpex.get_global_id(0)
+       x = X[i]
+       Y[i] = 3.0 + x * (1.0 + x * (-0.5 + 0.3 * x))
+
+Specifying the execution queue is done using Python context manager:
+
+.. code:: python
+
+   import numpy as np
+
+   X = np.random.randn(10**6)
+   Y = np.empty_like(X)
+
+   with dpctl.device_context(q):
+       # apply the kernel to elements of X, writing value into Y,
+       # while executing using given queue
+       numba_dpex_poly[X.size, numba_dpex.DEFAULT_LOCAL_SIZE](X, Y)
+
+The argument to ``device_context`` can be a queue object, a device object for
+which a temporary queue will be created, or a filter selector string. Thus we
+could have equally used ``dpctl.device_context(gpu_dev)`` or
+``dpctl.device_context("gpu")``.
+
+Note that in this examples data sharing was implicitly managed for us: in the
+case of calling a function from a precompiled oneAPI native extension data
+sharing was managed by DPC++ runtime, while in the case of using ``numba_dpex``
+kernel it was managed during execution of ``__call__`` method.
+
+Data sharing
+~~~~~~~~~~~~
+
+Implicit management of data is surely convenient, but its use in an interpreted
+code comes at a performance cost. A runtime must implicitly copy data from host
+to the device before the kernel execution commences and then copy some (or all)
+of it back after the execution completes for every Python API call.
+
+``dpctl`` provides for allocating memory directly accessible to kernels
+executing on a device using SYCL’s Unified Shared Memory (`USM
+<https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:usm>`__)
+feature. It also implements USM-based ND-array object
+``dpctl.tensor.usm_ndarray`` that conforms `array-API standard
+<https://data-apis.org/array-api/latest/>`__.
+
+.. code:: python
+
+   import dpctl.tensor as dpt
+
+   # allocate array of doubles using USM-device allocation on GPU device
+   X = dpt.arange(0.0, end=1.0, step=1e-6, device="gpu", usm_type="device")
+   # allocate array for the output
+   Y = dpt.empty_like(X)
+
+   # execution queue is inferred from allocation queues.
+   # Kernel is executed on the same device where arrays were allocated
+   numba_dpex_poly[X.size, numba_dpex.DEFAULT_LOCAL_SIZE](X, Y)
+
+The execution queue can be unambiguously determined in this case since both
+arguments are USM arrays with the same allocation queues and ``X.sycl_queue ==
+Y.sycl_queue`` evaluates to ``True``. Should allocation queues be different,
+such an inference becomes ambiguous and ``numba_dpex`` raises
+``IndeterminateExecutionQueueError`` advising user to explicitly migrate the
+data.
+
+Migration can be accomplished either by using ``dpctl.tensor.asarray(X,
+device=target_device)`` to create a copy, or by using
+``X.to_device(target_device)`` method.
+
+A USM array can be copied back into a NumPy array using ``dpt.asnumpy(Y)`` if
+needed.
+
+Compute follows data
+~~~~~~~~~~~~~~~~~~~~
+
+Automatic deduction of the execution queue from allocation queues is consitent
+with “`local control for data allocation target
+<https://data-apis.org/array-api/latest/design_topics/device_support.html>`__”
+in the array API standard. User has full control over memory allocation through
+three keyword arguments present in all `array creation functions
+<https://data-apis.org/array-api/latest/API_specification/creation_functions.html>`__.
+For example, consider
+
+.. code:: python
+   # Use usm_type = 'device' to get USM-device allocation (default),
+   #     usm_type = 'shared' to get USM-shared allocation,
+   #     usm_type = 'host'   to get USM-host allocation
+   # def dpt.empty(..., device=None, usm_type=None, sycl_queue=None) -> dpctl.tensor.usm_ndarray: ...
+
+The keyword ``device`` is `mandated by the array API
+<https://data-apis.org/array-api/latest/design_topics/device_support.html#syntax-for-device-assignment>`__.
+In ``dpctl.tensor`` the allowed values of the keyword are
+
+-  Filter selector string, e.g. ``device="gpu:0"``
+-  Existing ``dpctl.SyclDevice`` object, e.g. ``device=dev_gpu``
+-  Existing ``dpctl.SyclQueue`` object
+-  ``dpctl.tensor.Device`` object instance obtained from an existing USM array,
+   e.g. ``device=X.device``
+
+In all cases, an allocation queue object will be constructed as described
+`earlier <#specifying-offload-target>`__ and stored in the array instance,
+accessible with ``X.sycl_queue``. Instead of using ``device`` keyword, one can
+alternatively use ``sycl_queue`` keyword for readability to directly specify a
+``dpctl.SyclQueue`` object to be used as the allocation queue.
+
+The rationale for storing the allocation queue in the array is that kernels
+submitted to this queue are guaranteed to be able to correctly dereference (i.e.
+access) the USM pointer. Array operations that only involve this single USM
+array can thus execute on the allocation queue, and the output array can be
+allocated on this same allocation queue with the same usm type as the input
+array.
+
+.. note::
+   Reusing the allocation queue of the input
+   array ensures the computational tasks behind the API call can access the
+   array without making implicit copies and the output array is allocated
+   on the same device as the input.
+
+Compute follows data is the rule prescribing deduction of the execution and the
+allocation queue as well as the USM type for the result when multiple USM arrays
+are combined. It stipulates that arrays can be combined if and only if their
+allocation *queues are the same* as measured by ``==`` operator (i.e.
+``X.sycl_queue == Y.sycl_queue`` must evaluate to ``True``). Same queues refer
+to the same underlying task graphs and DPC++ schedulers.
+
+An attempt to combine USM arrays with unsame allocation queues raises an
+exception advising the user to migrate the data. Migration can be accomplished
+either by using ``dpctl.tensor.asarray(X, device=Y.device)`` to create a copy,
+or by using ``X.to_device(Y.device)`` method which can sometime do the migration
+more efficiently.
+
+.. warning::
+   ``dpctl`` and ``numba_dpex`` are both under heavy development. Feel free to file an
+   issue on GitHub or reach out on Gitter should you encounter any issues.
diff --git a/docs/source/user_guide/useful_links.rst b/docs/source/user_guide/useful_links.rst
new file mode 100644
index 0000000000..9d33d29615
--- /dev/null
+++ b/docs/source/user_guide/useful_links.rst
@@ -0,0 +1,41 @@
+.. _useful_links:
+.. include:: ./../ext_links.txt
+
+Useful links
+============
+
+.. list-table:: **Companion documentation**
+   :widths: 70 200
+   :header-rows: 1
+
+   * - Document
+     - Description
+   * - `Data Parallel Extension for Numpy*`_
+     - Documentation for programming NumPy-like codes on data parallel devices
+   * - `Data Parallel Extension for Numba*`_
+     - Documentation for programming Numba codes on data parallel devices the same way as you program Numba on CPU
+   * - `Data Parallel Control`_
+     - Documentation how to manage data and devices, how to interchange data between different tensor implementations,
+       and how to write data parallel extensions
+   * - `Intel VTune Profiler`_
+     - Performance profiler supporting  analysis of bottlenecks from function leve down to low level instructions.
+       Supports Python and Numba
+   * - `Intel Advisor`_
+     - Analyzes native and Python codes and provides an advice for better composition of heterogeneous algorithms
+   * - `Python* Array API Standard`_
+     - Standard for writing portable Numpy-like codes targeting different hardware vendors and frameworks
+       operating with tensor data
+   * - `SYCL*`_
+     - Standard for writing C++-like codes for heterogeneous computing
+   * - `DPC++`_
+     - Free e-book how to program data parallel devices using Data Parallel C++
+   * - `OpenCl*`_
+     - OpenCl* Standard for heterogeneous programming
+   * - `IEEE 754-2019 Standard for Floating-Point Arithmetic`_
+     - Standard for floating-point arithmetic, essential for writing robust numerical codes
+   * - `Numpy*`_
+     - Documentation for Numpy - foundational CPU library for array programming. Used in conjunction with
+       `Data Parallel Extension for Numpy*`_.
+   * - `Numba*`_
+     - Documentation for Numba - Just-In-Time compiler for Numpy-like codes. Used in conjunction with
+       `Data Parallel Extension for Numba*`_.
diff --git a/docs/user_guides/docker.rst b/docs/user_guides/docker.rst
deleted file mode 100644
index 05e2dca809..0000000000
--- a/docs/user_guides/docker.rst
+++ /dev/null
@@ -1,140 +0,0 @@
-.. _docker:
-
-Docker
-======
-
-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/environment.yml b/environment.yml
index 51c329c627..7b1aa00b33 100644
--- a/environment.yml
+++ b/environment.yml
@@ -14,7 +14,6 @@ dependencies:
   - dpctl >=0.14*
   - dpnp >=0.11*
   - mkl >=2021.3.0  # for dpnp
-  - spirv-tools
   - dpcpp-llvm-spirv
   - packaging
   - pytest
diff --git a/environment/coverage.yml b/environment/coverage.yml
index d214b07a77..315233f37d 100644
--- a/environment/coverage.yml
+++ b/environment/coverage.yml
@@ -4,29 +4,19 @@ channels:
   - numba
   - intel
   - numba/label/dev
-  - pkgs/main
+  - conda-forge
   - nodefaults
 dependencies:
-  - python=3.9
-  - defaults:libffi
+  - libffi
   - gxx_linux-64
   - dpcpp_linux-64
-  - numba 0.57*
-  - dppy/label/dev:dpctl
-  - dppy/label/dev:dpnp
-  - spirv-tools
+  - numba=0.57
+  - dpctl
+  - dpnp
   - dpcpp-llvm-spirv
-  - packaging
+  - opencl_rt
+  - coverage
   - pytest
-  - pip
-  - pip:
-    - coverage
-    - pre-commit
-    - flake8
-    - black==20.8b1
-    - pytest-cov
-    - pytest-xdist
-    - pexpect
-variables:
-  CHANNELS: -c defaults -c numba -c intel -c numba/label/dev -c dppy/label/dev --override-channels
-  CHANNELS_DEV: -c dppy/label/dev -c defaults -c numba -c intel -c numba/label/dev --override-channels
+  - pytest-cov
+  - pytest-xdist
+  - pexpect
diff --git a/environment/docs.yml b/environment/docs.yml
index 7d0573c203..14fa086643 100644
--- a/environment/docs.yml
+++ b/environment/docs.yml
@@ -3,29 +3,26 @@ channels:
   - dppy/label/dev
   - numba
   - intel
+  - numba/label/dev
   - conda-forge
   - nodefaults
 dependencies:
-  - python=3.10
+  - libffi
   - gxx_linux-64
   - dpcpp_linux-64
-  - numba 0.57*
-  - dpctl 0.14*
-  - dpnp >=0.12*
-  - spirv-tools
+  - numba=0.57
+  - dpctl
+  - dpnp
   - dpcpp-llvm-spirv
-  - packaging
-  - pytest
+  - opencl_rt
   - pip
   - pip:
-    - coverage
-    - pre-commit
-    - flake8
-    - black==20.8b1
-    - pexpect
     - sphinx
-    - autodoc
+    - autodoc # there is no conda package
     - recommonmark
     - sphinx-rtd-theme
     - sphinxcontrib-apidoc
     - sphinxcontrib-googleanalytics
+    - sphinxcontrib.programoutput
+    - pydata-sphinx-theme
+    - myst-parser
diff --git a/numba_dpex/config.py b/numba_dpex/config.py
index 609fbc2cc7..855c7748ac 100644
--- a/numba_dpex/config.py
+++ b/numba_dpex/config.py
@@ -53,9 +53,6 @@ def __getattr__(name):
 # Dump offload diagnostics
 OFFLOAD_DIAGNOSTICS = _readenv("NUMBA_DPEX_OFFLOAD_DIAGNOSTICS", int, 0)
 
-# Activate Native floating point atomcis support for supported devices.
-# Requires llvm-spirv supporting the FP atomics extension
-NATIVE_FP_ATOMICS = _readenv("NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE", int, 0)
 # Emit debug info
 DEBUG = _readenv("NUMBA_DPEX_DEBUG", int, config.DEBUG)
 DEBUGINFO_DEFAULT = _readenv(
diff --git a/numba_dpex/core/parfors/parfor_lowerer.py b/numba_dpex/core/parfors/parfor_lowerer.py
index f41811a964..07fbe50ba3 100644
--- a/numba_dpex/core/parfors/parfor_lowerer.py
+++ b/numba_dpex/core/parfors/parfor_lowerer.py
@@ -136,7 +136,7 @@ def _build_kernel_arglist(self, kernel_fn, lowerer):
             argtype = kernel_fn.kernel_arg_types[arg_num]
             llvm_val = _getvar(lowerer, arg)
             if isinstance(argtype, DpnpNdArray):
-                datamodel = dpex_dmm.lookup(arg_type)
+                datamodel = dpex_dmm.lookup(argtype)
                 self.kernel_builder.build_array_arg(
                     array_val=llvm_val,
                     array_data_model=datamodel,
diff --git a/numba_dpex/core/targets/kernel_target.py b/numba_dpex/core/targets/kernel_target.py
index 1f46168eb3..2fa8e05066 100644
--- a/numba_dpex/core/targets/kernel_target.py
+++ b/numba_dpex/core/targets/kernel_target.py
@@ -64,7 +64,6 @@ def resolve_argument_type(self, val):
         """
         try:
             numba_type = typeof(val)
-            py_type = type(numba_type)
 
             if isinstance(numba_type, NpArrayType) and not isinstance(
                 numba_type, USMNdArray
@@ -73,19 +72,28 @@ def resolve_argument_type(self, val):
                     type=str(type(val)), value=val
                 )
 
-            # XXX A kernel function has the spir_kernel ABI and requires
-            # pointers to have an address space attribute. For this reason, the
-            # UsmNdArray type uses a custom data model where the pointers are
-            # address space casted to have a SYCL-specific address space value.
-            # The DpnpNdArray type on the other hand is meant to be used inside
-            # host functions and has Numba's array model as its data model.
-            # If the value is a DpnpNdArray then use the ``to_usm_ndarray``
-            # function to convert it into a UsmNdArray type rather than passing
-            # it to the kernel as a DpnpNdArray. Thus, from a Numba typing
-            # perspective dpnp.ndarrays cannot be directly passed to a kernel.
-            if py_type is DpnpNdArray:
-                suai_attrs = get_info_from_suai(val)
-                return to_usm_ndarray(suai_attrs)
+            # A cast from DpnpNdArray type to USMNdArray is needed for all
+            # arguments of DpnpNdArray type. Although, DpnpNdArray derives from
+            # USMNdArray the two types use different data models. USMNdArray
+            # uses the numba_dpex.core.datamodel.models.ArrayModel data model
+            # that defines all CPointer type members in the GLOBAL address
+            # space. The DpnpNdArray uses Numba's default ArrayModel that does
+            # not define pointers in any specific address space. For OpenCL HD
+            # Graphics devices, defining a kernel function (spir_kernel calling
+            # convention) with pointer arguments that have no address space
+            # qualifier causes a run time crash. By casting the argument type
+            # for parfor arguments from DpnpNdArray type to the USMNdArray type
+            # the generated kernel always has an address space qualifier,
+            # avoiding the issue on OpenCL HD graphics devices.
+            if isinstance(numba_type, DpnpNdArray):
+                return USMNdArray(
+                    ndim=numba_type.ndim,
+                    layout=numba_type.layout,
+                    dtype=numba_type.dtype,
+                    usm_type=numba_type.usm_type,
+                    queue=numba_type.queue,
+                )
+
         except ValueError:
             # When an array-like kernel argument is not recognized by
             # numba-dpex, this additional check sees if the array-like object
diff --git a/numba_dpex/examples/kernel/atomic_op.py b/numba_dpex/examples/kernel/atomic_op.py
index 1ff8246dc6..49a5d74671 100644
--- a/numba_dpex/examples/kernel/atomic_op.py
+++ b/numba_dpex/examples/kernel/atomic_op.py
@@ -8,20 +8,28 @@
 
 
 @ndpx.kernel
-def atomic_reduction(a):
+def atomic_reduction(a, res):
+    """Summarize all the items in a and writes it into res using atomic.add.
+
+    :param a: array of values to get sum
+    :param res: result where to add all the items from a array. It must be preset to 0.
+    """
     idx = ndpx.get_global_id(0)
-    ndpx.atomic.add(a, 0, a[idx])
+    ndpx.atomic.add(res, 0, a[idx])
 
 
 def main():
     N = 10
-    a = np.arange(N)
+
+    # We are storing sum to the first element
+    a = np.arange(0, N)
+    res = np.zeros(1, dtype=a.dtype)
 
     print("Using device ...")
     print(a.device)
 
-    atomic_reduction[ndpx.Range(N)](a)
-    print("Reduction sum =", a[0])
+    atomic_reduction[ndpx.Range(N)](a, res)
+    print("Reduction sum =", res[0])
 
     print("Done...")
 
diff --git a/numba_dpex/ocl/__init__.py b/numba_dpex/ocl/__init__.py
index 9abb201da1..00fec5f515 100644
--- a/numba_dpex/ocl/__init__.py
+++ b/numba_dpex/ocl/__init__.py
@@ -1,5 +1,3 @@
 # SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
 #
 # SPDX-License-Identifier: Apache-2.0
-
-from .atomics import atomic_support_present
diff --git a/numba_dpex/ocl/atomics/__init__.py b/numba_dpex/ocl/atomics/__init__.py
index 2043302ed7..00fec5f515 100644
--- a/numba_dpex/ocl/atomics/__init__.py
+++ b/numba_dpex/ocl/atomics/__init__.py
@@ -1,32 +1,3 @@
 # SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
 #
 # SPDX-License-Identifier: Apache-2.0
-
-import os
-import os.path
-
-
-def atomic_support_present():
-    if os.path.isfile(
-        os.path.join(os.path.dirname(__file__), "atomic_ops.spir")
-    ):
-        return True
-    else:
-        return False
-
-
-def get_atomic_spirv_path():
-    if atomic_support_present():
-        return os.path.join(os.path.dirname(__file__), "atomic_ops.spir")
-    else:
-        return None
-
-
-def read_atomic_spirv_file():
-    path = get_atomic_spirv_path()
-    if path:
-        with open(path, "rb") as fin:
-            spirv = fin.read()
-        return spirv
-    else:
-        return None
diff --git a/numba_dpex/ocl/atomics/atomic_ops.cl b/numba_dpex/ocl/atomics/atomic_ops.cl
deleted file mode 100644
index 4e6ce925bd..0000000000
--- a/numba_dpex/ocl/atomics/atomic_ops.cl
+++ /dev/null
@@ -1,143 +0,0 @@
-// SPDX-FileCopyrightText: 2020 - 2022 Intel Corporation
-//
-// SPDX-License-Identifier: Apache-2.0
-
-/* OpenCL extension specification states extensions should define the extension
-   and to use it we need to enable is the following way using the pragma.
-   Link to specification: https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL_Ext.html
-*/
-#ifdef cl_khr_int64_base_atomics
-  #pragma OPENCL EXTENSION cl_khr_int64_base_atomics: enable
-  #ifdef cl_khr_fp64
-    #pragma OPENCL EXTENSION cl_khr_fp64: enable
-
-    double numba_dpex_atomic_cmpxchg_f64_local(volatile __generic double *p, double cmp, double val) {
-        union {
-            ulong  u64;
-            double f64;
-        } cmp_union, val_union, old_union;
-
-        cmp_union.f64 = cmp;
-        val_union.f64 = val;
-        old_union.u64 = atom_cmpxchg((volatile __local ulong *) p, cmp_union.u64, val_union.u64);
-        return old_union.f64;
-    }
-
-    double numba_dpex_atomic_cmpxchg_f64_global(volatile __generic double *p, double cmp, double val) {
-        union {
-            ulong  u64;
-            double f64;
-        } cmp_union, val_union, old_union;
-
-        cmp_union.f64 = cmp;
-        val_union.f64 = val;
-        old_union.u64 = atom_cmpxchg((volatile __global ulong *) p, cmp_union.u64, val_union.u64);
-        return old_union.f64;
-    }
-
-    double numba_dpex_atomic_add_f64_local(volatile __generic double *p, double val) {
-        double  found = *p;
-        double  expected;
-        do {
-            expected = found;
-            found = numba_dpex_atomic_cmpxchg_f64_local(p, expected, expected + val);
-        } while (found != expected);
-        return found;
-    }
-
-    double numba_dpex_atomic_add_f64_global(volatile __generic double *p, double val) {
-        double  found = *p;
-        double  expected;
-        do {
-            expected = found;
-            found = numba_dpex_atomic_cmpxchg_f64_global(p, expected, expected + val);
-        } while (found != expected);
-        return found;
-    }
-
-
-    double numba_dpex_atomic_sub_f64_local(volatile __generic double *p, double val) {
-        double  found = *p;
-        double  expected;
-        do {
-            expected = found;
-            found = numba_dpex_atomic_cmpxchg_f64_local(p, expected, expected - val);
-        } while (found != expected);
-        return found;
-    }
-
-    double numba_dpex_atomic_sub_f64_global(volatile __generic double *p, double val) {
-        double  found = *p;
-        double  expected;
-        do {
-            expected = found;
-            found = numba_dpex_atomic_cmpxchg_f64_global(p, expected, expected - val);
-        } while (found != expected);
-        return found;
-    }
-  #endif
-#endif
-
-float numba_dpex_atomic_cmpxchg_f32_local(volatile __generic float *p, float cmp, float val) {
-    union {
-        unsigned int u32;
-        float        f32;
-    } cmp_union, val_union, old_union;
-
-    cmp_union.f32 = cmp;
-    val_union.f32 = val;
-    old_union.u32 = atomic_cmpxchg((volatile __local unsigned int *) p, cmp_union.u32, val_union.u32);
-    return old_union.f32;
-}
-
-float numba_dpex_atomic_cmpxchg_f32_global(volatile __generic float *p, float cmp, float val) {
-    union {
-        unsigned int u32;
-        float        f32;
-    } cmp_union, val_union, old_union;
-
-    cmp_union.f32 = cmp;
-    val_union.f32 = val;
-    old_union.u32 = atomic_cmpxchg((volatile __global unsigned int *) p, cmp_union.u32, val_union.u32);
-    return old_union.f32;
-}
-
-float numba_dpex_atomic_add_f32_local(volatile __generic float *p, float val) {
-    float found = *p;
-    float expected;
-    do {
-        expected = found;
-        found = numba_dpex_atomic_cmpxchg_f32_local(p, expected, expected + val);
-    } while (found != expected);
-    return found;
-}
-
-float numba_dpex_atomic_add_f32_global(volatile __generic float *p, float val) {
-    float found = *p;
-    float expected;
-    do {
-        expected = found;
-        found = numba_dpex_atomic_cmpxchg_f32_global(p, expected, expected + val);
-    } while (found != expected);
-    return found;
-}
-
-float numba_dpex_atomic_sub_f32_local(volatile __generic float *p, float val) {
-    float found = *p;
-    float expected;
-    do {
-        expected = found;
-        found = numba_dpex_atomic_cmpxchg_f32_local(p, expected, expected - val);
-    } while (found != expected);
-    return found;
-}
-
-float numba_dpex_atomic_sub_f32_global(volatile __generic float *p, float val) {
-    float found = *p;
-    float expected;
-    do {
-        expected = found;
-        found = numba_dpex_atomic_cmpxchg_f32_global(p, expected, expected - val);
-    } while (found != expected);
-    return found;
-}
diff --git a/numba_dpex/ocl/oclimpl.py b/numba_dpex/ocl/oclimpl.py
index ff84b5bb93..e47bb215c4 100644
--- a/numba_dpex/ocl/oclimpl.py
+++ b/numba_dpex/ocl/oclimpl.py
@@ -143,63 +143,6 @@ def sub_group_barrier_impl(context, builder, sig, args):
     return _void_value
 
 
-def insert_and_call_atomic_fn(
-    context, builder, sig, fn_type, dtype, ptr, val, addrspace
-):
-    ll_p = None
-    name = ""
-    if dtype.name == "float32":
-        ll_val = llvmir.FloatType()
-        ll_p = ll_val.as_pointer()
-        if fn_type == "add":
-            name = "numba_dpex_atomic_add_f32"
-        elif fn_type == "sub":
-            name = "numba_dpex_atomic_sub_f32"
-        else:
-            raise TypeError("Operation type is not supported %s" % (fn_type))
-    elif dtype.name == "float64":
-        if True:
-            ll_val = llvmir.DoubleType()
-            ll_p = ll_val.as_pointer()
-            if fn_type == "add":
-                name = "numba_dpex_atomic_add_f64"
-            elif fn_type == "sub":
-                name = "numba_dpex_atomic_sub_f64"
-            else:
-                raise TypeError(
-                    "Operation type is not supported %s" % (fn_type)
-                )
-    else:
-        raise TypeError(
-            "Atomic operation is not supported for type %s" % (dtype.name)
-        )
-
-    if addrspace == address_space.LOCAL:
-        name = name + "_local"
-    else:
-        name = name + "_global"
-
-    assert ll_p is not None
-    assert name != ""
-    ll_p.addrspace = address_space.GENERIC
-
-    mod = builder.module
-    if sig.return_type == types.void:
-        llretty = llvmir.VoidType()
-    else:
-        llretty = context.get_value_type(sig.return_type)
-
-    llargs = [ll_p, context.get_value_type(sig.args[2])]
-    fnty = llvmir.FunctionType(llretty, llargs)
-
-    fn = cgutils.get_or_insert_function(mod, fnty, name)
-    fn.calling_convention = kernel_target.CC_SPIR_FUNC
-
-    generic_ptr = context.addrspacecast(builder, ptr, address_space.GENERIC)
-
-    return builder.call(fn, [generic_ptr, val])
-
-
 def native_atomic_add(context, builder, sig, args):
     aryty, indty, valty = sig.args
     ary, inds, val = args
@@ -282,27 +225,29 @@ def native_atomic_add(context, builder, sig, args):
     return builder.call(fn, fn_args)
 
 
+def support_atomic(dtype: types.Type) -> bool:
+    # This check should be the same as described in sycl documentation:
+    # https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:atomic-references
+    # If atomic is not supported, it will be emulated by the sycl compiler.
+    return (
+        dtype == types.int32
+        or dtype == types.uint32
+        or dtype == types.float32
+        or dtype == types.int64
+        or dtype == types.uint64
+        or dtype == types.float64
+    )
+
+
 @lower(stubs.atomic.add, types.Array, types.intp, types.Any)
 @lower(stubs.atomic.add, types.Array, types.UniTuple, types.Any)
 @lower(stubs.atomic.add, types.Array, types.Tuple, types.Any)
 def atomic_add_tuple(context, builder, sig, args):
-    device_type = dpctl.get_current_queue().sycl_device.device_type
     dtype = sig.args[0].dtype
-
-    if dtype == types.float32 or dtype == types.float64:
-        if (
-            device_type == dpctl.device_type.gpu
-            and config.NATIVE_FP_ATOMICS == 1
-        ):
-            return native_atomic_add(context, builder, sig, args)
-        else:
-            # Currently, DPCPP only supports native floating point
-            # atomics for GPUs.
-            return atomic_add(context, builder, sig, args, "add")
-    elif dtype == types.int32 or dtype == types.int64:
+    if support_atomic(dtype):
         return native_atomic_add(context, builder, sig, args)
     else:
-        raise TypeError("Atomic operation on unsupported type %s" % dtype)
+        raise TypeError(f"Atomic operation on unsupported type {dtype}")
 
 
 def atomic_sub_wrapper(context, builder, sig, args):
@@ -337,81 +282,11 @@ def atomic_sub_wrapper(context, builder, sig, args):
 @lower(stubs.atomic.sub, types.Array, types.UniTuple, types.Any)
 @lower(stubs.atomic.sub, types.Array, types.Tuple, types.Any)
 def atomic_sub_tuple(context, builder, sig, args):
-    device_type = dpctl.get_current_queue().sycl_device.device_type
     dtype = sig.args[0].dtype
-
-    if dtype == types.float32 or dtype == types.float64:
-        if (
-            device_type == dpctl.device_type.gpu
-            and config.NATIVE_FP_ATOMICS == 1
-        ):
-            return atomic_sub_wrapper(context, builder, sig, args)
-        else:
-            # Currently, DPCPP only supports native floating point
-            # atomics for GPUs.
-            return atomic_add(context, builder, sig, args, "sub")
-    elif dtype == types.int32 or dtype == types.int64:
+    if support_atomic(dtype):
         return atomic_sub_wrapper(context, builder, sig, args)
     else:
-        raise TypeError("Atomic operation on unsupported type %s" % dtype)
-
-
-def atomic_add(context, builder, sig, args, name):
-    from .atomics import atomic_support_present
-
-    if atomic_support_present():
-        context.extra_compile_options[kernel_target.LINK_ATOMIC] = True
-        aryty, indty, valty = sig.args
-        ary, inds, val = args
-        dtype = aryty.dtype
-
-        if indty == types.intp:
-            indices = [inds]  # just a single integer
-            indty = [indty]
-        else:
-            indices = cgutils.unpack_tuple(builder, inds, count=len(indty))
-            indices = [
-                context.cast(builder, i, t, types.intp)
-                for t, i in zip(indty, indices)
-            ]
-
-        if dtype != valty:
-            raise TypeError("expecting %s but got %s" % (dtype, valty))
-
-        if aryty.ndim != len(indty):
-            raise TypeError(
-                "indexing %d-D array with %d-D index" % (aryty.ndim, len(indty))
-            )
-
-        lary = context.make_array(aryty)(context, builder, ary)
-        ptr = cgutils.get_item_pointer(context, builder, aryty, lary, indices)
-
-        if isinstance(aryty, Array) and aryty.addrspace == address_space.LOCAL:
-            return insert_and_call_atomic_fn(
-                context,
-                builder,
-                sig,
-                name,
-                dtype,
-                ptr,
-                val,
-                address_space.LOCAL,
-            )
-        else:
-            return insert_and_call_atomic_fn(
-                context,
-                builder,
-                sig,
-                name,
-                dtype,
-                ptr,
-                val,
-                address_space.GLOBAL,
-            )
-    else:
-        raise ImportError(
-            "Atomic support is not present, can not perform atomic_add"
-        )
+        raise TypeError(f"Atomic operation on unsupported type {dtype}")
 
 
 @lower(stubs.private.array, types.IntegerLiteral, types.Any)
diff --git a/numba_dpex/spirv_generator.py b/numba_dpex/spirv_generator.py
index fdfe7f07e8..75627edb77 100644
--- a/numba_dpex/spirv_generator.py
+++ b/numba_dpex/spirv_generator.py
@@ -5,7 +5,6 @@
 """A wrapper to connect to the SPIR-V binaries (Tools, Translator)."""
 
 import os
-import shutil
 import tempfile
 from subprocess import CalledProcessError, check_call
 
@@ -75,8 +74,6 @@ def generate(self, llvm_spirv_args, ipath, opath):
         if config.DEBUG:
             llvm_spirv_flags.append("--spirv-debug-info-version=ocl-100")
 
-        if not config.NATIVE_FP_ATOMICS:
-            llvm_spirv_args = ["--spirv-max-version", "1.4"] + llvm_spirv_args
         llvm_spirv_tool = self._llvm_spirv()
 
         if config.DEBUG:
@@ -96,17 +93,6 @@ def _llvm_spirv():
         result = dls.get_llvm_spirv_path()
         return result
 
-    def link(self, opath, binaries):
-        """
-        Link spirv modules.
-
-        Args:
-            opath: Output file path of the linked final spirv.
-            binaries: Spirv modules to be linked.
-        """
-        flags = ["--allow-partial-linkage"]
-        check_call(["spirv-link", *flags, "-o", opath, *binaries])
-
 
 class Module(object):
     def __init__(self, context, llvmir, llvmbc):
@@ -156,15 +142,9 @@ def finalize(self):
         # Generate SPIR-V from "friendly" LLVM-based SPIR 2.0
         spirv_path = self._track_temp_file("generated-spirv")
 
-        binary_paths = [spirv_path]
-
         llvm_spirv_args = []
         for key in list(self.context.extra_compile_options.keys()):
-            if key == LINK_ATOMIC:
-                from .ocl.atomics import get_atomic_spirv_path
-
-                binary_paths.append(get_atomic_spirv_path())
-            elif key == LLVM_SPIRV_ARGS:
+            if key == LLVM_SPIRV_ARGS:
                 llvm_spirv_args = self.context.extra_compile_options[key]
             del self.context.extra_compile_options[key]
 
@@ -188,10 +168,6 @@ def finalize(self):
             opath=spirv_path,
         )
 
-        if len(binary_paths) > 1:
-            spirv_path = self._track_temp_file("linked-spirv")
-            self._cmd.link(spirv_path, binary_paths)
-
         if config.SAVE_IR_FILES != 0:
             # Dump the llvmir and llvmbc in file
             with open("generated_spirv.spir", "wb") as f1:
diff --git a/setup.py b/setup.py
index 4f1080560a..1801ceccdc 100644
--- a/setup.py
+++ b/setup.py
@@ -3,16 +3,12 @@
 # SPDX-License-Identifier: Apache-2.0
 
 import os
-import shutil
-import subprocess
 import sys
 import sysconfig
 
 import dpctl
 import numba
 import numpy
-import setuptools.command.develop as orig_develop
-import setuptools.command.install as orig_install
 from setuptools import Extension, find_packages, setup
 
 import versioneer
@@ -62,83 +58,6 @@ def get_ext_modules():
     return ext_modules
 
 
-class install(orig_install.install):
-    def run(self):
-        spirv_compile()
-        super().run()
-
-
-class develop(orig_develop.develop):
-    def run(self):
-        spirv_compile()
-        super().run()
-
-
-def _get_cmdclass():
-    cmdclass = versioneer.get_cmdclass()
-    cmdclass["install"] = install
-    cmdclass["develop"] = develop
-    return cmdclass
-
-
-def spirv_compile():
-    if IS_LIN:
-        compiler = "icx"
-    if IS_WIN:
-        compiler = "clang.exe"
-
-    clang_args = [
-        compiler,
-        "-flto",
-        "-fveclib=none",
-        "-target",
-        "spir64-unknown-unknown",
-        "-c",
-        "-x",
-        "cl",
-        "-emit-llvm",
-        "-cl-std=CL2.0",
-        "-Xclang",
-        "-finclude-default-header",
-        "numba_dpex/ocl/atomics/atomic_ops.cl",
-        "-o",
-        "numba_dpex/ocl/atomics/atomic_ops.bc",
-    ]
-    spirv_args = [
-        _llvm_spirv(),
-        "--spirv-max-version",
-        "1.4",
-        "numba_dpex/ocl/atomics/atomic_ops.bc",
-        "-o",
-        "numba_dpex/ocl/atomics/atomic_ops.spir",
-    ]
-    subprocess.check_call(
-        clang_args,
-        stderr=subprocess.DEVNULL,
-        stdout=subprocess.DEVNULL,
-        shell=False,
-    )
-    subprocess.check_call(
-        spirv_args,
-        stderr=subprocess.DEVNULL,
-        stdout=subprocess.DEVNULL,
-        shell=False,
-    )
-
-
-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")
-
-    result = dls.get_llvm_spirv_path()
-
-    return result
-
-
 packages = find_packages(
     include=["numba_dpex", "numba_dpex.*", "_dpexrt_python"]
 )
@@ -151,7 +70,6 @@ def _llvm_spirv():
 metadata = dict(
     name="numba-dpex",
     version=versioneer.get_version(),
-    cmdclass=_get_cmdclass(),
     description="An extension for Numba to add data-parallel offload capability",
     url="https://github.com/IntelPython/numba-dpex",
     packages=packages,


Test patch:
diff --git a/conda-recipe/run_test.sh b/conda-recipe/run_test.sh
index 4454e3abae..80c7f6ba35 100755
--- a/conda-recipe/run_test.sh
+++ b/conda-recipe/run_test.sh
@@ -5,15 +5,8 @@ unset ONEAPI_DEVICE_SELECTOR
 
 for selector in $(python -c "import dpctl; print(\" \".join([dev.backend.name+\":\"+dev.device_type.name for dev in dpctl.get_devices() if dev.device_type.name in [\"cpu\",\"gpu\"]]))")
 do
-    export "ONEAPI_DEVICE_SELECTOR=$selector"
-    unset NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1
-
+    ONEAPI_DEVICE_SELECTOR=$selector \
     pytest -q -ra --disable-warnings --pyargs numba_dpex -vv
-
-    export NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1
-
-    pytest -q -ra --disable-warnings -vv \
-        --pyargs numba_dpex.tests.kernel_tests.test_atomic_op::test_atomic_fp_native
 done
 
 exit 0
diff --git a/numba_dpex/tests/dpjit_tests/parfors/prange/__init__.py b/numba_dpex/tests/dpjit_tests/parfors/prange/__init__.py
new file mode 100644
index 0000000000..e69de29bb2
diff --git a/numba_dpex/tests/dpjit_tests/parfors/prange/test_pairwise_distance.py b/numba_dpex/tests/dpjit_tests/parfors/prange/test_pairwise_distance.py
new file mode 100644
index 0000000000..834cae54e3
--- /dev/null
+++ b/numba_dpex/tests/dpjit_tests/parfors/prange/test_pairwise_distance.py
@@ -0,0 +1,53 @@
+# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import dpctl
+import dpnp
+import numba as nb
+import pytest
+
+from numba_dpex import dpjit
+
+
+def test_pairwise_distance():
+    @dpjit
+    def pairwise_distance(X1, X2, D):
+        """Naïve pairwise distance impl - take an array representing M points in N
+        dimensions, and return the M x M matrix of Euclidean distances
+
+        Args:
+            X1 : Set of points
+            X2 : Set of points
+            D  : Outputted distance matrix
+        """
+        # Size of inputs
+        X1_rows = X1.shape[0]
+        X2_rows = X2.shape[0]
+        X1_cols = X1.shape[1]
+
+        # TODO: get rid of it once prange supports dtype
+        # https://github.com/IntelPython/numba-dpex/issues/1063
+        float0 = X1.dtype.type(0.0)
+
+        # Outermost parallel loop over the matrix X1
+        for i in nb.prange(X1_rows):
+            # Loop over the matrix X2
+            for j in range(X2_rows):
+                d = float0
+                # Compute exclidean distance
+                for k in range(X1_cols):
+                    tmp = X1[i, k] - X2[j, k]
+                    d += tmp * tmp
+                # Write computed distance to distance matrix
+                D[i, j] = dpnp.sqrt(d)
+
+    q = dpctl.SyclQueue()
+    X1 = dpnp.ones((100, 2), sycl_queue=q)
+    X2 = dpnp.ones((100, 2), sycl_queue=q)
+    D = dpnp.empty((100, 100), sycl_queue=q)
+
+    try:
+        pairwise_distance(X1, X2, D)
+    except:
+        pytest.fail("Failed to compile prange loop for pairwise distance calc")
diff --git a/numba_dpex/tests/kernel_tests/test_atomic_op.py b/numba_dpex/tests/kernel_tests/test_atomic_op.py
index 3f18b20e57..d4a28f8e1c 100644
--- a/numba_dpex/tests/kernel_tests/test_atomic_op.py
+++ b/numba_dpex/tests/kernel_tests/test_atomic_op.py
@@ -2,12 +2,10 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
-import dpctl
 import dpnp as np
 import pytest
 
 import numba_dpex as dpex
-from numba_dpex import config
 from numba_dpex.core.descriptor import dpex_kernel_target
 from numba_dpex.tests._helper import override_config
 
@@ -66,13 +64,6 @@ def f(a):
     return dpex.kernel(f), request.param[1]
 
 
-skip_no_atomic_support = pytest.mark.skipif(
-    not dpex.ocl.atomic_support_present(),
-    reason="No atomic support",
-)
-
-
-@skip_no_atomic_support
 def test_kernel_atomic_simple(input_arrays, kernel_result_pair):
     a, dtype = input_arrays()
     kernel, expected = kernel_result_pair
@@ -111,7 +102,6 @@ def f(a):
     return f
 
 
-@skip_no_atomic_support
 def test_kernel_atomic_local(input_arrays, return_list_of_op):
     a, dtype = input_arrays()
     op_type, expected = return_list_of_op
@@ -148,7 +138,6 @@ def f(a):
     return dpex.kernel(f)
 
 
-@skip_no_atomic_support
 def test_kernel_atomic_multi_dim(
     return_list_of_op, return_list_of_dim, return_dtype
 ):
@@ -160,23 +149,6 @@ def test_kernel_atomic_multi_dim(
     assert a[0] == expected
 
 
-skip_NATIVE_FP_ATOMICS_0 = pytest.mark.skipif(
-    not config.NATIVE_FP_ATOMICS, reason="Native FP atomics disabled"
-)
-
-
-def skip_if_disabled(*args):
-    return pytest.param(*args, marks=skip_NATIVE_FP_ATOMICS_0)
-
-
-@skip_no_atomic_support
-@pytest.mark.parametrize(
-    "NATIVE_FP_ATOMICS, expected_native_atomic_for_device",
-    [
-        skip_if_disabled(1, lambda device: device != "opencl:cpu:0"),
-        (0, lambda device: False),
-    ],
-)
 @pytest.mark.parametrize(
     "function_generator", [get_func_global, get_func_local]
 )
@@ -189,8 +161,6 @@ def skip_if_disabled(*args):
 )
 @pytest.mark.parametrize("dtype", list_of_f_dtypes)
 def test_atomic_fp_native(
-    NATIVE_FP_ATOMICS,
-    expected_native_atomic_for_device,
     function_generator,
     operator_name,
     expected_spirv_function,
@@ -206,16 +176,13 @@ def test_atomic_fp_native(
         for arg in args
     ]
 
-    with override_config("NATIVE_FP_ATOMICS", NATIVE_FP_ATOMICS):
-        kernel.compile(
-            args=argtypes,
-            debug=False,
-            compile_flags=None,
-            target_ctx=dpex_kernel_target.target_context,
-            typing_ctx=dpex_kernel_target.typing_context,
-        )
-
-        is_native_atomic = expected_spirv_function in kernel._llvm_module
-        assert is_native_atomic == expected_native_atomic_for_device(
-            dpctl.select_default_device().filter_string
-        )
+    kernel.compile(
+        args=argtypes,
+        debug=False,
+        compile_flags=None,
+        target_ctx=dpex_kernel_target.target_context,
+        typing_ctx=dpex_kernel_target.typing_context,
+    )
+
+    # TODO: this may fail if code is generated for platform that emulates atomic support?
+    assert expected_spirv_function in kernel._llvm_module