From 5bea1e0b660d3df72405169b85230aade20031a5 Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Wed, 15 Apr 2026 13:35:11 -0700
Subject: [PATCH 01/24] add new pre-commit linters

---
 .pre-commit-config.yaml | 53 +++++++++++++++++++++++++++++++++++++++++
 1 file changed, 53 insertions(+)

diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
index a3777a0975..f2cb4fbd86 100644
--- a/.pre-commit-config.yaml
+++ b/.pre-commit-config.yaml
@@ -10,8 +10,45 @@ repos:
 -   repo: https://github.com/pre-commit/pre-commit-hooks
     rev: v6.0.0
     hooks:
+  # Git
+    -   id: check-added-large-files
+    -   id: no-commit-to-branch
+        name: "ensure no direct commit to master/maintenance branches"
+        args: [--branch, "master", --pattern, "maintenance/.*"]
+    -   id: check-case-conflict
+    -   id: check-illegal-windows-names
+    -   id: check-ast
+    -   id: check-builtin-literals
+    -   id: check-case-conflict
+    -   id: check-executables-have-shebangs
+    -   id: check-merge-conflict
+    -   id: check-toml
+    -   id: debug-statements
+    -   id: destroyed-symlinks
     -   id: end-of-file-fixer
+    -   id: fix-byte-order-marker
+    -   id: mixed-line-ending
     -   id: trailing-whitespace
+-   repo: https://github.com/pre-commit/pygrep-hooks
+    rev: v1.10.0
+    hooks:
+    -   id: python-check-blanket-noqa
+    -   id: python-check-blanket-type-ignore
+    -   id: python-check-mock-methods
+    -   id: python-no-eval
+    -   id: python-no-log-warn
+    -   id: python-use-type-annotations
+    -   id: rst-backticks
+    -   id: rst-directive-colons
+    -   id: rst-inline-touching-normal
+    -   id: text-unicode-replacement-char
+-   repo: https://github.com/codespell-project/codespell
+    rev: v2.4.1
+    hooks:
+    - id: codespell
+      args: ["-L", "nd,hart,elemente,wirth"]
+      additional_dependencies:
+      - tomli
 -   repo: https://github.com/psf/black
     rev: 26.3.1
     hooks:
@@ -47,3 +84,19 @@ repos:
   hooks:
   - id: cython-lint
   - id: double-quote-cython-strings
+
+- repo: https://github.com/jumanjihouse/pre-commit-hooks
+  rev: 3.0.0
+  hooks:
+  - id: shellcheck
+
+- repo: https://github.com/gitleaks/gitleaks
+  rev: v8.30.0
+  hooks:
+  - id: gitleaks
+
+- repo: https://github.com/rhysd/actionlint
+  rev: v1.7.11
+  hooks:
+  - id: actionlint
+    args: ["-ignore", "SC2317"]

From 27a91367b772a502e38bc208fc3bb58dfe74015b Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Wed, 15 Apr 2026 14:01:28 -0700
Subject: [PATCH 02/24] pre-commit and codespell tweaks

---
 .pre-commit-config.yaml | 4 +++-
 pyproject.toml          | 6 ++++++
 2 files changed, 9 insertions(+), 1 deletion(-)

diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
index f2cb4fbd86..08fe5cee1a 100644
--- a/.pre-commit-config.yaml
+++ b/.pre-commit-config.yaml
@@ -1,5 +1,6 @@
 # See https://pre-commit.com for more information
 # See https://pre-commit.com/hooks.html for more hooks
+exclude: "docs/_legacy/"
 repos:
 -   repo: https://github.com/PyCQA/bandit
     rev: '1.9.4'
@@ -46,9 +47,10 @@ repos:
     rev: v2.4.1
     hooks:
     - id: codespell
-      args: ["-L", "nd,hart,elemente,wirth"]
+      args: ["-L", "ba,som,xWindows"]  # ignore some variable names
       additional_dependencies:
       - tomli
+      exclude: "docs/doxyrest-config.lua.in"
 -   repo: https://github.com/psf/black
     rev: 26.3.1
     hooks:
diff --git a/pyproject.toml b/pyproject.toml
index 92be4c0771..019dd178c6 100644
--- a/pyproject.toml
+++ b/pyproject.toml
@@ -86,6 +86,12 @@ Repository = "https://github.com/IntelPython/dpctl.git"
 exclude = "dpctl/_version.py"
 line-length = 80
 
+[tool.codespell]
+builtin = "clear,rare,informal,names"
+check-filenames = true
+ignore-words-list = "nd"
+quiet-level = 3
+
 [tool.coverage.report]
 omit = [
   "dpctl/tests/*",

From 21491c42a42eb23d8062a4031d371deebe2e2572 Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Wed, 15 Apr 2026 14:02:53 -0700
Subject: [PATCH 03/24] update pre-commit workflow

---
 .github/workflows/pre-commit.yml | 9 ++++++++-
 1 file changed, 8 insertions(+), 1 deletion(-)

diff --git a/.github/workflows/pre-commit.yml b/.github/workflows/pre-commit.yml
index 7a136e5903..8c8f6fc9a9 100644
--- a/.github/workflows/pre-commit.yml
+++ b/.github/workflows/pre-commit.yml
@@ -9,13 +9,20 @@ permissions: read-all
 
 jobs:
   pre-commit:
+    name: pre-commit
+
     runs-on: ubuntu-24.04
     timeout-minutes: 30
+
     steps:
     - uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
     - uses: actions/setup-python@a309ff8b426b58ec0e2a45f0f869d46889d02405 # v6.2.0
       with:
-        python-version: '3.12'
+        python-version: '3.14'
+    - name: Install from PyPI
+      shell: bash -l {0}
+      run: |
+        pip install codespell pylint
     - name: Version of clang-format
       run: |
         clang-format --version

From 5cf074899f226726d07026db9899bb7632d8c1de Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Wed, 15 Apr 2026 18:51:03 -0700
Subject: [PATCH 04/24] make adjustments to pass pre-commit

---
 .github/workflows/conda-package.yml           | 196 +++++++++---------
 .github/workflows/generate-coverage.yaml      |  16 +-
 .github/workflows/generate-docs.yml           |  18 +-
 .github/workflows/os-llvm-sycl-build.yml      |  45 ++--
 .../workflows/run-tests-from-dppy-bits.yaml   |  20 +-
 README.md                                     |   2 +-
 cmake/IntelSYCLConfig.cmake                   |   4 +-
 conda-recipe/build.sh                         |   6 +-
 docs/CMakeLists.txt                           |   4 +-
 .../docfiles/user_guides/QuickStart.rst       |   2 +-
 docs/doc_sources/_static/dpctl.svg            |   0
 .../dpctl/filter_selector_string.rst          |   2 +-
 .../doc_sources/api_reference/dpctl_cmake.rst |   2 +-
 .../beginners_guides/installation.rst         |   2 +-
 .../beginners_guides/managing_devices.rst     |   2 +-
 .../contributor_guides/building.rst           |   2 +-
 docs/doc_sources/extlinks_gen.py              |   2 +-
 docs/doc_sources/index.rst                    |   2 +-
 .../user_guides/basic_concepts.rst            |   2 +-
 dpctl/_host_task_util.hpp                     |   2 +-
 dpctl/_sycl_device.pyx                        |   8 +-
 dpctl/_sycl_device_factory.pyx                |   2 +-
 dpctl/_sycl_queue.pyx                         |   2 +-
 dpctl/_sycl_queue_manager.pyx                 |   2 +-
 dpctl/_sycl_timer.py                          |   2 +-
 dpctl/apis/include/dpctl4pybind11.hpp         |   6 +-
 dpctl/tests/helper/_helper.py                 |   2 +-
 dpctl/tests/test_sycl_device.py               |   4 +-
 dpctl/tests/test_sycl_event.py                |   2 +-
 dpctl/tests/test_sycl_queue.py                |  10 +-
 dpctl/tests/test_sycl_usm.py                  |  12 +-
 dpctl/tests/test_utils.py                     |   2 +-
 dpctl/utils/CMakeLists.txt                    |   2 +-
 dpctl/utils/_intel_device_info.py             |   2 +-
 examples/cython/README.md                     |   2 +-
 examples/cython/sycl_buffer/README.md         |   2 +-
 examples/cython/sycl_buffer/scripts/bench.py  |   2 +-
 .../pybind11/onemkl_gemv/sycl_gemm/build.sh   |  19 +-
 examples/python/_runner.py                    |   2 +-
 examples/python/subdevices.py                 |   4 +-
 examples/python/usm_memory_host_access.py     |   2 +-
 .../cmake/modules/GetProjectVersion.cmake     |   6 +-
 libsyclinterface/dbg_build.bat                |   2 +-
 libsyclinterface/dbg_build.sh                 |   9 +-
 libsyclinterface/dbg_build_custom.sh          |  18 +-
 .../helper/source/dpctl_error_handlers.cpp    |   2 +-
 .../include/syclinterface/Support/DllExport.h |   2 +-
 .../syclinterface/dpctl_sycl_device_manager.h |   2 +-
 .../dpctl_sycl_platform_manager.h             |   2 +-
 libsyclinterface/tests/test_service.cpp       |   2 +-
 50 files changed, 246 insertions(+), 221 deletions(-)
 mode change 100755 => 100644 cmake/IntelSYCLConfig.cmake
 mode change 100755 => 100644 docs/doc_sources/_static/dpctl.svg

diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml
index ae6ff28610..8924245c78 100644
--- a/.github/workflows/conda-package.yml
+++ b/.github/workflows/conda-package.yml
@@ -50,13 +50,13 @@ jobs:
             ${{ runner.os }}-conda-${{ env.CACHE_NUMBER }}-python-${{ matrix.python }}-
             ${{ runner.os }}-conda-${{ env.CACHE_NUMBER }}-
       - name: Add conda to system path
-        run: echo $CONDA/bin >> $GITHUB_PATH
+        run: echo "$CONDA/bin" >> "$GITHUB_PATH"
       - name: Install conda-build
         run: conda install conda-build -c conda-forge --override-channels
       - name: Store conda paths as envs
         shell: bash -l {0}
         run: |
-          echo "WHEELS_OUTPUT_FOLDER=$GITHUB_WORKSPACE${{ runner.os == 'Linux' && '/' || '\\' }}" >> $GITHUB_ENV
+          echo "WHEELS_OUTPUT_FOLDER=$GITHUB_WORKSPACE${{ runner.os == 'Linux' && '/' || '\\' }}" >> "$GITHUB_ENV"
       - name: Build conda package
         run: |
           # use bootstrap channel to pull NumPy linked with OpenBLAS
@@ -64,9 +64,9 @@ jobs:
           VERSIONS="--python ${{ matrix.python }} --numpy 2.0"
           TEST="--no-test"
           conda build \
-            $TEST \
-            $VERSIONS \
-            $CHANNELS \
+            "$TEST" \
+            "$VERSIONS" \
+            "$CHANNELS" \
             conda-recipe
       - name: Upload artifact
         uses: actions/upload-artifact@043fb46d1a93c77aae656e7c1c64a875d1fc6a0a # v7.0.1
@@ -126,8 +126,8 @@ jobs:
       - name: Store conda paths as envs
         shell: bash -l {0}
         run: |
-          echo "CONDA_BLD=$CONDA/conda-bld/win-64/" | tr "\\\\" '/' >> $GITHUB_ENV
-          echo "WHEELS_OUTPUT_FOLDER=$GITHUB_WORKSPACE${{ runner.os == 'Linux' && '/' || '\\' }}" >> $GITHUB_ENV
+          echo "CONDA_BLD=$CONDA/conda-bld/win-64/" | tr "\\\\" '/' >> "$GITHUB_ENV"
+          echo "WHEELS_OUTPUT_FOLDER=$GITHUB_WORKSPACE${{ runner.os == 'Linux' && '/' || '\\' }}" >> "$GITHUB_ENV"
 
       - name: Build conda package
         env:
@@ -163,16 +163,16 @@ jobs:
     steps:
       - name: Construct channels line
         run: |
-          echo "CHANNELS=-c ${{ env.INTEL_CHANNEL }} -c conda-forge --override-channels" >> $GITHUB_ENV
+          echo "CHANNELS=-c ${{ env.INTEL_CHANNEL }} -c conda-forge --override-channels" >> "$GITHUB_ENV"
       - name: Display channels line
         run: |
-          echo ${{ env.CHANNELS }}
+          echo "${{ env.CHANNELS }}"
       - name: Download artifact
         uses: actions/download-artifact@3e5f45b2cfb9172054b4087a40e8e0b5a5461e7c # v8.0.1
         with:
           name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Python ${{ matrix.python }}
       - name: Add conda to system path
-        run: echo $CONDA/bin >> $GITHUB_PATH
+        run: echo "$CONDA/bin" >> "$GITHUB_PATH"
       - name: Install conda-index
         # Needed to be able to run conda index
         run: |
@@ -180,18 +180,20 @@ jobs:
           conda install conda-index -c conda-forge --override-channels
       - name: Create conda channel
         run: |
-          mkdir -p $GITHUB_WORKSPACE/channel/linux-64
-          conda index $GITHUB_WORKSPACE/channel || exit 1
-          mv ${PACKAGE_NAME}-*.conda $GITHUB_WORKSPACE/channel/linux-64 || exit 1
-          conda index $GITHUB_WORKSPACE/channel || exit 1
+          mkdir -p "$GITHUB_WORKSPACE/channel/linux-64"
+          conda index "$GITHUB_WORKSPACE/channel" || exit 1
+          mv "${PACKAGE_NAME}"-*.conda "$GITHUB_WORKSPACE/channel/linux-64" || exit 1
+          conda index "$GITHUB_WORKSPACE/channel" || exit 1
           # Test channel
-          conda search $PACKAGE_NAME -c $GITHUB_WORKSPACE/channel --override-channels --info --json > $GITHUB_WORKSPACE/ver.json
+          conda search "$PACKAGE_NAME" -c "$GITHUB_WORKSPACE/channel" --override-channels --info --json > "$GITHUB_WORKSPACE/ver.json"
           cat ver.json
       - name: Collect dependencies
         run: |
           CHANNELS="-c $GITHUB_WORKSPACE/channel ${{ env.CHANNELS }}"
-          export PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
-          conda create -n ${{ env.TEST_ENV_NAME }} $PACKAGE_NAME=${PACKAGE_VERSION} python=${{ matrix.python }} $CHANNELS --only-deps --dry-run > lockfile
+          export CHANNELS
+          PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
+          export PACKAGE_VERSION
+          conda create -n "${{ env.TEST_ENV_NAME }}" "${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION}" python=${{ matrix.python }} "$CHANNELS" --only-deps --dry-run > lockfile
           cat lockfile
       - name: Set pkgs_dirs
         run: |
@@ -209,16 +211,19 @@ jobs:
             ${{ runner.os }}-conda-${{ env.CACHE_NUMBER }}-
       - name: Install dpctl
         run: |
-          export CHANNELS="-c $GITHUB_WORKSPACE/channel ${{ env.CHANNELS }}"
-          export TEST_DEPENDENCIES="pytest pytest-cov cython setuptools"
-          export PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
-          conda create -n ${{ env.TEST_ENV_NAME }} $PACKAGE_NAME=${PACKAGE_VERSION} ${TEST_DEPENDENCIES} python=${{ matrix.python }} ${CHANNELS}
+          CHANNELS="-c $GITHUB_WORKSPACE/channel ${{ env.CHANNELS }}"
+          export CHANNELS
+          TEST_DEPENDENCIES="pytest pytest-cov cython setuptools"
+          export TEST_DEPENDENCIES
+          PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
+          export PACKAGE_VERSION
+          conda create -n "${{ env.TEST_ENV_NAME }}" "${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION}" "${TEST_DEPENDENCIES}" python=${{ matrix.python }} "${CHANNELS}"
           # Test installed packages
-          conda list -n ${{ env.TEST_ENV_NAME }}
+          conda list -n "${{ env.TEST_ENV_NAME }}"
       - name: Smoke test
         run: |
-          . $CONDA/etc/profile.d/conda.sh
-          conda activate ${{ env.TEST_ENV_NAME }}
+          . "$CONDA/etc/profile.d/conda.sh"
+          conda activate "${{ env.TEST_ENV_NAME }}"
           python -c "import dpctl; dpctl.lsplatform(verbosity=2)"
       - name: Install gdb
         run: |
@@ -226,21 +231,21 @@ jobs:
           sudo apt-get install -y gdb
       - name: Run test_elementwise under gdb
         run: |
-          . $CONDA/etc/profile.d/conda.sh
-          conda activate ${{ env.TEST_ENV_NAME }}
-          gdb --batch -ex r -ex 'info sharedlibrary' -ex 'set print elements 1000' -ex bt --args ${CONDA_PREFIX}/bin/python -m pytest -q -ra --disable-warnings --pyargs dpctl.tests.elementwise.test_trigonometric::test_trig_order -vv || true
+          . "$CONDA/etc/profile.d/conda.sh"
+          conda activate "${{ env.TEST_ENV_NAME }}"
+          gdb --batch -ex r -ex 'info sharedlibrary' -ex 'set print elements 1000' -ex bt --args "${CONDA_PREFIX}/bin/python" -m pytest -q -ra --disable-warnings --pyargs dpctl.tests.elementwise.test_trigonometric::test_trig_order -vv || true
       - name: Create test temp dir
         # create temporary empty folder to runs tests from
         # https://github.com/pytest-dev/pytest/issues/11904
-        run: mkdir -p ${GITHUB_WORKSPACE}/test_tmp
+        run: mkdir -p "${GITHUB_WORKSPACE}/test_tmp"
       - name: Run tests
         working-directory: ${{ github.workspace }}/test_tmp
         env:
           SYCL_CACHE_PERSISTENT: 1
         run: |
-          . $CONDA/etc/profile.d/conda.sh
-          conda activate ${{ env.TEST_ENV_NAME }}
-          python -m pytest -v --pyargs $MODULE_NAME
+          . "$CONDA/etc/profile.d/conda.sh"
+          conda activate "${{ env.TEST_ENV_NAME }}"
+          python -m pytest -v --pyargs "${{ env.MODULE_NAME }}"
 
   test_windows:
     needs: build_windows
@@ -266,7 +271,7 @@ jobs:
 
       - name: Display channels line
         run: |
-          echo ${{ env.CHANNELS }}
+          echo "${{ env.CHANNELS }}"
 
       - name: Download artifact
         uses: actions/download-artifact@3e5f45b2cfb9172054b4087a40e8e0b5a5461e7c # v8.0.1
@@ -366,8 +371,8 @@ jobs:
       - name: Report content of test environment
         shell: cmd /C CALL {0}
         run: |
-          echo "Value of CONDA enviroment variable was: " %CONDA%
-          echo "Value of CONDA_PREFIX enviroment variable was: " %CONDA_PREFIX%
+          echo "Value of CONDA environment variable was: " %CONDA%
+          echo "Value of CONDA_PREFIX environment variable was: " %CONDA_PREFIX%
           conda info && conda list -n ${{ env.TEST_ENV_NAME }}
 
       - name: Configure Intel OpenCL CPU RT
@@ -436,9 +441,9 @@ jobs:
       - name: Install anaconda-client
         run: conda install anaconda-client -c conda-forge --override-channels
       - name: Add conda to system path
-        run: echo $CONDA/bin >> $GITHUB_PATH
+        run: echo "$CONDA/bin" >> "$GITHUB_PATH"
       - name: Package version
-        run: echo "PACKAGE_VERSION=$(basename ${{ env.PACKAGE_NAME }}-*.conda | sed 's/^${{ env.PACKAGE_NAME }}-\([^-]*\).*/\1/')" >> $GITHUB_ENV
+        run: echo "PACKAGE_VERSION=$(basename ${{ env.PACKAGE_NAME }}-*.conda | sed 's/^${{ env.PACKAGE_NAME }}-\([^-]*\).*/\1/')" >> "$GITHUB_ENV"
 
       - name: Upload
         env:
@@ -447,14 +452,14 @@ jobs:
           ANACONDA_CLIENT_FORCE_STANDALONE: true
         if: ${{ env.ANACONDA_TOKEN != '' }}
         run: |
-          anaconda --token ${{ env.ANACONDA_TOKEN }} upload --user dppy --label dev ${PACKAGE_NAME}-*.conda
+          anaconda --token "${{ env.ANACONDA_TOKEN }}" upload --user dppy --label dev "${{ env.PACKAGE_NAME }}"-*.conda
 
       - name: Upload Wheels
         env:
           ANACONDA_TOKEN: ${{ secrets.ANACONDA_TOKEN }}
           # remove when anaconda-client 1.13.1 is available
           ANACONDA_CLIENT_FORCE_STANDALONE: true
-        run: anaconda --token ${{ env.ANACONDA_TOKEN }} upload --user dppy --label dev ${{ env.PACKAGE_NAME }}-*.whl --version ${{ env.PACKAGE_VERSION }}
+        run: anaconda --token "${{ env.ANACONDA_TOKEN }}" upload --user dppy --label dev "${{ env.PACKAGE_NAME }}"-*.whl --version "${{ env.PACKAGE_VERSION }}"
 
   upload_windows:
     needs: test_windows
@@ -491,18 +496,18 @@ jobs:
 
       - name: Package version
         shell: bash -el {0}
-        run: echo "PACKAGE_VERSION=$(basename ${{ env.PACKAGE_NAME }}-*.conda | sed 's/^${{ env.PACKAGE_NAME }}-\([^-]*\).*/\1/')" >> $GITHUB_ENV
+        run: echo "PACKAGE_VERSION=$(basename ${{ env.PACKAGE_NAME }}-*.conda | sed 's/^${{ env.PACKAGE_NAME }}-\([^-]*\).*/\1/')" >> "$GITHUB_ENV"
 
       - name: Upload
         env:
           ANACONDA_TOKEN: ${{ secrets.ANACONDA_TOKEN }}
         run: |
-          anaconda --token ${{ env.ANACONDA_TOKEN }} upload --user dppy --label dev ${{ env.PACKAGE_NAME }}-*.conda
+          anaconda --token "${{ env.ANACONDA_TOKEN }}" upload --user dppy --label dev "${{ env.PACKAGE_NAME }}"-*.conda
 
       - name: Upload Wheels
         env:
           ANACONDA_TOKEN: ${{ secrets.ANACONDA_TOKEN }}
-        run: anaconda --token ${{ env.ANACONDA_TOKEN }} upload --user dppy --label dev ${{ env.PACKAGE_NAME }}-*.whl --version ${{ env.PACKAGE_VERSION }}
+        run: anaconda --token "${{ env.ANACONDA_TOKEN }}" upload --user dppy --label dev "${{ env.PACKAGE_NAME }}"-*.whl --version "${{ env.PACKAGE_VERSION }}"
 
   test_examples_linux:
     needs: build_linux
@@ -521,7 +526,7 @@ jobs:
     steps:
       - name: Construct channels line
         run: |
-          echo "CHANNELS=-c ${{ env.INTEL_CHANNEL }} -c conda-forge --override-channels" >> $GITHUB_ENV
+          echo "CHANNELS=-c ${{ env.INTEL_CHANNEL }} -c conda-forge --override-channels" >> "$GITHUB_ENV"
       - name: Display channels line
         run: |
           echo ${{ env.CHANNELS }}
@@ -539,21 +544,23 @@ jobs:
         with:
           name: ${{ env.PACKAGE_NAME }} ${{ runner.os }} Python ${{ matrix.python }}
       - name: Add conda to system path
-        run: echo $CONDA/bin >> $GITHUB_PATH
+        run: echo "$CONDA/bin" >> "$GITHUB_PATH"
       - name: Create conda channel
         run: |
-          mkdir -p $GITHUB_WORKSPACE/channel/linux-64
-          conda index $GITHUB_WORKSPACE/channel || exit 1
-          mv ${PACKAGE_NAME}-*.conda $GITHUB_WORKSPACE/channel/linux-64 || exit 1
-          conda index $GITHUB_WORKSPACE/channel || exit 1
+          mkdir -p "$GITHUB_WORKSPACE/channel/linux-64"
+          conda index "$GITHUB_WORKSPACE/channel" || exit 1
+          mv "${{ env.PACKAGE_NAME }}"-*.conda "$GITHUB_WORKSPACE/channel/linux-64" || exit 1
+          conda index "$GITHUB_WORKSPACE/channel" || exit 1
           # Test channel
-          conda search $PACKAGE_NAME -c $GITHUB_WORKSPACE/channel --override-channels --info --json > $GITHUB_WORKSPACE/ver.json
+          conda search "$PACKAGE_NAME" -c "$GITHUB_WORKSPACE/channel" --override-channels --info --json > "$GITHUB_WORKSPACE/ver.json"
           cat ver.json
       - name: Collect dependencies
         run: |
           CHANNELS="-c $GITHUB_WORKSPACE/channel ${{ env.CHANNELS }}"
-          export PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
-          conda create -n ${{ env.EXAMPLES_ENV_NAME }} $PACKAGE_NAME=${PACKAGE_VERSION} python=${{ matrix.python }} $CHANNELS --only-deps --dry-run > lockfile
+          export CHANNELS
+          PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
+          export PACKAGE_VERSION
+          conda create -n "${{ env.EXAMPLES_ENV_NAME }}" "${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION}" python=${{ matrix.python }} "$CHANNELS" --only-deps --dry-run > lockfile
           cat lockfile
       - name: Set pkgs_dirs
         run: |
@@ -574,96 +581,93 @@ jobs:
         env:
           DPCPP_CMPLR: "dpcpp_linux-64>=2025.0"
         run: |
-          CHANNELS="${{ env.CHANNELS }}"
-          . $CONDA/etc/profile.d/conda.sh
+          export CHANNELS="${{ env.CHANNELS }}"
+          . "$CONDA/etc/profile.d/conda.sh"
           DPCTL_DEPENDS="$(python -c "${VER_SCRIPT1} ${VER_SCRIPT3}")"
           echo "Dpctl dependencies: ${DPCTL_DEPENDS}"
-          conda create -n ${{ env.EXAMPLES_ENV_NAME }} -y pytest python=${{ matrix.python }} "setuptools<72.2.0" $CHANNELS
+          conda create -n "${{ env.EXAMPLES_ENV_NAME }}" -y pytest python="${{ matrix.python }}" "setuptools<72.2.0" "$CHANNELS"
           echo "Environment created"
-          conda install -n ${{ env.EXAMPLES_ENV_NAME }} -y cmake ninja $CHANNELS || exit 1
+          conda install -n "${{ env.EXAMPLES_ENV_NAME }}" -y cmake ninja "$CHANNELS" || exit 1
           echo "Cmake and Ninja installed"
-          conda install -n ${{ env.EXAMPLES_ENV_NAME }} -y pybind11 cython scikit-build $CHANNELS || exit 1
+          conda install -n "${{ env.EXAMPLES_ENV_NAME }}" -y pybind11 cython scikit-build "$CHANNELS" || exit 1
           echo "scikit-build installed"
-          conda install -n ${{ env.EXAMPLES_ENV_NAME }} -y mkl-dpcpp \
+          conda install -n "${{ env.EXAMPLES_ENV_NAME }}" -y mkl-dpcpp \
               mkl-devel-dpcpp dpcpp_cpp_rt "${DPCTL_DEPENDS}" \
-              $CHANNELS || exit 1
+              "$CHANNELS" || exit 1
           echo "IPL installed"
-          conda create -y -n ${{ env.BUILD_ENV_NAME }} $CHANNELS gcc_linux-64 gxx_linux-64 \
-              ${{ env.DPCPP_CMPLR }} "${DPCTL_DEPENDS}" \
+          conda create -y -n "${{ env.BUILD_ENV_NAME }}" "$CHANNELS" gcc_linux-64 gxx_linux-64 \
+              "${{ env.DPCPP_CMPLR }}" "${DPCTL_DEPENDS}" \
               "sysroot_linux-64>=2.28"
           echo "Compiler installed"
-          conda list -n ${{ env.BUILD_ENV_NAME }}
+          conda list -n "${{ env.BUILD_ENV_NAME }}"
       - name: Install dpctl
         shell: bash -l {0}
         run: |
-          source $CONDA/etc/profile.d/conda.sh
+          source "$CONDA/etc/profile.d/conda.sh"
           CHANNELS="-c $GITHUB_WORKSPACE/channel -c dppy/label/dev -c ${{ env.INTEL_CHANNEL }} -c conda-forge --override-channels"
-          export PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
-          conda install -n ${{ env.EXAMPLES_ENV_NAME }} -y ${CHANNELS} dpctl=${PACKAGE_VERSION} dpnp || exit 1
+          export CHANNELS
+          PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
+          export PACKAGE_VERSION
+          conda install -n "${{ env.EXAMPLES_ENV_NAME }}" -y "${CHANNELS}" "${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION}" dpnp || exit 1
       - name: Build and run examples of pybind11 extensions
         shell: bash -l {0}
         run: |
-          source $CONDA/etc/profile.d/conda.sh
-          conda activate ${{ env.EXAMPLES_ENV_NAME }}
+          source "$CONDA/etc/profile.d/conda.sh"
+          conda activate "${{ env.EXAMPLES_ENV_NAME }}"
           conda list
           cd examples/pybind11
-          for d in $(find . -maxdepth 1 -type d -not -path ".")
-          do
-            pushd $d
+          while IFS= read -r d; do
+            pushd "$d" > /dev/null
             conda activate --stack build_env
             CC=icx CXX=icpx python setup.py build_ext --inplace -G Ninja || exit 1
             conda deactivate
-            if [ -e tests ]
-            then
-               LD_LIBRARY_PATH=${CONDA_PREFIX}/lib python -m pytest tests || exit 1
+            if [ -e tests ]; then
+               LD_LIBRARY_PATH="${CONDA_PREFIX}/lib" python -m pytest tests || exit 1
             else
-               LD_LIBRARY_PATH=${CONDA_PREFIX}/lib python example.py || exit 1
+               LD_LIBRARY_PATH="${CONDA_PREFIX}/lib" python example.py || exit 1
             fi
-            popd
-          done
+            popd > /dev/null
+          done < <(find . -maxdepth 1 -type d -not -path ".")
       - name: Build and run examples of Cython extensions
         shell: bash -l {0}
         run: |
-          source $CONDA/etc/profile.d/conda.sh
-          conda activate ${{ env.EXAMPLES_ENV_NAME }}
+          source "$CONDA/etc/profile.d/conda.sh"
+          conda activate "${{ env.EXAMPLES_ENV_NAME }}"
           conda list
           cd examples/cython
-          for d in $(find . -maxdepth 1 -type d -not -path ".")
-          do
-            pushd $d
-            conda activate --stack ${{ env.BUILD_ENV_NAME }}
+          while IFS= read -r d; do
+            pushd "$d" > /dev/null
+            conda activate --stack "${{ env.BUILD_ENV_NAME }}"
             CC=icx CXX=icpx python setup.py build_ext --inplace -G Ninja  || exit 1
             conda deactivate
             python -m pytest tests || exit 1
-            popd
-          done
+            popd > /dev/null
+          done < <(find . -maxdepth 1 -type d -not -path ".")
       - name: Build and run examples of C-extensions
         shell: bash -l {0}
         run: |
-          source $CONDA/etc/profile.d/conda.sh
-          conda activate ${{ env.EXAMPLES_ENV_NAME }}
+          source "$CONDA/etc/profile.d/conda.sh"
+          conda activate "${{ env.EXAMPLES_ENV_NAME }}"
           conda list
           cd examples/c
-          for d in $(find . -maxdepth 1 -type d -not -path ".")
-          do
-            pushd $d
-            conda activate --stack ${{ env.BUILD_ENV_NAME }}
+          while IFS= read -r d; do
+            pushd "$d" > /dev/null
+            conda activate --stack "${{ env.BUILD_ENV_NAME }}"
             python setup.py build_ext --inplace || exit 1
             conda deactivate
             python -m pytest tests || exit 1
-            popd
-          done
+            popd > /dev/null
+          done < <(find . -maxdepth 1 -type d -not -path ".")
       - name: Run Python examples
         shell: bash -l {0}
         run: |
           cd examples/python
-          source $CONDA/etc/profile.d/conda.sh
-          conda activate ${{ env.EXAMPLES_ENV_NAME }}
-          for script in $(find . \( -not -name "_*" -and -name "*.py" \))
-          do
+          source "$CONDA/etc/profile.d/conda.sh"
+          conda activate "${{ env.EXAMPLES_ENV_NAME }}"
+          while IFS= read -r script; do
             echo "Executing ${script}"
-            python ${script} || exit 1
-          done
+            python "${script}" || exit 1
+          done < <(find . \( -not -name "_*" -and -name "*.py" \))
 
   cleanup_packages:
     name: Clean up anaconda packages
diff --git a/.github/workflows/generate-coverage.yaml b/.github/workflows/generate-coverage.yaml
index 0032966518..d29cba85a3 100644
--- a/.github/workflows/generate-coverage.yaml
+++ b/.github/workflows/generate-coverage.yaml
@@ -111,15 +111,17 @@ jobs:
         shell: bash -l {0}
         run: |
           echo "Processing c-api-coverage"
-          export DPCTL_LCOV_FN=$(find _skbuild -name dpctl.lcov)
-          grep "/tmp" $DPCTL_LCOV_FN
+          DPCTL_LCOV_FN=$(find _skbuild -name dpctl.lcov)
+          export DPCTL_LCOV_FN
+          grep "/tmp" "$DPCTL_LCOV_FN"
           coveralls-lcov -v -n \
-              $DPCTL_LCOV_FN > dpctl-c-api-coverage.json
+              "$DPCTL_LCOV_FN" > dpctl-c-api-coverage.json
           echo "Processing pytest-coverage"
-          export DPCTL_PYTEST_LCOV=$(find . -name dpctl_pytest.lcov)
-          grep "/tmp" $DPCTL_PYTEST_LCOV
+          DPCTL_PYTEST_LCOV=$(find . -name dpctl_pytest.lcov)
+          export DPCTL_PYTEST_LCOV
+          grep "/tmp" "$DPCTL_PYTEST_LCOV"
           coveralls-lcov -v -n \
-              $DPCTL_PYTEST_LCOV > pytest-dpctl-c-api-coverage.json
+              "$DPCTL_PYTEST_LCOV" > pytest-dpctl-c-api-coverage.json
           echo "Merging JSON files"
           python -c "import json; \
               fh1 = open('dpctl-c-api-coverage.json', 'r'); \
@@ -132,7 +134,7 @@ jobs:
           # merge combined file with coverage data and upload
           ls -lh dpctl-c-api-coverage.json pytest-dpctl-c-api-coverage.json \
                combined-dpctl-c-api-coverage.json \
-               $(find _skbuild -name dpctl.lcov) $(find . -name dpctl_pytest.lcov)
+               "$DPCTL_LCOV_FN" "$DPCTL_PYTEST_LCOV"
           echo "Merging combined files with coverage data"
           coveralls --service=github --merge=combined-dpctl-c-api-coverage.json
         env:
diff --git a/.github/workflows/generate-docs.yml b/.github/workflows/generate-docs.yml
index 6295e9aa3b..cd47d7a3ea 100644
--- a/.github/workflows/generate-docs.yml
+++ b/.github/workflows/generate-docs.yml
@@ -75,7 +75,7 @@ jobs:
           source /opt/intel/oneapi/setvars.sh
           wget https://github.com/vovkos/doxyrest/releases/download/doxyrest-2.1.2/doxyrest-2.1.2-linux-amd64.tar.xz
           tar xf doxyrest-2.1.2-linux-amd64.tar.xz
-          python scripts/gen_docs.py --doxyrest-root=`pwd`/doxyrest-2.1.2-linux-amd64 --verbose --multiversion --clean || exit 1
+          python scripts/gen_docs.py --doxyrest-root="$(pwd)/doxyrest-2.1.2-linux-amd64" --verbose --multiversion --clean || exit 1
           python -c "import dpctl; print(dpctl.__version__)" || exit 1
           mv "$(find _skbuild -type d -path "*/cmake-install/docs/docs" | head -n 1)" ~/docs
           git clean -dfx
@@ -87,9 +87,9 @@ jobs:
           git remote add tokened_docs https://IntelPython:${{ secrets.GITHUB_TOKEN }}@github.com/IntelPython/dpctl.git
           git fetch tokened_docs
           git checkout --track tokened_docs/gh-pages
-          echo `pwd`
+          pwd
           cd master
-          git rm -rf *
+          git rm -rf ./*
           mv ~/docs/* . || exit 1
           git add .
           git config --global user.name 'github-actions[doc-deploy-bot]'
@@ -112,10 +112,10 @@ jobs:
           git remote add tokened_docs https://IntelPython:${{ secrets.GITHUB_TOKEN }}@github.com/IntelPython/dpctl.git
           git fetch tokened_docs
           git checkout --track tokened_docs/gh-pages
-          echo `pwd`
-          [ -d pulls/${PR_NUM} ] && git rm -rf pulls/${PR_NUM}
-          mkdir -p pulls/${PR_NUM}
-          cd pulls/${PR_NUM}
+          pwd
+          [ -d "pulls/${PR_NUM}" ] && git rm -rf "pulls/${PR_NUM}"
+          mkdir -p "pulls/${PR_NUM}"
+          cd "pulls/${PR_NUM}"
           mv ~/docs/* .
           git add .
           git config --global user.name 'github-actions[doc-deploy-bot]'
@@ -131,10 +131,10 @@ jobs:
           git remote add tokened_docs https://IntelPython:${{ secrets.GITHUB_TOKEN }}@github.com/IntelPython/dpctl.git
           git fetch tokened_docs
           git checkout --track tokened_docs/gh-pages
-          echo `pwd`
+          pwd
           ls
           [ -d pulls ] && ls pulls && echo "This is pull/${PR_NUM}"
-          [ -d pulls/${PR_NUM} ] && git rm -rf pulls/${PR_NUM}
+          [ -d "pulls/${PR_NUM}" ] && git rm -rf "pulls/${PR_NUM}"
           git config --global user.name 'github-actions[doc-deploy-bot]'
           git config --global user.email 'github-actions[doc-deploy-bot]@users.noreply.github.com'
           git commit -m "Removing docs for closed pull request ${PR_NUM}"
diff --git a/.github/workflows/os-llvm-sycl-build.yml b/.github/workflows/os-llvm-sycl-build.yml
index 571b651524..b0bbe9f2df 100644
--- a/.github/workflows/os-llvm-sycl-build.yml
+++ b/.github/workflows/os-llvm-sycl-build.yml
@@ -59,28 +59,35 @@ jobs:
           cd sycl_bundle
           if [[ "${USE_LATEST_SYCLOS:-0}" -eq "1" ]]; then
               # get list of shas and tags from remote, filter nightly tags and reverse order
-              export LLVM_TAGS=$(git -c 'versionsort.suffix=-' ls-remote --tags --sort='v:refname' https://github.com/intel/llvm.git | \
+              LLVM_TAGS=$(git -c 'versionsort.suffix=-' ls-remote --tags --sort='v:refname' https://github.com/intel/llvm.git | \
                           grep 'refs/tags/nightly-' | awk '{a[i++]=$0} END {for (j=i-1; j>=0;) print a[j--] }')
+              export LLVM_TAGS
               # initialize
               unset DEPLOY_NIGHTLY_TAG
               unset DEPLOY_NIGHTLY_TAG_SHA
 
               # go through tags and find the most recent one where nighly build binary is available
               while IFS= read -r NEXT_LLVM_TAG; do
-                  export NEXT_LLVM_TAG_SHA=$(echo ${NEXT_LLVM_TAG} | awk '{print $1}')
-                  export NEXT_NIGHTLY_TAG=$(python3 -c "import sys, urllib.parse as ul; print (ul.quote_plus(sys.argv[1]))" \
-                                              $(echo ${NEXT_LLVM_TAG} | awk '{gsub(/^refs\/tags\//, "", $2)} {print $2}'))
-                  if [[ `wget -S --spider ${DOWNLOAD_URL_PREFIX}/${NEXT_NIGHTLY_TAG}/${ARTIFACT_NAME}.tar.gz  2>&1 | grep 'HTTP/1.1 200 OK'` ]];
+                  NEXT_LLVM_TAG_SHA=$(echo "${NEXT_LLVM_TAG}" | awk '{print $1}')
+                  export NEXT_LLVM_TAG_SHA
+                  NEXT_NIGHTLY_TAG=$(python3 -c "import sys, urllib.parse as ul; print(ul.quote_plus(sys.argv[1]))" \
+                                      "$(echo "${NEXT_LLVM_TAG}" | awk '{gsub(/^refs\/tags\//, "", $2)} {print $2}')")
+                  export NEXT_NIGHTLY_TAG
+                  if wget -S --spider "${DOWNLOAD_URL_PREFIX}/${NEXT_NIGHTLY_TAG}/${ARTIFACT_NAME}.tar.gz" 2>&1 | grep -q 'HTTP/1.1 200 OK';
                   then
-                      export DEPLOY_NIGHTLY_TAG=${NEXT_NIGHTLY_TAG}
-                      export DEPLOY_LLVM_TAG_SHA=${NEXT_LLVM_TAG_SHA}
+                      DEPLOY_NIGHTLY_TAG="${NEXT_NIGHTLY_TAG}"
+                      export DEPLOY_NIGHTLY_TAG
+                      DEPLOY_LLVM_TAG_SHA="${NEXT_LLVM_TAG_SHA}"
+                      export DEPLOY_LLVM_TAG_SHA
                       break
                   fi
               done <<< "${LLVM_TAGS}"
           else
               # Use latest known to work tag instead
-              export DEPLOY_NIGHTLY_TAG="sycl-nightly%2F20230606"
-              export DEPLOY_LLVM_TAG_SHA=f44d0133d4b0077298f034697a1f3818ff1d6134
+              DEPLOY_NIGHTLY_TAG="sycl-nightly%2F20230606"
+              export DEPLOY_NIGHTLY_TAG
+              DEPLOY_LLVM_TAG_SHA=f44d0133d4b0077298f034697a1f3818ff1d6134
+              export DEPLOY_LLVM_TAG_SHA
           fi
 
           [[ -n "${DEPLOY_NIGHTLY_TAG}" ]] || exit 1
@@ -90,16 +97,20 @@ jobs:
           if [[ -f bundle_id.txt && ( "$(cat bundle_id.txt)" == "${DEPLOY_LLVM_TAG_SHA}" ) ]]; then
               echo "Using cached download of ${DEPLOY_LLVM_TAG_SHA}"
           else
-              rm -rf ${ARTIFACT_NAME}.tar.gz
-              wget ${DOWNLOAD_URL_PREFIX}/${DEPLOY_NIGHTLY_TAG}/${ARTIFACT_NAME}.tar.gz && echo ${DEPLOY_LLVM_TAG_SHA} > bundle_id.txt || rm -rf bundle_id.txt
-              [ -f ${OCLCPUEXP_FN} ] || wget ${DOWNLOAD_URL_PREFIX}/${DRIVER_PATH}/${OCLCPUEXP_FN} || rm -rf bundle_id.txt
-              [ -f ${TBB_FN} ] || wget ${TBB_URL}/${TBB_FN} || rm -rf bundle_id.txt
+              rm -rf "${ARTIFACT_NAME}.tar.gz"
+              if wget "${DOWNLOAD_URL_PREFIX}/${DEPLOY_NIGHTLY_TAG}/${ARTIFACT_NAME}.tar.gz"; then
+                echo "${DEPLOY_LLVM_TAG_SHA}" > bundle_id.txt
+              else
+                rm -rf bundle_id.txt
+              fi
+              [ -f "${OCLCPUEXP_FN}" ] || wget "${DOWNLOAD_URL_PREFIX}/${DRIVER_PATH}/${OCLCPUEXP_FN}" || rm -rf bundle_id.txt
+              [ -f "${TBB_FN}" ] || wget "${TBB_URL}/${TBB_FN}" || rm -rf bundle_id.txt
               rm -rf dpcpp_compiler
               mkdir -p dpcpp_compiler
-              tar xf ${ARTIFACT_NAME}.tar.gz -C dpcpp_compiler
+              tar xf "${ARTIFACT_NAME}.tar.gz" -C dpcpp_compiler
               mkdir -p oclcpuexp
-              [ -d oclcpuexp/x64 ] || tar xf ${OCLCPUEXP_FN} -C oclcpuexp
-              [ -d ${TBB_INSTALL_DIR}/lib ] || tar xf ${TBB_FN}
+              [ -d oclcpuexp/x64 ] || tar xf "${OCLCPUEXP_FN}" -C oclcpuexp
+              [ -d "${TBB_INSTALL_DIR}/lib" ] || tar xf "${TBB_FN}"
               cp oclcpuexp/x64/libOpenCL.so* dpcpp_compiler/lib/
           fi
 
@@ -152,7 +163,7 @@ jobs:
         run: |
           source set_allvars.sh
           python scripts/build_locally.py --c-compiler=clang --cxx-compiler=clang++ \
-              --compiler-root=${SYCL_BUNDLE_FOLDER}/dpcpp_compiler/bin || exit 1
+              --compiler-root="${SYCL_BUNDLE_FOLDER}/dpcpp_compiler/bin" || exit 1
 
       - name: Run lsplatforms
         shell: bash -l {0}
diff --git a/.github/workflows/run-tests-from-dppy-bits.yaml b/.github/workflows/run-tests-from-dppy-bits.yaml
index 3a883cf1f1..9172902b51 100644
--- a/.github/workflows/run-tests-from-dppy-bits.yaml
+++ b/.github/workflows/run-tests-from-dppy-bits.yaml
@@ -35,11 +35,11 @@ jobs:
     steps:
       - name: Construct channels line
         run: |
-          echo "CHANNELS=-c ${{ env.INTEL_CHANNEL }} -c conda-forge --override-channels" >> $GITHUB_ENV
+          echo "CHANNELS=-c ${{ env.INTEL_CHANNEL }} -c conda-forge --override-channels" >> "$GITHUB_ENV"
 
       - name: Display channels line
         run: |
-          echo ${{ env.CHANNELS }}
+          echo "${{ env.CHANNELS }}"
 
       - name: Set pkgs_dirs
         run: |
@@ -47,27 +47,27 @@ jobs:
 
       - name: Install dpctl
         run: |
-          conda create -n ${{ env.TEST_ENV_NAME }} -c dppy/label/dev ${{ env.CHANNELS }} dpctl pytest pytest-cov cython setuptools c-compiler cxx-compiler
+          conda create -n "${{ env.TEST_ENV_NAME }}" -c dppy/label/dev "${{ env.CHANNELS }}" dpctl pytest pytest-cov cython setuptools c-compiler cxx-compiler
 
       - name: Smoke test
         run: |
-          . $CONDA/etc/profile.d/conda.sh
-          conda activate ${{ env.TEST_ENV_NAME }}
+          . "$CONDA/etc/profile.d/conda.sh"
+          conda activate "${{ env.TEST_ENV_NAME }}"
           python -m dpctl -f
 
       - name: Create test temp dir
         # create temporary empty folder to runs tests from
         # https://github.com/pytest-dev/pytest/issues/11904
-        run: mkdir -p ${GITHUB_WORKSPACE}/test_tmp
+        run: mkdir -p "${GITHUB_WORKSPACE}/test_tmp"
 
       - name: Run tests
         working-directory: ${{ github.workspace }}/test_tmp
         env:
           SYCL_CACHE_PERSISTENT: 1
         run: |
-          . $CONDA/etc/profile.d/conda.sh
-          conda activate ${{ env.TEST_ENV_NAME }}
-          python -m pytest -v --pyargs $MODULE_NAME
+          . "$CONDA/etc/profile.d/conda.sh"
+          conda activate "${{ env.TEST_ENV_NAME }}"
+          python -m pytest -v --pyargs "${{ env.MODULE_NAME }}"
 
 
   test_windows:
@@ -94,7 +94,7 @@ jobs:
 
       - name: Display channels line
         run: |
-          echo ${{ env.CHANNELS }}
+          echo "${{ env.CHANNELS }}"
 
       - uses: conda-incubator/setup-miniconda@fc2d68f6413eb2d87b895e92f8584b5b94a10167 # v3.3.0
         with:
diff --git a/README.md b/README.md
index 428fd0410f..04507fe063 100644
--- a/README.md
+++ b/README.md
@@ -57,7 +57,7 @@ To get the library from the latest oneAPI release, follow the
 instructions from Intel(R) [oneAPI installation
 guide](https://www.intel.com/content/www/us/en/developer/articles/guide/installation-guide-for-oneapi-toolkits.html).
 
-> **NOTE:** You need to install the Intel(R) oneAPI AI Analytics Tookit to get
+> **NOTE:** You need to install the Intel(R) oneAPI AI Analytics Toolkit to get
 >IDP and `dpctl`.
 
 
diff --git a/cmake/IntelSYCLConfig.cmake b/cmake/IntelSYCLConfig.cmake
old mode 100755
new mode 100644
index c51e47290c..7a4721dd6a
--- a/cmake/IntelSYCLConfig.cmake
+++ b/cmake/IntelSYCLConfig.cmake
@@ -18,7 +18,7 @@
 IntelSYCLConfig
 -------
 
-Library to verify SYCL compatability of CMAKE_CXX_COMPILER
+Library to verify SYCL compatibility of CMAKE_CXX_COMPILER
 and passes relevant compiler flags.
 
 Result Variables
@@ -290,7 +290,7 @@ if(nosycllang)
   set(IntelSYCL_NOT_FOUND_MESSAGE "${SYCL_REASON_FAILURE}")
 endif()
 
-# Placeholder for identifying various implemenations of SYCL compilers.
+# Placeholder for identifying various implementations of SYCL compilers.
 # for now, set to the CMAKE_CXX_COMPILER_ID
 set(SYCL_IMPLEMENTATION_ID "${CMAKE_CXX_COMPILER_ID}")
 
diff --git a/conda-recipe/build.sh b/conda-recipe/build.sh
index 1ca6329e3a..e892602ede 100755
--- a/conda-recipe/build.sh
+++ b/conda-recipe/build.sh
@@ -5,8 +5,10 @@ export LIBRARY_PATH="$LIBRARY_PATH:${BUILD_PREFIX}/lib"
 
 # Intel LLVM must cooperate with compiler and sysroot from conda
 echo "--gcc-toolchain=${BUILD_PREFIX} --sysroot=${BUILD_PREFIX}/${HOST}/sysroot -target ${HOST}" > icpx_for_conda.cfg
-export ICPXCFG="$(pwd)/icpx_for_conda.cfg"
-export ICXCFG="$(pwd)/icpx_for_conda.cfg"
+ICPXCFG="$(pwd)/icpx_for_conda.cfg"
+export ICPXCFG
+ICXCFG="$(pwd)/icpx_for_conda.cfg"
+export ICXCFG
 
 read -r GLIBC_MAJOR GLIBC_MINOR <<<"$(conda list '^sysroot_linux-64$' \
     | tail -n 1 | awk '{print $2}' | grep -oP '\d+' | head -n 2 | tr '\n' ' ')"
diff --git a/docs/CMakeLists.txt b/docs/CMakeLists.txt
index d2053d106c..0bd31e059c 100644
--- a/docs/CMakeLists.txt
+++ b/docs/CMakeLists.txt
@@ -4,10 +4,10 @@ option(DPCTL_ENABLE_DOXYREST
     OFF
 )
 
-# Option to add verion links to the side bar. This option is primarily
+# Option to add version links to the side bar. This option is primarily
 # intended to generate dpctl's docs for our github.io page.
 option(DPCTL_USE_MULTIVERSION_TEMPLATE
-    "Enable adding verion links to side bar"
+    "Enable adding version links to side bar"
     OFF
 )
 
diff --git a/docs/_legacy/docfiles/user_guides/QuickStart.rst b/docs/_legacy/docfiles/user_guides/QuickStart.rst
index 57c443e543..382528cc98 100644
--- a/docs/_legacy/docfiles/user_guides/QuickStart.rst
+++ b/docs/_legacy/docfiles/user_guides/QuickStart.rst
@@ -262,7 +262,7 @@ library.
     pushd build || exit 1
 
     INSTALL_PREFIX=$(pwd)/../install
-    rm -rf ${INSTALL_PREFIX}
+    rm -rf "${INSTALL_PREFIX}"
     export ONEAPI_ROOT=/opt/intel/oneapi
     DPCPP_ROOT=${ONEAPI_ROOT}/compiler/latest/linux
 
diff --git a/docs/doc_sources/_static/dpctl.svg b/docs/doc_sources/_static/dpctl.svg
old mode 100755
new mode 100644
diff --git a/docs/doc_sources/api_reference/dpctl/filter_selector_string.rst b/docs/doc_sources/api_reference/dpctl/filter_selector_string.rst
index b31edbf789..f96a8764df 100644
--- a/docs/doc_sources/api_reference/dpctl/filter_selector_string.rst
+++ b/docs/doc_sources/api_reference/dpctl/filter_selector_string.rst
@@ -24,7 +24,7 @@ least one component.
 ``Backend`` specifies the desired backend of targeted devices, while
 ``DeviceType`` specifies the type of targeted devices.
 ``RelativeDeviceNumber`` refers to the number of the device that matches
-any other given requirements, starting from `0` to marking the
+any other given requirements, starting from 0 to marking the
 "first device that matches the requirements".
 
 Attempting to use a non-conforming string in places where filter selector
diff --git a/docs/doc_sources/api_reference/dpctl_cmake.rst b/docs/doc_sources/api_reference/dpctl_cmake.rst
index 62a9135e5f..e72b5dc23c 100644
--- a/docs/doc_sources/api_reference/dpctl_cmake.rst
+++ b/docs/doc_sources/api_reference/dpctl_cmake.rst
@@ -3,7 +3,7 @@
 CMake support
 =============
 
-:py:mod:`dpctl` comes with the configuration file `dpctl-config.cmake` which is installed
+:py:mod:`dpctl` comes with the configuration file :file:`dpctl-config.cmake` which is installed
 on the `standard search path CMake uses to search for packages <cmake_find_package_search_paths_>`_.
 
 To build your extension that leverages :py:mod:`dpctl`, include the following line in your CMake script:
diff --git a/docs/doc_sources/beginners_guides/installation.rst b/docs/doc_sources/beginners_guides/installation.rst
index 75fab31e02..c3d41bf658 100644
--- a/docs/doc_sources/beginners_guides/installation.rst
+++ b/docs/doc_sources/beginners_guides/installation.rst
@@ -211,7 +211,7 @@ AMD build
 
     python scripts/build_locally.py --verbose --target-hip=<arch>
 
-Note that the `oneAPI for AMD GPUs` plugin requires the architecture be specified and only
+Note that the oneAPI for AMD GPUs plugin requires the architecture be specified and only
 one architecture can be specified at a time.
 
 To determine the architecture code (``<arch>``) for your AMD GPU, run:
diff --git a/docs/doc_sources/beginners_guides/managing_devices.rst b/docs/doc_sources/beginners_guides/managing_devices.rst
index 4e7ec12cca..9389f47fc3 100644
--- a/docs/doc_sources/beginners_guides/managing_devices.rst
+++ b/docs/doc_sources/beginners_guides/managing_devices.rst
@@ -310,7 +310,7 @@ For Intel GPU devices, additional architectural information can be access with :
 
     In [2]: d_gpu = dpctl.SyclDevice()
 
-    # Output for Iris Xe integerate GPU, with PCI ID 0x9a49
+    # Output for Iris Xe integrate GPU, with PCI ID 0x9a49
     # (corresponding decimal value: 39497)
     In [3]: dpctl.utils.intel_device_info(d_gpu)
     Out[3]:
diff --git a/docs/doc_sources/contributor_guides/building.rst b/docs/doc_sources/contributor_guides/building.rst
index f1681c6c6a..df2a8c35b2 100644
--- a/docs/doc_sources/contributor_guides/building.rst
+++ b/docs/doc_sources/contributor_guides/building.rst
@@ -209,7 +209,7 @@ library.
     pushd build || exit 1
 
     INSTALL_PREFIX=$(pwd)/../install
-    rm -rf ${INSTALL_PREFIX}
+    rm -rf "${INSTALL_PREFIX}"
     export ONEAPI_ROOT=/opt/intel/oneapi
     # Values are set as appropriate for oneAPI DPC++ 2024.0
     # or later.
diff --git a/docs/doc_sources/extlinks_gen.py b/docs/doc_sources/extlinks_gen.py
index 6e61624cda..789e3a433a 100644
--- a/docs/doc_sources/extlinks_gen.py
+++ b/docs/doc_sources/extlinks_gen.py
@@ -19,7 +19,7 @@
 
 def create_extlinks():
     """Reads a JSON file to create a dictionary of urls in the format supported
-    by the sphinx.ect.extlinks extension.
+    by the sphinx.ext.extlinks extension.
 
     Returns:
         dict: A dictionary that is understood by the extlinks Sphinx extension.
diff --git a/docs/doc_sources/index.rst b/docs/doc_sources/index.rst
index e69559f7f6..3923b716b5 100644
--- a/docs/doc_sources/index.rst
+++ b/docs/doc_sources/index.rst
@@ -56,7 +56,7 @@ Intel(R) oneAPI :dpcpp_compiler:`DPC++ compiler <>`.
 
             Access API Reference
 
-    .. grid-item-card:: Contibutor Guides
+    .. grid-item-card:: Contributor Guides
 
         The contributing guidelines will suggest a process of
         contributing to :mod:`dpctl`.
diff --git a/docs/doc_sources/user_guides/basic_concepts.rst b/docs/doc_sources/user_guides/basic_concepts.rst
index d2438e4a38..ebec1d6377 100644
--- a/docs/doc_sources/user_guides/basic_concepts.rst
+++ b/docs/doc_sources/user_guides/basic_concepts.rst
@@ -111,7 +111,7 @@ backend. The context is required to map unified address space pointer to the dev
 where it was allocated unambiguously.
 
 In order for two DPC++-based Python extensions to share USM allocations, they each
-must use the `same` SYCL context when submitting for execution programs that would
+must use the *same* SYCL context when submitting for execution programs that would
 access this allocation.
 
 Since ``sycl::context`` is dynamically constructed by each extension  sharing a USM allocation,
diff --git a/dpctl/_host_task_util.hpp b/dpctl/_host_task_util.hpp
index 4db07e1cb1..5a2ea19645 100644
--- a/dpctl/_host_task_util.hpp
+++ b/dpctl/_host_task_util.hpp
@@ -65,7 +65,7 @@ DPCTLSyclEventRef async_dec_ref(DPCTLSyclQueueRef QRef,
 #else
                 const bool finalizing = Py_IsFinalizing();
 #endif
-                // if the main thread has not finilized the interpreter yet
+                // if the main thread has not finalized the interpreter yet
                 if (initialized && !finalizing) {
                     PyGILState_STATE gstate;
                     gstate = PyGILState_Ensure();
diff --git a/dpctl/_sycl_device.pyx b/dpctl/_sycl_device.pyx
index b887c313c2..6c6638e2e3 100644
--- a/dpctl/_sycl_device.pyx
+++ b/dpctl/_sycl_device.pyx
@@ -1210,7 +1210,7 @@ cdef class SyclDevice(_SyclDevice):
         sg_sizes = DPCTLDevice_GetSubGroupSizes(
             self._device_ref, &sg_sizes_len)
         if (sg_sizes is not NULL and sg_sizes_len > 0):
-            res = list()
+            res = []
             for i in range(sg_sizes_len):
                 res.append(sg_sizes[i])
             DPCTLSize_t_Array_Delete(sg_sizes)
@@ -1691,7 +1691,7 @@ cdef class SyclDevice(_SyclDevice):
         return _get_devices(DVRef)
 
     def create_sub_devices(self, **kwargs):
-        """create_sub_devices(partition=parition_spec)
+        """create_sub_devices(partition=partition_spec)
         Creates a list of sub-devices by partitioning a root device based on the
         provided partition specifier.
 
@@ -1755,12 +1755,12 @@ cdef class SyclDevice(_SyclDevice):
         """
         if "partition" not in kwargs:
             raise TypeError(
-                "create_sub_devices(partition=parition_spec) is expected."
+                "create_sub_devices(partition=partition_spec) is expected."
             )
         partition = kwargs.pop("partition")
         if kwargs:
             raise TypeError(
-                "create_sub_devices(partition=parition_spec) is expected."
+                "create_sub_devices(partition=partition_spec) is expected."
             )
         if isinstance(partition, int) and partition >= 0:
             return self.create_sub_devices_equally(partition)
diff --git a/dpctl/_sycl_device_factory.pyx b/dpctl/_sycl_device_factory.pyx
index 1689249bc9..7a9794ede0 100644
--- a/dpctl/_sycl_device_factory.pyx
+++ b/dpctl/_sycl_device_factory.pyx
@@ -416,7 +416,7 @@ cdef class _DefaultDeviceCache:
     cdef dict __device_map__
 
     def __cinit__(self):
-        self.__device_map__ = dict()
+        self.__device_map__ = {}
 
     cdef get_or_create(self):
         """Return instance of SyclDevice and indicator if cache
diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx
index b0b8336788..74b3b49d71 100644
--- a/dpctl/_sycl_queue.pyx
+++ b/dpctl/_sycl_queue.pyx
@@ -1786,7 +1786,7 @@ cdef class RawKernelArg:
             Variadic argument, see class documentation.
 
     Raises:
-        TypeError: In case of incorrect arguments given to constructurs,
+        TypeError: In case of incorrect arguments given to constructors,
                    unexpected types of input arguments.
     """
     def __cinit__(self, *args):
diff --git a/dpctl/_sycl_queue_manager.pyx b/dpctl/_sycl_queue_manager.pyx
index bf720e9402..9638ed211f 100644
--- a/dpctl/_sycl_queue_manager.pyx
+++ b/dpctl/_sycl_queue_manager.pyx
@@ -35,7 +35,7 @@ cdef class _DeviceDefaultQueueCache:
     cdef dict __device_queue_map__
 
     def __cinit__(self):
-        self.__device_queue_map__ = dict()
+        self.__device_queue_map__ = {}
 
     def get_or_create(self, key):
         """Return instance of SyclQueue and indicator if cache
diff --git a/dpctl/_sycl_timer.py b/dpctl/_sycl_timer.py
index 09e668281a..9acca40747 100644
--- a/dpctl/_sycl_timer.py
+++ b/dpctl/_sycl_timer.py
@@ -190,7 +190,7 @@ def __init__(
         self.queue = None
         self.host_times = []
         self.bracketing_events = []
-        self._context_data = list()
+        self._context_data = []
         if device_timer is None:
             device_timer = "queue_barrier"
         if device_timer == "queue_barrier":
diff --git a/dpctl/apis/include/dpctl4pybind11.hpp b/dpctl/apis/include/dpctl4pybind11.hpp
index 6e1a9fe3da..416227d9ac 100644
--- a/dpctl/apis/include/dpctl4pybind11.hpp
+++ b/dpctl/apis/include/dpctl4pybind11.hpp
@@ -157,11 +157,11 @@ class dpctl_capi
         {
             const bool initialized = Py_IsInitialized();
 #if PY_VERSION_HEX < 0x30d0000
-            const bool finilizing = _Py_IsFinalizing();
+            const bool finalizing = _Py_IsFinalizing();
 #else
-            const bool finilizing = Py_IsFinalizing();
+            const bool finalizing = Py_IsFinalizing();
 #endif
-            const bool guard = initialized && !finilizing;
+            const bool guard = initialized && !finalizing;
 
             if (guard) {
                 delete p;
diff --git a/dpctl/tests/helper/_helper.py b/dpctl/tests/helper/_helper.py
index 475669976c..c79fddd305 100644
--- a/dpctl/tests/helper/_helper.py
+++ b/dpctl/tests/helper/_helper.py
@@ -43,7 +43,7 @@ def create_invalid_capsule():
     return ctor(id(ctor), b"invalid", 0)
 
 
-def get_queue_or_skip(args=tuple()):
+def get_queue_or_skip(args=()):
     try:
         q = dpctl.SyclQueue(*args)
     except dpctl.SyclQueueCreationError:
diff --git a/dpctl/tests/test_sycl_device.py b/dpctl/tests/test_sycl_device.py
index c8828ee68f..e9eee7e223 100644
--- a/dpctl/tests/test_sycl_device.py
+++ b/dpctl/tests/test_sycl_device.py
@@ -192,7 +192,7 @@ def test_supported_aspect(supported_aspect):
         assert getattr(d_wa, "has_aspect_" + supported_aspect)
     except dpctl.SyclDeviceCreationError:
         # ValueError may be raised if no device with
-        # requested aspect charateristics is available
+        # requested aspect characteristics is available
         assert not has_it
 
 
@@ -405,7 +405,7 @@ def test_peer_device_arg_validation(method):
         dev = dpctl.SyclDevice()
     except dpctl.SyclDeviceCreationError:
         pytest.skip("No default device available")
-    bad_dev = dict()
+    bad_dev = {}
     callable = getattr(dev, method)
     with pytest.raises(TypeError):
         callable(bad_dev)
diff --git a/dpctl/tests/test_sycl_event.py b/dpctl/tests/test_sycl_event.py
index 2985bf3af0..584c8932da 100644
--- a/dpctl/tests/test_sycl_event.py
+++ b/dpctl/tests/test_sycl_event.py
@@ -70,7 +70,7 @@ def test_create_event_from_capsule():
 
 def test_invalid_constructor_arg():
     with pytest.raises(TypeError):
-        dpctl.SyclEvent(list())
+        dpctl.SyclEvent([])
 
 
 def test_wait_with_event():
diff --git a/dpctl/tests/test_sycl_queue.py b/dpctl/tests/test_sycl_queue.py
index 61af60da34..613a13adb6 100644
--- a/dpctl/tests/test_sycl_queue.py
+++ b/dpctl/tests/test_sycl_queue.py
@@ -209,7 +209,7 @@ def test_queue_invalid_property():
     with pytest.raises(ValueError):
         dpctl.SyclQueue(property=4.5)
     with pytest.raises(ValueError):
-        dpctl.SyclQueue(property=["abc", tuple()])
+        dpctl.SyclQueue(property=["abc", ()])
 
 
 def test_queue_capsule():
@@ -345,13 +345,13 @@ def test_queue_memops():
     q.prefetch(m1, 512)
     q.mem_advise(m1, 512, 0)
     with pytest.raises(TypeError):
-        q.memcpy(m1, list(), 512)
+        q.memcpy(m1, [], 512)
     with pytest.raises(TypeError):
-        q.memcpy(list(), m2, 512)
+        q.memcpy([], m2, 512)
     with pytest.raises(TypeError):
-        q.prefetch(list(), 512)
+        q.prefetch([], 512)
     with pytest.raises(TypeError):
-        q.mem_advise(list(), 512, 0)
+        q.mem_advise([], 512, 0)
 
 
 @pytest.fixture(scope="session")
diff --git a/dpctl/tests/test_sycl_usm.py b/dpctl/tests/test_sycl_usm.py
index c17951f993..b1aefa325c 100644
--- a/dpctl/tests/test_sycl_usm.py
+++ b/dpctl/tests/test_sycl_usm.py
@@ -101,7 +101,7 @@ def test_memory_without_context():
 def test_memory_cpu_context():
     mobj = _create_memory()
 
-    # type respective to the context in which
+    # USM type respective to the context in which
     # memory was created
     usm_type = mobj.get_usm_type()
     assert usm_type == "shared"
@@ -110,9 +110,9 @@ def test_memory_cpu_context():
         cpu_queue = dpctl.SyclQueue("cpu")
     except dpctl.SyclQueueCreationError:
         pytest.skip("SyclQueue('cpu') failed, skip further testing")
-    # type as view from CPU queue
+    # USM type as view from CPU queue
     usm_type = mobj.get_usm_type(cpu_queue)
-    # type can be unknown if current queue is
+    # USM type can be unknown if current queue is
     # not in the same SYCL context
     assert usm_type in ["unknown", "shared"]
 
@@ -265,7 +265,7 @@ def test_usm_type_exceptions():
     with pytest.raises(TypeError):
         m.get_usm_type(syclobj=Ellipsis)
     with pytest.raises(TypeError):
-        m.get_usm_type_enum(syclobj=list())
+        m.get_usm_type_enum(syclobj=[])
 
 
 def test_sycl_usm_array_interface(memory_ctor):
@@ -419,7 +419,7 @@ def invalid_version(suai_iface):
     # data validation
     def invalid_data(suai_iface):
         "Set data to invalid"
-        suai_iface["data"] = tuple()
+        suai_iface["data"] = ()
         return suai_iface
 
     v = View(
@@ -439,7 +439,7 @@ def invalid_data(suai_iface):
     with pytest.raises(ValueError):
         MemoryUSMShared(v)
 
-    # typestr validation
+    # validate typestring
     def invalid_typestr(suai_iface):
         suai_iface["typestr"] = "invalid"
         return suai_iface
diff --git a/dpctl/tests/test_utils.py b/dpctl/tests/test_utils.py
index 2114346dbc..472edd2ff2 100644
--- a/dpctl/tests/test_utils.py
+++ b/dpctl/tests/test_utils.py
@@ -61,7 +61,7 @@ def test_intel_device_info():
 
 
 def test_intel_device_info_validation():
-    invalid_device = dict()
+    invalid_device = {}
     with pytest.raises(TypeError):
         dpctl.utils.intel_device_info(invalid_device)
 
diff --git a/dpctl/utils/CMakeLists.txt b/dpctl/utils/CMakeLists.txt
index 40569da8a4..8f27fb0566 100644
--- a/dpctl/utils/CMakeLists.txt
+++ b/dpctl/utils/CMakeLists.txt
@@ -57,7 +57,7 @@ foreach(python_module_name ${_pybind11_targets})
             ${_dpctl_sycl_target_link_options}
         )
     endif()
-    # TODO: update source so they refernece individual libraries instead of
+    # TODO: update source so they reference individual libraries instead of
     #   dpctl4pybind11.hpp. It will allow to simplify dependency tree
     target_link_libraries(${python_module_name} PRIVATE DpctlCAPI)
     if (DPCTL_WITH_REDIST)
diff --git a/dpctl/utils/_intel_device_info.py b/dpctl/utils/_intel_device_info.py
index f4de00a8f4..c25d4575db 100644
--- a/dpctl/utils/_intel_device_info.py
+++ b/dpctl/utils/_intel_device_info.py
@@ -105,7 +105,7 @@ def intel_device_info(dev, /):
         if mbw:
             res["memory_bus_width"] = mbw
         return res
-    return dict()
+    return {}
 
 
 __all__ = [
diff --git a/examples/cython/README.md b/examples/cython/README.md
index 010e871867..e2daa39be6 100644
--- a/examples/cython/README.md
+++ b/examples/cython/README.md
@@ -4,5 +4,5 @@ The `dpctl` package provides Cython definition files for types it defines.
 
 Use `cimport dpctl as c_dpctl` or `cimport dpctl.memory as c_dpm` to use these definitions.
 
-Cython definition fille `dpctl.sycl` provides incomplete definitions of core SYCL runtime classes as
+Cython definition file `dpctl.sycl` provides incomplete definitions of core SYCL runtime classes as
 well as conversion routine between `SyclInterface` reference types and SYCL runtime classes.
diff --git a/examples/cython/sycl_buffer/README.md b/examples/cython/sycl_buffer/README.md
index 89d57f0d40..6cda697dcf 100644
--- a/examples/cython/sycl_buffer/README.md
+++ b/examples/cython/sycl_buffer/README.md
@@ -1,6 +1,6 @@
 # SYCL Extension Working NumPy Array Input via SYCL Buffers
 
-## Decription
+## Description
 
 Cython function expecting a 2D array in a C-contiguous layout that
 computes column-wise total by using SYCL oneMKL (as GEMV call with
diff --git a/examples/cython/sycl_buffer/scripts/bench.py b/examples/cython/sycl_buffer/scripts/bench.py
index 0589c7ad36..60be049e65 100644
--- a/examples/cython/sycl_buffer/scripts/bench.py
+++ b/examples/cython/sycl_buffer/scripts/bench.py
@@ -79,7 +79,7 @@ def run_offload(selector_string, X):
 for ss in ["opencl:cpu", "opencl:gpu", "level_zero:gpu"]:
     print(f"Result for '{ss}': {run_offload(ss, X)}")
 
-print("=" * 10 + " Running bechmarks " + "=" * 10)
+print("=" * 10 + " Running benchmarks " + "=" * 10)
 
 for ss in ["opencl:cpu", "opencl:gpu", "level_zero:gpu"]:
     print(f"Timing offload to '{ss}': {bench_offload(ss, X)}")
diff --git a/examples/pybind11/onemkl_gemv/sycl_gemm/build.sh b/examples/pybind11/onemkl_gemv/sycl_gemm/build.sh
index 221f484ae1..7e3a770d3b 100755
--- a/examples/pybind11/onemkl_gemv/sycl_gemm/build.sh
+++ b/examples/pybind11/onemkl_gemv/sycl_gemm/build.sh
@@ -1,15 +1,20 @@
 #!/bin/bash -x
 
-export PYBIND11_INCLUDES=$(python3 -m pybind11 --includes)
-export DPCTL_INCLUDE_DIR=$(python -c "import dpctl; print(dpctl.get_include())")
-export DPCTL_LIB_DIR=${DPCTL_INCLUDE_DIR}/..
-export PY_EXT_SUFFIX=$(python3-config --extension-suffix)
-export HOST_COMPILER_FLAGS="-g -std=c++2a -O3 -Wno-return-type -Wno-deprecated-declarations -fPIC ${PYBIND11_INCLUDES} -I${DPCTL_INCLUDE_DIR}"
+PYBIND11_INCLUDES=$(python3 -m pybind11 --includes)
+export PYBIND11_INCLUDES
+DPCTL_INCLUDE_DIR=$(python -c "import dpctl; print(dpctl.get_include())")
+export DPCTL_INCLUDE_DIR
+DPCTL_LIB_DIR=${DPCTL_INCLUDE_DIR}/..
+export DPCTL_LIB_DIR
+PY_EXT_SUFFIX=$(python3-config --extension-suffix)
+export PY_EXT_SUFFIX
+HOST_COMPILER_FLAGS="-g -std=c++2a -O3 -Wno-return-type -Wno-deprecated-declarations -fPIC ${PYBIND11_INCLUDES} -I${DPCTL_INCLUDE_DIR}"
+export HOST_COMPILER_FLAGS
 
 #    -fsycl-host-compiler=g++ \
 #    -fsycl-host-compiler-options="${HOST_COMPILER_FLAGS}" \
 
 dpcpp -O3 -fsycl -Wno-deprecated-declarations \
     -fpic -fPIC -shared \
-     ${PYBIND11_INCLUDES} -I${DPCTL_INCLUDE_DIR} \
-    sycl_gemm.cpp -o _sycl_gemm${PY_EXT_SUFFIX}
+     "${PYBIND11_INCLUDES}" -I"${DPCTL_INCLUDE_DIR}" \
+    sycl_gemm.cpp -o _sycl_gemm"${PY_EXT_SUFFIX}"
diff --git a/examples/python/_runner.py b/examples/python/_runner.py
index 94ab6b4077..452f0207b9 100644
--- a/examples/python/_runner.py
+++ b/examples/python/_runner.py
@@ -75,7 +75,7 @@ def run_examples(example_description, glbls_dict):
                 if has_nondefault_params(sgn):
                     if not args.quiet:
                         print(
-                            f"INFO: Skip exectution of {fn} as it "
+                            f"INFO: Skip execution of {fn} as it "
                             "requires arguments"
                         )
                 else:
diff --git a/examples/python/subdevices.py b/examples/python/subdevices.py
index e5e69373b8..282674f297 100644
--- a/examples/python/subdevices.py
+++ b/examples/python/subdevices.py
@@ -35,7 +35,7 @@ def subdivide_root_cpu_device():
     """
     Create root CPU device, and equally partition it
     into smaller CPU devices 4 execution units each,
-    and then further parition those subdevice into
+    and then further partition those sub-devices into
     smaller sub-devices
     """
     cpu_d = dpctl.SyclDevice("cpu")
@@ -84,7 +84,7 @@ def create_subdevice_queue():
     cpu_count = cpu_d.max_compute_units
     sub_devs = cpu_d.create_sub_devices(partition=cpu_count // 2)
     multidevice_ctx = dpctl.SyclContext(sub_devs)
-    # create a SyclQueue for each sub-device, using commont
+    # create a SyclQueue for each sub-device, using common
     # multi-device context
     q0, q1 = [dpctl.SyclQueue(multidevice_ctx, d) for d in sub_devs]
     # for each sub-device allocate 26 bytes
diff --git a/examples/python/usm_memory_host_access.py b/examples/python/usm_memory_host_access.py
index 21092a01a5..f6af9c4305 100644
--- a/examples/python/usm_memory_host_access.py
+++ b/examples/python/usm_memory_host_access.py
@@ -23,7 +23,7 @@
 
 # USM-shared and USM-host pointers are host-accessible,
 # meaning they are accessible from Python, therefore
-# they implement Pyton buffer protocol
+# they implement Python buffer protocol
 
 # allocate 1K of USM-shared buffer
 ms = dpmem.MemoryUSMShared(1024)
diff --git a/libsyclinterface/cmake/modules/GetProjectVersion.cmake b/libsyclinterface/cmake/modules/GetProjectVersion.cmake
index 8968eec729..a45ad0ee96 100644
--- a/libsyclinterface/cmake/modules/GetProjectVersion.cmake
+++ b/libsyclinterface/cmake/modules/GetProjectVersion.cmake
@@ -45,7 +45,7 @@ function(get_version)
     if(NOT result EQUAL 0)
         message(WARNING
             "Something went wrong when executing \"git describe\". "
-            "Seting all version values to 0."
+            "Setting all version values to 0."
         )
         set(VERSION_MAJOR 0     PARENT_SCOPE)
         set(VERSION_MINOR 0     PARENT_SCOPE)
@@ -63,7 +63,7 @@ function(get_version)
     else()
         message(WARNING
             "The last git tag does not use proper semantic versioning. "
-            "Seting all version values to 0."
+            "Setting all version values to 0."
         )
         set(VERSION_MAJOR 0     PARENT_SCOPE)
         set(VERSION_MINOR 0     PARENT_SCOPE)
@@ -86,7 +86,7 @@ function(get_version)
     if(NOT result EQUAL 0)
         message(WARNING
             "Something went wrong when executing \"git describe\". "
-            "Seting all version values to 0."
+            "Setting all version values to 0."
         )
         set(VERSION_MAJOR 0     PARENT_SCOPE)
         set(VERSION_MINOR 0     PARENT_SCOPE)
diff --git a/libsyclinterface/dbg_build.bat b/libsyclinterface/dbg_build.bat
index fd591d9399..6101f41471 100644
--- a/libsyclinterface/dbg_build.bat
+++ b/libsyclinterface/dbg_build.bat
@@ -5,7 +5,7 @@ if errorlevel 1 (
     call "%ONEAPI_ROOT%\compiler\latest\env\vars.bat"
     if %ERRORLEVEL% neq 0 exit 1
 )
-@REM conda uses %ERRORLEVEL% but FPGA scripts can set it. So it should be reseted.
+@REM conda uses %ERRORLEVEL% but FPGA scripts can set it. So it should be reset.
 set ERRORLEVEL=
 
 rmdir /S /Q build
diff --git a/libsyclinterface/dbg_build.sh b/libsyclinterface/dbg_build.sh
index cba4e71d71..c090bc700c 100755
--- a/libsyclinterface/dbg_build.sh
+++ b/libsyclinterface/dbg_build.sh
@@ -5,10 +5,11 @@ mkdir build
 pushd build || exit 1
 
 INSTALL_PREFIX=$(pwd)/../install
-rm -rf ${INSTALL_PREFIX}
+rm -rf "${INSTALL_PREFIX}"
 
-# With DPC++ 2024.0 adn newer set these to ensure that
+# With DPC++ 2024.0 and newer set these to ensure that
 # cmake can find llvm-cov and other utilities
+# shellcheck disable=SC2034
 LLVM_TOOLS_HOME=${CMPLR_ROOT}/bin/compiler
 PATH=$PATH:${CMPLR_ROOT}/bin/compiler
 
@@ -17,8 +18,8 @@ cmake                                                       \
     -DCMAKE_C_COMPILER=icx                                  \
     -DCMAKE_CXX_COMPILER=icpx                               \
     -DCMAKE_CXX_FLAGS=-fsycl                                \
-    -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX}                \
-    -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX}                   \
+    -DCMAKE_INSTALL_PREFIX="${INSTALL_PREFIX}"              \
+    -DCMAKE_PREFIX_PATH="${INSTALL_PREFIX}"                 \
     -DDPCTL_ENABLE_L0_PROGRAM_CREATION=ON                   \
     -DDPCTL_BUILD_CAPI_TESTS=ON                             \
     -DDPCTL_GENERATE_COVERAGE=OFF                           \
diff --git a/libsyclinterface/dbg_build_custom.sh b/libsyclinterface/dbg_build_custom.sh
index ad738a0c8a..4c30cca1ce 100755
--- a/libsyclinterface/dbg_build_custom.sh
+++ b/libsyclinterface/dbg_build_custom.sh
@@ -2,10 +2,10 @@
 set +xe
 rm -rf build
 mkdir build
-pushd build
+pushd build || exit 1
 
-INSTALL_PREFIX=`pwd`/../install
-rm -rf ${INSTALL_PREFIX}
+INSTALL_PREFIX=$(pwd)/../install
+rm -rf "${INSTALL_PREFIX}"
 
 if [[ -z "${DPCPP_HOME}" ]]; then
     echo "Set the DPCPP_HOME environment variable to root directory."
@@ -13,11 +13,11 @@ fi
 
 cmake                                                       \
     -DCMAKE_BUILD_TYPE=Debug                                \
-    -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX}                \
-    -DCMAKE_PREFIX_PATH=${INSTALL_PREFIX}                   \
-    -DDPCTL_CUSTOM_DPCPP_INSTALL_DIR=${DPCPP_HOME}          \
-    -DCMAKE_LINKER:PATH=${DPCPP_HOME}/bin/lld               \
-    -DDPCTL_ENABLE_L0_PROGRAM_CREATION=${USE_LO_HEADERS}    \
+    -DCMAKE_INSTALL_PREFIX="${INSTALL_PREFIX}"              \
+    -DCMAKE_PREFIX_PATH="${INSTALL_PREFIX}"                 \
+    -DDPCTL_CUSTOM_DPCPP_INSTALL_DIR="${DPCPP_HOME}"        \
+    -DCMAKE_LINKER:PATH="${DPCPP_HOME}/bin/lld"             \
+    -DDPCTL_ENABLE_L0_PROGRAM_CREATION="${USE_LO_HEADERS}"  \
     -DDPCTL_BUILD_CAPI_TESTS=ON                             \
     -DDPCTL_GENERATE_COVERAGE=ON                            \
     ..
@@ -32,4 +32,4 @@ make lcov-genhtml
 # ctest -V --progress --output-on-failure -j 4
 # cd ..
 
-popd
+popd || exit 1
diff --git a/libsyclinterface/helper/source/dpctl_error_handlers.cpp b/libsyclinterface/helper/source/dpctl_error_handlers.cpp
index 149eb26692..8ee9626f00 100644
--- a/libsyclinterface/helper/source/dpctl_error_handlers.cpp
+++ b/libsyclinterface/helper/source/dpctl_error_handlers.cpp
@@ -20,7 +20,7 @@
 ///
 /// \file
 /// A functor to use for passing an error handler callback function to sycl
-/// context and queue contructors.
+/// context and queue constructors.
 //===----------------------------------------------------------------------===//
 
 #include "dpctl_error_handlers.h"
diff --git a/libsyclinterface/include/syclinterface/Support/DllExport.h b/libsyclinterface/include/syclinterface/Support/DllExport.h
index 0578387145..f2eb627bc7 100644
--- a/libsyclinterface/include/syclinterface/Support/DllExport.h
+++ b/libsyclinterface/include/syclinterface/Support/DllExport.h
@@ -1,4 +1,4 @@
-//===---------  DllExport.h - Decalres dllexport for Windows     -*-C++-*- ===//
+//===---------  DllExport.h - Declares dllexport for Windows     -*-C++-*- ===//
 //
 //                      Data Parallel Control (dpctl)
 //
diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_device_manager.h b/libsyclinterface/include/syclinterface/dpctl_sycl_device_manager.h
index cec8ec076c..3f4f08e18f 100644
--- a/libsyclinterface/include/syclinterface/dpctl_sycl_device_manager.h
+++ b/libsyclinterface/include/syclinterface/dpctl_sycl_device_manager.h
@@ -41,7 +41,7 @@ DPCTL_C_EXTERN_C_BEGIN
  * @{
  */
 
-// Declares a set of types abd functions to deal with vectors of
+// Declares a set of types and functions to deal with vectors of
 // DPCTLSyclDeviceRef. Refer dpctl_vector_macros.h
 DPCTL_DECLARE_VECTOR(Device)
 
diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_platform_manager.h b/libsyclinterface/include/syclinterface/dpctl_sycl_platform_manager.h
index 99ff224317..5fce58bd19 100644
--- a/libsyclinterface/include/syclinterface/dpctl_sycl_platform_manager.h
+++ b/libsyclinterface/include/syclinterface/dpctl_sycl_platform_manager.h
@@ -39,7 +39,7 @@ DPCTL_C_EXTERN_C_BEGIN
  * @{
  */
 
-// Declares a set of types abd functions to deal with vectors of
+// Declares a set of types and functions to deal with vectors of
 // DPCTLSyclPlatformRef. Refer dpctl_vector_macros.h
 DPCTL_DECLARE_VECTOR(Platform)
 
diff --git a/libsyclinterface/tests/test_service.cpp b/libsyclinterface/tests/test_service.cpp
index b8581092d5..4b75b34e4c 100644
--- a/libsyclinterface/tests/test_service.cpp
+++ b/libsyclinterface/tests/test_service.cpp
@@ -1,4 +1,4 @@
-//===--- test_service.cpp - Test cases for sevice functions  ===//
+//===--- test_service.cpp - Test cases for service functions  ===//
 //
 //                      Data Parallel Control (dpctl)
 //

From 22954af8926518dc3d47d73eb3d6b309b3a668af Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Wed, 15 Apr 2026 20:23:07 -0700
Subject: [PATCH 05/24] ignore blame for linting commit

---
 .git-blame-ignore-revs | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/.git-blame-ignore-revs b/.git-blame-ignore-revs
index 654afad97a..c60ead8a6f 100644
--- a/.git-blame-ignore-revs
+++ b/.git-blame-ignore-revs
@@ -17,3 +17,6 @@ b37657e6ad9af16eaec2982d8e2397acd2af2881
 
 # Add cython-lint to pre-commit config
 0ce1aef210ffb88b7d2ea3a89e861486498f652f
+
+# Add new linters and rules to pre-commit
+5cf074899f226726d07026db9899bb7632d8c1de

From a871555698dc0c0fbbe19dce17689801bbbe6f9a Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Wed, 15 Apr 2026 20:27:12 -0700
Subject: [PATCH 06/24] fix CI after linting changes

---
 .github/workflows/conda-package.yml      | 9 +--------
 .github/workflows/generate-coverage.yaml | 8 ++++----
 .github/workflows/generate-docs.yml      | 8 ++++----
 3 files changed, 9 insertions(+), 16 deletions(-)

diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml
index 8924245c78..e1e24eb7ad 100644
--- a/.github/workflows/conda-package.yml
+++ b/.github/workflows/conda-package.yml
@@ -60,14 +60,7 @@ jobs:
       - name: Build conda package
         run: |
           # use bootstrap channel to pull NumPy linked with OpenBLAS
-          CHANNELS="-c conda-forge --override-channels"
-          VERSIONS="--python ${{ matrix.python }} --numpy 2.0"
-          TEST="--no-test"
-          conda build \
-            "$TEST" \
-            "$VERSIONS" \
-            "$CHANNELS" \
-            conda-recipe
+          conda build --no-test --python ${{ matrix.python }} --numpy 2.0 -c conda-forge --override-channels conda-recipe
       - name: Upload artifact
         uses: actions/upload-artifact@043fb46d1a93c77aae656e7c1c64a875d1fc6a0a # v7.0.1
         with:
diff --git a/.github/workflows/generate-coverage.yaml b/.github/workflows/generate-coverage.yaml
index d29cba85a3..a1f4e91c22 100644
--- a/.github/workflows/generate-coverage.yaml
+++ b/.github/workflows/generate-coverage.yaml
@@ -26,10 +26,10 @@ jobs:
 
       - name: Add Intel repository
         run: |
-          wget https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB
-          cat GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB | gpg --dearmor | sudo tee /usr/share/keyrings/oneapi-archive-keyring.gpg > /dev/null
-          rm GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB
-          echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main" | sudo tee /etc/apt/sources.list.d/oneAPI.list
+          wget -qO- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB \
+            | gpg --dearmor | sudo tee /usr/share/keyrings/oneapi-archive-keyring.gpg > /dev/null
+          echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main" \
+            | sudo tee /etc/apt/sources.list.d/oneAPI.list
           sudo apt update
 
       - name: Install latest Intel OneAPI
diff --git a/.github/workflows/generate-docs.yml b/.github/workflows/generate-docs.yml
index cd47d7a3ea..fa81f764de 100644
--- a/.github/workflows/generate-docs.yml
+++ b/.github/workflows/generate-docs.yml
@@ -24,10 +24,10 @@ jobs:
       - name: Add Intel repository
         if: ${{ !github.event.pull_request || github.event.action != 'closed' }}
         run: |
-          wget https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB
-          cat GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB | gpg --dearmor | sudo tee /usr/share/keyrings/oneapi-archive-keyring.gpg > /dev/null
-          rm GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB
-          echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main" | sudo tee /etc/apt/sources.list.d/oneAPI.list
+          wget -qO- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB \
+            | gpg --dearmor | sudo tee /usr/share/keyrings/oneapi-archive-keyring.gpg > /dev/null
+          echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main" \
+            | sudo tee /etc/apt/sources.list.d/oneAPI.list
           sudo apt update
       - name: Install Intel OneAPI
         if: ${{ !github.event.pull_request || github.event.action != 'closed' }}

From e0e8f27e5ca0bda3bf444d7857731575b1fbb5da Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Wed, 15 Apr 2026 20:30:54 -0700
Subject: [PATCH 07/24] use mirrors-clang-format

also moves to clang-format 22.1.3
---
 .pre-commit-config.yaml | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
index 08fe5cee1a..58174eeafd 100644
--- a/.pre-commit-config.yaml
+++ b/.pre-commit-config.yaml
@@ -71,8 +71,8 @@ repos:
     rev: 7.3.0
     hooks:
     -   id: flake8
--   repo: https://github.com/pocc/pre-commit-hooks
-    rev: v1.3.5
+-   repo: https://github.com/pre-commit/mirrors-clang-format
+    rev: v22.1.3
     hooks:
     -   id: clang-format
         args: ["-i"]

From 759dcdc80d0d99790205e87bdd7c5d4f2e93a1ff Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Wed, 15 Apr 2026 20:37:05 -0700
Subject: [PATCH 08/24] reformat for clang-format 22.1.3

---
 .../external_usm_allocation/_usm_alloc_example.cpp              | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/examples/pybind11/external_usm_allocation/external_usm_allocation/_usm_alloc_example.cpp b/examples/pybind11/external_usm_allocation/external_usm_allocation/_usm_alloc_example.cpp
index b982fcaa20..f92a99026e 100644
--- a/examples/pybind11/external_usm_allocation/external_usm_allocation/_usm_alloc_example.cpp
+++ b/examples/pybind11/external_usm_allocation/external_usm_allocation/_usm_alloc_example.cpp
@@ -49,7 +49,7 @@ struct DMatrix
         : n_(rows), m_(columns), q_(q), alloc_(q), vec_(n_ * m_, alloc_)
     {
     }
-    ~DMatrix(){};
+    ~DMatrix() {};
     DMatrix(const DMatrix &) = default;
     DMatrix(DMatrix &&) = default;
 

From 0bde0486266b1baab04a700f01ab83d4d1d21ebc Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Wed, 15 Apr 2026 21:12:25 -0700
Subject: [PATCH 09/24] add clang-format change to blame ignore

---
 .git-blame-ignore-revs | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/.git-blame-ignore-revs b/.git-blame-ignore-revs
index c60ead8a6f..dccd02b48a 100644
--- a/.git-blame-ignore-revs
+++ b/.git-blame-ignore-revs
@@ -20,3 +20,6 @@ b37657e6ad9af16eaec2982d8e2397acd2af2881
 
 # Add new linters and rules to pre-commit
 5cf074899f226726d07026db9899bb7632d8c1de
+
+# Transition from clang-format 18 to clang-format 22
+759dcdc80d0d99790205e87bdd7c5d4f2e93a1ff

From 434c07882adf1d01b6cc7a8d4bea22bd07eb84e3 Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Wed, 15 Apr 2026 21:53:56 -0700
Subject: [PATCH 10/24] add disable shellcheck

---
 .github/workflows/conda-package.yml | 36 ++++++++++++++++++-----------
 1 file changed, 23 insertions(+), 13 deletions(-)

diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml
index e1e24eb7ad..55b2ee5d4b 100644
--- a/.github/workflows/conda-package.yml
+++ b/.github/workflows/conda-package.yml
@@ -185,8 +185,9 @@ jobs:
           CHANNELS="-c $GITHUB_WORKSPACE/channel ${{ env.CHANNELS }}"
           export CHANNELS
           PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
-          export PACKAGE_VERSION
-          conda create -n "${{ env.TEST_ENV_NAME }}" "${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION}" python=${{ matrix.python }} "$CHANNELS" --only-deps --dry-run > lockfile
+            export PACKAGE_VERSION
+            # shellcheck disable=SC2086
+          conda create -n ${{ env.TEST_ENV_NAME }} ${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION} python=${{ matrix.python }} $CHANNELS --only-deps --dry-run > lockfile
           cat lockfile
       - name: Set pkgs_dirs
         run: |
@@ -210,7 +211,8 @@ jobs:
           export TEST_DEPENDENCIES
           PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
           export PACKAGE_VERSION
-          conda create -n "${{ env.TEST_ENV_NAME }}" "${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION}" "${TEST_DEPENDENCIES}" python=${{ matrix.python }} "${CHANNELS}"
+          # shellcheck disable=SC2086
+          conda create -n ${{ env.TEST_ENV_NAME }} ${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION} $TEST_DEPENDENCIES python=${{ matrix.python }} $CHANNELS
           # Test installed packages
           conda list -n "${{ env.TEST_ENV_NAME }}"
       - name: Smoke test
@@ -553,7 +555,8 @@ jobs:
           export CHANNELS
           PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
           export PACKAGE_VERSION
-          conda create -n "${{ env.EXAMPLES_ENV_NAME }}" "${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION}" python=${{ matrix.python }} "$CHANNELS" --only-deps --dry-run > lockfile
+          # shellcheck disable=SC2086
+          conda create -n ${{ env.EXAMPLES_ENV_NAME }} ${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION} python=${{ matrix.python }} $CHANNELS --only-deps --dry-run > lockfile
           cat lockfile
       - name: Set pkgs_dirs
         run: |
@@ -577,20 +580,26 @@ jobs:
           export CHANNELS="${{ env.CHANNELS }}"
           . "$CONDA/etc/profile.d/conda.sh"
           DPCTL_DEPENDS="$(python -c "${VER_SCRIPT1} ${VER_SCRIPT3}")"
+          export DPCTL_DEPENDS
           echo "Dpctl dependencies: ${DPCTL_DEPENDS}"
-          conda create -n "${{ env.EXAMPLES_ENV_NAME }}" -y pytest python="${{ matrix.python }}" "setuptools<72.2.0" "$CHANNELS"
+          # shellcheck disable=SC2086
+          conda create -n ${{ env.EXAMPLES_ENV_NAME }} -y pytest python=${{ matrix.python }} "setuptools<72.2.0" $CHANNELS
           echo "Environment created"
-          conda install -n "${{ env.EXAMPLES_ENV_NAME }}" -y cmake ninja "$CHANNELS" || exit 1
+          # shellcheck disable=SC2086
+          conda install -n ${{ env.EXAMPLES_ENV_NAME }} -y cmake ninja $CHANNELS || exit 1
           echo "Cmake and Ninja installed"
-          conda install -n "${{ env.EXAMPLES_ENV_NAME }}" -y pybind11 cython scikit-build "$CHANNELS" || exit 1
+          # shellcheck disable=SC2086
+          conda install -n ${{ env.EXAMPLES_ENV_NAME }} -y pybind11 cython scikit-build $CHANNELS || exit 1
           echo "scikit-build installed"
-          conda install -n "${{ env.EXAMPLES_ENV_NAME }}" -y mkl-dpcpp \
+          # shellcheck disable=SC2086
+          conda install -n ${{ env.EXAMPLES_ENV_NAME }} -y mkl-dpcpp \
               mkl-devel-dpcpp dpcpp_cpp_rt "${DPCTL_DEPENDS}" \
-              "$CHANNELS" || exit 1
+              $CHANNELS || exit 1
           echo "IPL installed"
-          conda create -y -n "${{ env.BUILD_ENV_NAME }}" "$CHANNELS" gcc_linux-64 gxx_linux-64 \
-              "${{ env.DPCPP_CMPLR }}" "${DPCTL_DEPENDS}" \
-              "sysroot_linux-64>=2.28"
+          # shellcheck disable=SC2086
+          conda create -y -n ${{ env.BUILD_ENV_NAME }} $CHANNELS gcc_linux-64 gxx_linux-64 \
+              ${{ env.DPCPP_CMPLR }} ${DPCTL_DEPENDS} \
+              sysroot_linux-64>=2.28
           echo "Compiler installed"
           conda list -n "${{ env.BUILD_ENV_NAME }}"
       - name: Install dpctl
@@ -601,7 +610,8 @@ jobs:
           export CHANNELS
           PACKAGE_VERSION=$(python -c "${VER_SCRIPT1} ${VER_SCRIPT2}")
           export PACKAGE_VERSION
-          conda install -n "${{ env.EXAMPLES_ENV_NAME }}" -y "${CHANNELS}" "${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION}" dpnp || exit 1
+          # shellcheck disable=SC2086
+          conda install -n ${{ env.EXAMPLES_ENV_NAME }} -y $CHANNELS ${{ env.PACKAGE_NAME }}=${PACKAGE_VERSION} dpnp || exit 1
       - name: Build and run examples of pybind11 extensions
         shell: bash -l {0}
         run: |

From 0cfc3a7905228eb2c6aa9a0cba7ba5e368f73a97 Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Thu, 16 Apr 2026 01:25:37 -0700
Subject: [PATCH 11/24] remove gdb test run from workflow

element-wise tests no longer present after migration of tensor
---
 .github/workflows/conda-package.yml | 9 ---------
 1 file changed, 9 deletions(-)

diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml
index 55b2ee5d4b..ff53b958b5 100644
--- a/.github/workflows/conda-package.yml
+++ b/.github/workflows/conda-package.yml
@@ -220,15 +220,6 @@ jobs:
           . "$CONDA/etc/profile.d/conda.sh"
           conda activate "${{ env.TEST_ENV_NAME }}"
           python -c "import dpctl; dpctl.lsplatform(verbosity=2)"
-      - name: Install gdb
-        run: |
-          sudo apt-get update --fix-missing
-          sudo apt-get install -y gdb
-      - name: Run test_elementwise under gdb
-        run: |
-          . "$CONDA/etc/profile.d/conda.sh"
-          conda activate "${{ env.TEST_ENV_NAME }}"
-          gdb --batch -ex r -ex 'info sharedlibrary' -ex 'set print elements 1000' -ex bt --args "${CONDA_PREFIX}/bin/python" -m pytest -q -ra --disable-warnings --pyargs dpctl.tests.elementwise.test_trigonometric::test_trig_order -vv || true
       - name: Create test temp dir
         # create temporary empty folder to runs tests from
         # https://github.com/pytest-dev/pytest/issues/11904

From a2d0a2e4c75d9b99dd0a8d70d143ffea21d92eaa Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Thu, 16 Apr 2026 18:57:23 -0700
Subject: [PATCH 12/24] fix comment for cmd shell

---
 .github/workflows/conda-package.yml | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml
index ff53b958b5..a541ae940c 100644
--- a/.github/workflows/conda-package.yml
+++ b/.github/workflows/conda-package.yml
@@ -126,7 +126,7 @@ jobs:
         env:
           OVERRIDE_INTEL_IPO: 1   # IPO requires more resources that GH actions VM provides
         run: |
-          # TODO: roll back use of Intel channel when 2025.1 is available on conda-forge
+          :: TODO: roll back use of Intel channel when 2025.1 is available on conda-forge
           conda build --no-test --python ${{ matrix.python }} --numpy 2.0 -c ${{ env.INTEL_CHANNEL }} -c conda-forge --override-channels conda-recipe
 
       - name: Upload artifact

From 39e7abb83f2a2552df93b934554f955539132dd3 Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Thu, 16 Apr 2026 18:57:49 -0700
Subject: [PATCH 13/24] remove w/a for old anaconda-client

---
 .github/workflows/conda-package.yml | 4 ----
 1 file changed, 4 deletions(-)

diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml
index a541ae940c..40c8d60464 100644
--- a/.github/workflows/conda-package.yml
+++ b/.github/workflows/conda-package.yml
@@ -434,8 +434,6 @@ jobs:
       - name: Upload
         env:
           ANACONDA_TOKEN: ${{ secrets.ANACONDA_TOKEN }}
-          # remove when anaconda-client 1.13.1 is available
-          ANACONDA_CLIENT_FORCE_STANDALONE: true
         if: ${{ env.ANACONDA_TOKEN != '' }}
         run: |
           anaconda --token "${{ env.ANACONDA_TOKEN }}" upload --user dppy --label dev "${{ env.PACKAGE_NAME }}"-*.conda
@@ -443,8 +441,6 @@ jobs:
       - name: Upload Wheels
         env:
           ANACONDA_TOKEN: ${{ secrets.ANACONDA_TOKEN }}
-          # remove when anaconda-client 1.13.1 is available
-          ANACONDA_CLIENT_FORCE_STANDALONE: true
         run: anaconda --token "${{ env.ANACONDA_TOKEN }}" upload --user dppy --label dev "${{ env.PACKAGE_NAME }}"-*.whl --version "${{ env.PACKAGE_VERSION }}"
 
   upload_windows:

From 6c303c993f3ec8388686cbafd14bcf556a81e9cb Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Fri, 24 Apr 2026 10:35:47 -0700
Subject: [PATCH 14/24] add  0.22.1 changelog

---
 CHANGELOG.md | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/CHANGELOG.md b/CHANGELOG.md
index 3d16c02dbc..1d2035dd76 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -12,6 +12,10 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 
 ### Fixed
 
+## [0.22.1] - Apr. 24, 2026
+
+This is a bug-fix release which fixes a memory leak in `dpctl.RawKernelArg` [gh-2294](https://github.com/IntelPython/dpctl/pull/2294).
+
 ## [0.22.0] - Apr. 14, 2026
 
 The highlight of this release is the full migration of `dpctl.tensor` submodule to sister project [`dpnp`](https://github.com/IntelPython/dpnp), shrinking the size of the package tremendously, by between 93% and 96%. The `__sycl_usm_array_interface__` is still supported, with `dpctl` serving as curator of the protocol.

From 8707aa9f74c4fdf1c9972f1580fa5978b125b053 Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Fri, 24 Apr 2026 11:07:34 -0700
Subject: [PATCH 15/24] Reformat changelog for 0.22.1

---
 CHANGELOG.md | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/CHANGELOG.md b/CHANGELOG.md
index 1d2035dd76..0a9e67c252 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -14,13 +14,15 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 
 ## [0.22.1] - Apr. 24, 2026
 
-This is a bug-fix release which fixes a memory leak in `dpctl.RawKernelArg` [gh-2294](https://github.com/IntelPython/dpctl/pull/2294).
+This is a bug-fix release which fixes a memory leak in `dpctl.RawKernelArg`.
+
+### Fixed
+* Fixed a memory leak in `dpctl.RawKernelArg` [gh-2294](https://github.com/IntelPython/dpctl/pull/2294)
 
 ## [0.22.0] - Apr. 14, 2026
 
 The highlight of this release is the full migration of `dpctl.tensor` submodule to sister project [`dpnp`](https://github.com/IntelPython/dpnp), shrinking the size of the package tremendously, by between 93% and 96%. The `__sycl_usm_array_interface__` is still supported, with `dpctl` serving as curator of the protocol.
 
-
 Additionally, `dpctl` build scripts were updated, removing use of `python setup.py develop` and `python setup.py install`, and `dpctl` [documentation page](https://intelpython.github.io/dpctl/latest/index.html) now supports a version dropdown.
 
 **NOTE**: Changes below which reference `tensor` were added to the `tensor` submodule prior to release, and therefore are included in the migrated `tensor` submodule in [`dpnp`](https://github.com/IntelPython/dpnp). They are included here for transparency and continuity of the submodule's history in the changelog.

From 117eac712eaa285f75675c22510754f549a0a956 Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Fri, 24 Apr 2026 11:46:31 -0700
Subject: [PATCH 16/24] add more pre-commit hooks

---
 .pre-commit-config.yaml  | 2 ++
 conda-recipe/run_test.sh | 0
 2 files changed, 2 insertions(+)
 mode change 100644 => 100755 conda-recipe/run_test.sh

diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
index 58174eeafd..9b194e8d3d 100644
--- a/.pre-commit-config.yaml
+++ b/.pre-commit-config.yaml
@@ -23,6 +23,8 @@ repos:
     -   id: check-case-conflict
     -   id: check-executables-have-shebangs
     -   id: check-merge-conflict
+    -   id: check-shebang-scripts-are-executable
+    -   id: check-symlinks
     -   id: check-toml
     -   id: debug-statements
     -   id: destroyed-symlinks
diff --git a/conda-recipe/run_test.sh b/conda-recipe/run_test.sh
old mode 100644
new mode 100755

From 1d634c715904e838e926cc0b0d451fb13c78a789 Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Mon, 27 Apr 2026 19:15:23 -0700
Subject: [PATCH 17/24] Implement SpecializationConstant class

---
 dpctl/_backend.pxd                            |   4 +
 dpctl/_sycl_platform.pyx                      |   2 +-
 dpctl/program/__init__.py                     |   2 +
 dpctl/program/_program.pyx                    | 170 ++++++++++++++++++
 .../dpctl_sycl_kernel_bundle_interface.h      |   7 +
 5 files changed, 184 insertions(+), 1 deletion(-)

diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd
index 93d9b5ef97..d2f1a0decb 100644
--- a/dpctl/_backend.pxd
+++ b/dpctl/_backend.pxd
@@ -431,6 +431,10 @@ cdef extern from "syclinterface/dpctl_sycl_context_interface.h":
 
 
 cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h":
+    ctypedef struct _spec_const "DPCTLSpecConst":
+        uint32_t id
+        size_t size
+        const void *value
     cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromSpirv(
         const DPCTLSyclContextRef Ctx,
         const DPCTLSyclDeviceRef Dev,
diff --git a/dpctl/_sycl_platform.pyx b/dpctl/_sycl_platform.pyx
index 41eff7b5d3..ba78226e50 100644
--- a/dpctl/_sycl_platform.pyx
+++ b/dpctl/_sycl_platform.pyx
@@ -236,7 +236,7 @@ cdef class SyclPlatform(_SyclPlatform):
           and filter string for each device is printed.
 
         Args:
-            verbosity (Literal[0, 1, 2], optional):.
+            verbosity (Literal[0, 1, 2], optional):
                 The verbosity controls how much information is printed by the
                 function. Value ``0`` is the lowest level set by default and
                 ``2`` is the highest level to print the most verbose output.
diff --git a/dpctl/program/__init__.py b/dpctl/program/__init__.py
index 71302e4186..bb16f9611b 100644
--- a/dpctl/program/__init__.py
+++ b/dpctl/program/__init__.py
@@ -22,6 +22,7 @@
 """
 
 from ._program import (
+    SpecializationConstant,
     SyclKernel,
     SyclKernelBundle,
     SyclKernelBundleCompilationError,
@@ -41,6 +42,7 @@
     "SyclKernelBundleCompilationError",
     "SyclProgram",
     "SyclProgramCompilationError",
+    "SpecializationConstant",
 ]
 
 
diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx
index 8737be4762..7cecc3fa4d 100644
--- a/dpctl/program/_program.pyx
+++ b/dpctl/program/_program.pyx
@@ -26,7 +26,17 @@ an OpenCL source string or a SPIR-V binary file.
 
 """
 
+from cpython.buffer cimport (
+    Py_buffer,
+    PyBUF_ANY_CONTIGUOUS,
+    PyBUF_SIMPLE,
+    PyBuffer_Release,
+    PyObject_CheckBuffer,
+    PyObject_GetBuffer,
+)
+from cpython.bytes cimport PyBytes_FromStringAndSize
 from libc.stdint cimport uint32_t
+from libc.string cimport memcmp
 
 import warnings
 
@@ -51,14 +61,20 @@ from dpctl._backend cimport (  # noqa: E211, E402;
     DPCTLSyclDeviceRef,
     DPCTLSyclKernelBundleRef,
     DPCTLSyclKernelRef,
+    _spec_const,
 )
 
+import numbers
+
+import numpy as np
+
 __all__ = [
     "create_kernel_bundle_from_source",
     "create_kernel_bundle_from_spirv",
     "SyclKernel",
     "SyclKernelBundle",
     "SyclKernelBundleCompilationError",
+    "SpecializationConstant",
 ]
 
 cdef class SyclKernelBundleCompilationError(Exception):
@@ -252,6 +268,160 @@ cdef api SyclKernelBundle SyclKernelBundle_Make(DPCTLSyclKernelBundleRef KBRef):
     return SyclKernelBundle._create(copied_KBRef)
 
 
+cdef class SpecializationConstant:
+    """
+    SpecializationConstant(spec_id, *args)
+
+    Python class representing SYCL specialization constants that can be used
+    when creating a :class:`dpctl.program.SyclKernelBundle` from SPIR-V.
+
+    There are multiple ways to create a :class:`.SpecializationConstant`:
+
+    - ``SpecializationConstant(spec_id, obj)``
+      If the constructor is invoked with a single variadic argument, the
+      argument is expected to either expose the Python buffer protocol or be
+      coercible to a NumPy array. If the argument is coercible to a NumPy array
+      or is one, it must have a supported data type (bool, integral, floating
+      point, or void). The specialization constant will be constructed from the
+      data in the buffer
+
+    - ``SpecializationConstant(spec_id, dtype, obj)``
+    If the constructor is invoked with two variadic arguments, and the first
+    argument is a string, it is interpreted as a NumPy ``dtype`` string and the
+    second argument will be coerced to a NumPy array with that data type.
+    The data type specified by the first argument must be a supported data
+    type (bool, integral, floating point, or void).
+
+    - ``SpecializationConstant(spec_id, nbytes, raw_ptr)``
+      If the constructor is invoked with two variadic arguments where both are
+      integers, the first argument is interpreted as the number of bytes and
+      the second argument is interpreted as a pointer to the data.
+
+    Note that when constructing from a buffer, the
+    :class:`.SpecializationConstant`, shares memory with the original object.
+    Modifications to the original object's data after construction will be
+    reflected when the :class:`.SpecializationConstant` is used to create a
+    :class:`.SyclKernelBundle`. This is not the case when constructing from a
+    raw pointer, as the data is copied.
+
+    Args:
+        spec_id (int):
+            The SPIR-V specialization ID.
+        args:
+            Variadic argument, see class documentation.
+
+    Raises:
+        TypeError: In case of incorrect arguments given to constructor,
+                   failure to coerce to a buffer, or unsupported data type when
+                   coercing to a buffer.
+        ValueError: If the provided object fails to construct a buffer.
+    """
+
+    cdef _spec_const _spec_const
+    cdef Py_buffer _buffer
+
+    def __cinit__(self, spec_id, *args):
+        cdef int ret_code = 0
+        cdef object target_obj = None
+
+        if not isinstance(spec_id, numbers.Integral):
+            raise TypeError(
+                "Specialization constant ID must be of type `int`, got "
+                f"{type(spec_id)}"
+            )
+
+        if len(args) == 0 or len(args) > 2:
+            raise TypeError(
+                f"Constructor takes 2 or 3 arguments, got {len(args)}."
+            )
+
+        self._spec_const.id = <uint32_t>spec_id
+
+        if len(args) == 2:
+            if (
+                isinstance(args[0], numbers.Integral) and
+                isinstance(args[1], numbers.Integral)
+            ):
+                target_obj = PyBytes_FromStringAndSize(
+                    <const char *><size_t>args[1], <Py_ssize_t>args[0]
+                )
+            elif isinstance(args[0], str):
+                target_obj = np.ascontiguousarray(args[1], dtype=args[0])
+
+        elif len(args) == 1:
+            target_obj = args[0]
+            if not PyObject_CheckBuffer(target_obj):
+                # attempt to coerce to a numpy array
+                target_obj = np.ascontiguousarray(target_obj)
+        else:
+            raise TypeError(
+                "Invalid arguments."
+            )
+
+        if isinstance(target_obj, np.ndarray):
+            if target_obj.dtype.kind not in ("b", "i", "u", "f", "c", "V"):
+                raise TypeError(
+                    "Coercion of input to buffer resulted in an unsupported "
+                    f"data type '{target_obj.dtype}'. When coercing objects, "
+                    "`SpecializationConstant` expects the data to coerce to a "
+                    "supported type: bool, integral, real or complex floating "
+                    "point, or void. To pass arbitrary data, use a "
+                    "`memoryview` or `bytes` object, or pass the pointer and "
+                    "size directly."
+                )
+
+        ret_code = PyObject_GetBuffer(
+            target_obj, &(self._buffer), PyBUF_SIMPLE | PyBUF_ANY_CONTIGUOUS
+        )
+        if ret_code != 0:
+            raise ValueError(
+                "Failed to get buffer view for the provided object."
+            )
+        self._spec_const.value = <void*>self._buffer.buf
+        self._spec_const.size = <size_t>self._buffer.len
+
+    def __dealloc__(self):
+        PyBuffer_Release(&(self._buffer))
+
+    def __repr__(self):
+        return f"SpecializationConstant({self._spec_const.id})"
+
+    def __eq__(self, other):
+        if not isinstance(other, SpecializationConstant):
+            return False
+        cdef SpecializationConstant _other = <SpecializationConstant>other
+        if (
+            self._spec_const.id != _other._spec_const.id or
+            self._spec_const.size != _other._spec_const.size or
+            self._spec_const.value != _other._spec_const.value
+        ):
+            return False
+        return memcmp(
+            self._spec_const.value,
+            _other._spec_const.value,
+            self._spec_const.size
+        ) == 0
+
+    @property
+    def id(self):
+        """Returns the specialization ID for this specialization constant."""
+        return self._spec_const.id
+
+    @property
+    def size(self):
+        """
+        Returns the size in bytes of the data for this specialization constant.
+        """
+        return self._spec_const.size
+
+    cdef size_t addressof(self):
+        """
+        Returns the address of the _spec_const for this
+        :class:`.SpecializationConstant` cast to ``size_t``.
+        """
+        return <size_t>&(self._spec_const)
+
+
 cpdef create_kernel_bundle_from_source(SyclQueue q, str src, str copts=""):
     """
         Creates a Sycl interoperability kernel bundle from an OpenCL source
diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h
index 07a76c3fd8..de7543016e 100644
--- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h
+++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h
@@ -35,6 +35,13 @@
 
 DPCTL_C_EXTERN_C_BEGIN
 
+typedef struct DPCTLSpecConstTy
+{
+    uint32_t id;
+    size_t size;
+    const void *value;
+} DPCTLSpecConst;
+
 /**
  * @defgroup KernelBundleInterface Kernel_bundle class C wrapper
  */

From 977f32646fc6cb7807b649360c9c5ce93262b16f Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Tue, 28 Apr 2026 14:55:36 -0700
Subject: [PATCH 18/24] hook specialization constants into kernel bundle
 interface

---
 dpctl/_backend.pxd                            |  4 +-
 dpctl/program/_program.pxd                    |  5 +-
 dpctl/program/_program.pyx                    | 53 +++++++++++++--
 .../dpctl_sycl_kernel_bundle_interface.h      |  6 +-
 .../dpctl_sycl_kernel_bundle_interface.cpp    | 66 +++++++++++++++++--
 5 files changed, 118 insertions(+), 16 deletions(-)

diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd
index d2f1a0decb..d23edaf506 100644
--- a/dpctl/_backend.pxd
+++ b/dpctl/_backend.pxd
@@ -440,7 +440,9 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h":
         const DPCTLSyclDeviceRef Dev,
         const void *IL,
         size_t Length,
-        const char *CompileOpts)
+        const char *CompileOpts,
+        size_t NumSpecConsts,
+        const _spec_const *SpecConsts)
     cdef DPCTLSyclKernelBundleRef DPCTLKernelBundle_CreateFromOCLSource(
         const DPCTLSyclContextRef Ctx,
         const DPCTLSyclDeviceRef Dev,
diff --git a/dpctl/program/_program.pxd b/dpctl/program/_program.pxd
index 435ef68521..41f781cecd 100644
--- a/dpctl/program/_program.pxd
+++ b/dpctl/program/_program.pxd
@@ -63,7 +63,10 @@ cpdef create_kernel_bundle_from_source (
     SyclQueue q, unicode source, unicode copts=*
 )
 cpdef create_kernel_bundle_from_spirv (
-    SyclQueue q, const unsigned char[:] IL, unicode copts=*
+    SyclQueue q,
+    const unsigned char[:] IL,
+    unicode copts=*,
+    list specializations=*,
 )
 cpdef create_program_from_source (SyclQueue q, unicode source, unicode copts=*)
 cpdef create_program_from_spirv (
diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx
index 7cecc3fa4d..50ad8aac02 100644
--- a/dpctl/program/_program.pyx
+++ b/dpctl/program/_program.pyx
@@ -36,6 +36,7 @@ from cpython.buffer cimport (
 )
 from cpython.bytes cimport PyBytes_FromStringAndSize
 from libc.stdint cimport uint32_t
+from libc.stdlib cimport free, malloc
 from libc.string cimport memcmp
 
 import warnings
@@ -469,7 +470,10 @@ cpdef create_kernel_bundle_from_source(SyclQueue q, str src, str copts=""):
 
 
 cpdef create_kernel_bundle_from_spirv(
-    SyclQueue q, const unsigned char[:] IL, str copts=""
+    SyclQueue q,
+    const unsigned char[:] IL,
+    str copts="",
+    list specializations=None,
 ):
     """
         Creates a Sycl interoperability kernel bundle from an SPIR-V binary.
@@ -487,7 +491,9 @@ cpdef create_kernel_bundle_from_spirv(
             copts (str, optional)
                 Optional compilation flags that will be used
                 when compiling the kernel bundle. Default: ``""``.
-
+            specializations (list, optional)
+                A list of :class:`.SpecializationConstant` objects to be used
+                when creating the kernel bundle. Default: ``None``.
         Returns:
             kernel_bundle (:class:`.SyclKernelBundle`)
                 A :class:`.SyclKernelBundle` object wrapping the
@@ -506,11 +512,44 @@ cpdef create_kernel_bundle_from_spirv(
     cdef size_t length = IL.shape[0]
     cdef bytes bCOpts = copts.encode("utf8")
     cdef const char *COpts = <const char*>bCOpts
-    KBref = DPCTLKernelBundle_CreateFromSpirv(
-        CRef, DRef, <const void*>dIL, length, COpts
-    )
-    if KBref is NULL:
-        raise SyclKernelBundleCompilationError()
+    cdef size_t num_spconsts
+    cdef _spec_const *spconsts
+    cdef SpecializationConstant spconst
+
+    if specializations is not None:
+        num_spconsts = len(specializations)
+        spconsts = <_spec_const *>(
+            malloc(num_spconsts * sizeof(_spec_const))
+        )
+        if spconsts == NULL:
+            raise MemoryError(
+                "Failed to allocate memory for specialization constants."
+            )
+        for i, spconst in enumerate(specializations):
+            if not isinstance(spconst, SpecializationConstant):
+                free(spconsts)
+                raise TypeError(
+                    "All items in specializations must be of type "
+                    f"`SpecializationConstant`, got {type(spconst)}"
+                )
+            spconsts[i] = spconst._spec_const
+    else:
+        num_spconsts = 0
+        spconsts = NULL
+    try:
+        KBref = DPCTLKernelBundle_CreateFromSpirv(
+            CRef,
+            DRef,
+            <const void*>dIL,
+            length, COpts,
+            num_spconsts,
+            spconsts,
+        )
+        if KBref is NULL:
+            raise SyclKernelBundleCompilationError()
+    finally:
+        if spconsts != NULL:
+            free(spconsts)
 
     return SyclKernelBundle._create(KBref)
 
diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h
index de7543016e..3909a1c3d1 100644
--- a/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h
+++ b/libsyclinterface/include/syclinterface/dpctl_sycl_kernel_bundle_interface.h
@@ -58,6 +58,8 @@ typedef struct DPCTLSpecConstTy
  * @param    Length         The size of the IL binary in bytes.
  * @param    CompileOpts    Optional compiler flags used when compiling the
  *                          SPIR-V binary.
+ * @param    NumSpecConsts  The number of specialization constants.
+ * @param    SpecConsts     An array of specialization constants.
  * @return   A new SyclKernelBundleRef pointer if the kernel_bundle creation
  * succeeded, else returns NULL.
  * @ingroup KernelBundleInterface
@@ -68,7 +70,9 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef Ctx,
                                   __dpctl_keep const DPCTLSyclDeviceRef Dev,
                                   __dpctl_keep const void *IL,
                                   size_t Length,
-                                  const char *CompileOpts);
+                                  const char *CompileOpts,
+                                  size_t NumSpecConsts,
+                                  const DPCTLSpecConst *SpecConsts);
 
 /*!
  * @brief Create a Sycl kernel bundle from an OpenCL kernel source string.
diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp
index 78c714ecbb..73e5d9ef87 100644
--- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp
+++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp
@@ -31,6 +31,7 @@
 #include "dpctl_error_handlers.h"
 #include "dpctl_sycl_type_casters.hpp"
 #include <CL/cl.h> /* OpenCL headers     */
+#include <cstdint>
 #include <sstream>
 #include <stddef.h>
 #include <sycl/backend/opencl.hpp>
@@ -170,6 +171,21 @@ std::string _GetErrorCode_ocl_impl(cl_int code)
     }
 }
 
+typedef cl_int (*clSetProgramSpecializationConstantFT)(cl_program,
+                                                       cl_uint,
+                                                       size_t,
+                                                       const void *);
+const char *clSetProgramSpecializationConstant_Name =
+    "clSetProgramSpecializationConstant";
+clSetProgramSpecializationConstantFT get_clSetProgramSpecializationConstant()
+{
+    static auto st_clSetProgramSpecializationConstantF =
+        cl_loader::get().getSymbol<clSetProgramSpecializationConstantFT>(
+            clSetProgramSpecializationConstant_Name);
+
+    return st_clSetProgramSpecializationConstantF;
+}
+
 DPCTLSyclKernelBundleRef
 _CreateKernelBundle_common_ocl_impl(cl_program clProgram,
                                     const context &ctx,
@@ -235,7 +251,9 @@ _CreateKernelBundleWithIL_ocl_impl(const context &ctx,
                                    const device &dev,
                                    const void *IL,
                                    size_t il_length,
-                                   const char *CompileOpts)
+                                   const char *CompileOpts,
+                                   size_t NumSpecConsts,
+                                   const DPCTLSpecConst *SpecConsts)
 {
     auto clCreateProgramWithILF = get_clCreateProgramWithIL();
     if (clCreateProgramWithILF == nullptr) {
@@ -257,6 +275,22 @@ _CreateKernelBundleWithIL_ocl_impl(const context &ctx,
         return nullptr;
     }
 
+    if (SpecConsts != nullptr && NumSpecConsts > 0) {
+        auto clSetProgramSpecConstF = get_clSetProgramSpecializationConstant();
+        if (clSetProgramSpecConstF) {
+            for (size_t i = 0; i < NumSpecConsts; ++i) {
+                clSetProgramSpecConstF(clProgram, SpecConsts[i].id,
+                                       SpecConsts[i].size, SpecConsts[i].value);
+            }
+        }
+        else {
+            error_handler("clSetProgramSpecializationConstant is not available "
+                          "in the OpenCL implementation.",
+                          __FILE__, __func__, __LINE__);
+            return nullptr;
+        }
+    }
+
     return _CreateKernelBundle_common_ocl_impl(clProgram, ctx, dev,
                                                CompileOpts);
 }
@@ -428,7 +462,9 @@ _CreateKernelBundleWithIL_ze_impl(const context &SyclCtx,
                                   const device &SyclDev,
                                   const void *IL,
                                   size_t il_length,
-                                  const char *CompileOpts)
+                                  const char *CompileOpts,
+                                  size_t NumSpecConsts,
+                                  const DPCTLSpecConst *SpecConsts)
 {
     auto zeModuleCreateFn = get_zeModuleCreate();
     if (zeModuleCreateFn == nullptr) {
@@ -444,8 +480,22 @@ _CreateKernelBundleWithIL_ze_impl(const context &SyclCtx,
     ZeDevice = get_native<ze_be>(SyclDev);
 
     // Specialization constants are not supported by DPCTL at the moment
+    std::vector<std::uint32_t> spec_ids;
+    std::vector<const void *> spec_values;
+
+    if (SpecConsts != nullptr && NumSpecConsts > 0) {
+        spec_ids.reserve(NumSpecConsts);
+        spec_values.reserve(NumSpecConsts);
+        for (size_t i = 0; i < NumSpecConsts; ++i) {
+            spec_ids.push_back(SpecConsts[i].id);
+            spec_values.push_back(SpecConsts[i].value);
+        }
+    }
     ze_module_constants_t ZeSpecConstants = {};
-    ZeSpecConstants.numConstants = 0;
+    ZeSpecConstants.numConstants = static_cast<std::uint32_t>(NumSpecConsts);
+    ZeSpecConstants.pConstantIds = spec_ids.empty() ? nullptr : spec_ids.data();
+    ZeSpecConstants.pConstantValues =
+        spec_values.empty() ? nullptr : spec_values.data();
 
     // Populate the Level Zero module descriptions
     ze_module_desc_t ZeModuleDesc = {};
@@ -583,7 +633,9 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef,
                                   __dpctl_keep const DPCTLSyclDeviceRef DevRef,
                                   __dpctl_keep const void *IL,
                                   size_t length,
-                                  const char *CompileOpts)
+                                  const char *CompileOpts,
+                                  size_t NumSpecConsts,
+                                  const DPCTLSpecConst *SpecConsts)
 {
     DPCTLSyclKernelBundleRef KBRef = nullptr;
     if (!CtxRef) {
@@ -611,12 +663,14 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef,
     switch (BE) {
     case backend::opencl:
         KBRef = _CreateKernelBundleWithIL_ocl_impl(*SyclCtx, *SyclDev, IL,
-                                                   length, CompileOpts);
+                                                   length, CompileOpts,
+                                                   NumSpecConsts, SpecConsts);
         break;
     case backend::ext_oneapi_level_zero:
 #ifdef DPCTL_ENABLE_L0_PROGRAM_CREATION
         KBRef = _CreateKernelBundleWithIL_ze_impl(*SyclCtx, *SyclDev, IL,
-                                                  length, CompileOpts);
+                                                  length, CompileOpts,
+                                                  NumSpecConsts, SpecConsts);
         break;
 #endif
     default:

From 0a680788ce0f96e6cc4a7130250b8ce172cf20e3 Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Tue, 28 Apr 2026 15:29:29 -0700
Subject: [PATCH 19/24] add  test for SpecializationConstant use in kernel

---
 .../specialization_constant_kernel.spv        | Bin 0 -> 2288 bytes
 dpctl/tests/test_sycl_program.py              |  38 ++++++++++++++++++
 2 files changed, 38 insertions(+)
 create mode 100644 dpctl/tests/input_files/specialization_constant_kernel.spv

diff --git a/dpctl/tests/input_files/specialization_constant_kernel.spv b/dpctl/tests/input_files/specialization_constant_kernel.spv
new file mode 100644
index 0000000000000000000000000000000000000000..d696fa755f7e56d61ceeab707c468f0c937b8ed6
GIT binary patch
literal 2288
zcmZ{kTTc^F5XVoUih_cOco&O!Ll6NGwRj6qN~MBQgNZMjrQPTzWlP${fbqd6KZmcr
z`Q=R11mpj=yT@!uJ<Vk9b7uZKr=1-mT_&3`z2>laW#ZFsI!sm~{7}LVC%h-&*Aspt
z;roNebek?ydAZ@NEU!D>?zUGA9k<om+42<Aqe>a^wQVn0C{MJ)n$ew|y1_)$XNJ^)
zF(*aN+t-!Vnbuyle%INUnyq=E>({4e)*CxP&8w}u!8>ooYX)BZjUT>WSuU+sCY>2S
z^tMW6f2vU4oGKIwNYj(^Zp*J$)X734XoYc7p`a_Gxwx-SoxvAH+}l1Uj_*{z;PUlL
zM)^#n9H+JIH$OP@JAOSZ1;u(})2)|+4~?oD`i-EZF%@?-x}$Au#s3IChs+WEjH%WI
z*#xua^BO*%Z!#|aZdGaHkI~8J-DcC>Bkq_v{x8mP-g$x7bVIM^w8O%fqN15&RMfV=
z1GDfAi^74N(~+EwWcK1>q{n#vw_Y-@tpoP94l;EvM|HAhK>WIh9VrTV^l=e5p`Q{@
z`_a?>ymap2(DP&Qo&GqtAemUcC$*Li=w<2Y+%@Us4l0&@o{2c<z2fx5LC=XpoH3z8
ze7DHfvo_|mXDrDZP7GG(eUWE#9>o67u<0a+IXPCRcg`94JH@#lo|DeJ)I%STr%U{z
ztl)0@9Fmt^!-mXk=;R+0CkGpPpE$b!21C->J!}yE6XL{UA6BFwjCAVYBj#LU<9=$e
zN5~NVOX6{#v3^y0e815SbeP*o+`Krw<V4)HAWj~3Xil}5pO{QaXYbHQH7m6yM6-$_
zm;WOUeczX#zS%=$`e)~Ust5i3nNO<DEayeUJl3mch6{-enZAi*rXdmeONt}!N@8O#
zhDA1iB=On&UFCzpH4#2rkGGJncT+kx>fI38db~}12kYIE4Igpr$CwDLh|7syNkbkN
z;p2QuX1pUA-0*RqtuZMXd=J)`k_{in)|i$|4Sc+p@4AogaaTlraQvl8{9k}M{D;00
zweN|D;XBNT&WP}le_uome6u2azl9-o<g+5`A=^72N~R|7?ayF`L0?io;65h;d;H9q
z6MrPaPA}}(XT|2k`Om?}8@2gQB!e~iPZWnvTsjx|sfZYGv3=1$cCfIxF|W0QK^h}w
z+t+qK5ZRqxl#RptV2|g;3n_U?GI$il?RQ+3+-?_2vg6~}ex6I_JK|#(Z0{?Q@mWmK
n!G!x?i0A>o#dI~Xr*X;2wkE<?7TNrlksRlPDLQdB*JAJ&rApd7

literal 0
HcmV?d00001

diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py
index 1c09adc28e..7b5f8db20e 100644
--- a/dpctl/tests/test_sycl_program.py
+++ b/dpctl/tests/test_sycl_program.py
@@ -18,6 +18,7 @@
 
 import os
 
+import numpy as np
 import pytest
 
 import dpctl
@@ -262,3 +263,40 @@ def test_create_kernel_bundle_from_invalid_src_ocl():
     }"
     with pytest.raises(dpctl_prog.SyclKernelBundleCompilationError):
         dpctl_prog.create_kernel_bundle_from_source(q, invalid_oclSrc)
+
+
+def test_create_kernel_bundle_with_spec_const():
+    try:
+        q = dpctl.SyclQueue()
+    except dpctl.SyclQueueCreationError:
+        pytest.skip("Could not create default queue")
+
+    spec_id = 0
+    sp = dpctl_prog.SpecializationConstant(spec_id, "i4", 42)
+
+    spirv_file = get_spirv_abspath("specialization_constant_kernel.spv")
+    with open(spirv_file, "br") as spv:
+        spv_bytes = spv.read()
+
+    kb = dpctl_prog.create_kernel_bundle_from_spirv(
+        q, spv_bytes, specializations=[sp]
+    )
+    kernel = kb.get_sycl_kernel("_ZTS20BasicSpecConstKernel")
+
+    n = 128
+    x = np.ones(n, dtype="i4")
+    y = np.zeros_like(x)
+
+    x_usm = dpctl.memory.MemoryUSMDevice(x.nbytes, queue=q)
+    y_usm = dpctl.memory.MemoryUSMDevice(y.nbytes, queue=q)
+
+    e1 = q.memcpy_async(x_usm, x, x.nbytes)
+    e2 = q.submit(kernel, [x_usm, y_usm], [n], dEvents=[e1])
+    e3 = q.memcpy_async(y, y_usm, y.nbytes, [e2])
+
+    ht_e = q._submit_keep_args_alive([x_usm], [e3])
+
+    e3.wait()
+    ht_e.wait()
+
+    assert np.all(y == 43)

From aba560be460cb5a1c34c5ca22f4aa15c10cd285a Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Tue, 28 Apr 2026 16:25:17 -0700
Subject: [PATCH 20/24] fix libsyclinterface tests

---
 .../test_sycl_kernel_bundle_interface.cpp     | 24 +++++++++++--------
 .../tests/test_sycl_queue_submit.cpp          |  6 +++--
 ...t_sycl_queue_submit_local_accessor_arg.cpp |  6 +++--
 .../test_sycl_queue_submit_raw_kernel_arg.cpp |  6 +++--
 ...ycl_queue_submit_work_group_memory_arg.cpp |  6 +++--
 5 files changed, 30 insertions(+), 18 deletions(-)

diff --git a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp
index a835c277b9..5793e983d9 100644
--- a/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp
+++ b/libsyclinterface/tests/test_sycl_kernel_bundle_interface.cpp
@@ -69,7 +69,8 @@ struct TestDPCTLSyclKernelBundleInterface
             spirvFile.seekg(0, std::ios::beg);
             spirvFile.read(spirvBuffer.data(), spirvFileSize);
             KBRef = DPCTLKernelBundle_CreateFromSpirv(
-                CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr);
+                CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr, 0,
+                nullptr);
         }
     }
 
@@ -132,18 +133,21 @@ TEST_P(TestDPCTLSyclKernelBundleInterface, ChkCreateFromSpirvNull)
     const void *null_spirv = nullptr;
     DPCTLSyclKernelBundleRef KBRef = nullptr;
     // Null context
-    EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv(
-                                Null_CRef, Null_DRef, null_spirv, 0, nullptr));
+    EXPECT_NO_FATAL_FAILURE(
+        KBRef = DPCTLKernelBundle_CreateFromSpirv(
+            Null_CRef, Null_DRef, null_spirv, 0, nullptr, 0, nullptr));
     ASSERT_TRUE(KBRef == nullptr);
 
     // Null device
-    EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv(
-                                CRef, Null_DRef, null_spirv, 0, nullptr));
+    EXPECT_NO_FATAL_FAILURE(
+        KBRef = DPCTLKernelBundle_CreateFromSpirv(CRef, Null_DRef, null_spirv,
+                                                  0, nullptr, 0, nullptr));
     ASSERT_TRUE(KBRef == nullptr);
 
     // Null IL
-    EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv(
-                                CRef, DRef, null_spirv, 0, nullptr));
+    EXPECT_NO_FATAL_FAILURE(
+        KBRef = DPCTLKernelBundle_CreateFromSpirv(CRef, DRef, null_spirv, 0,
+                                                  nullptr, 0, nullptr));
     ASSERT_TRUE(KBRef == nullptr);
 }
 
@@ -350,8 +354,8 @@ TEST_F(TestKernelBundleUnsupportedBackend, CheckCreateFromSpirv)
     spirvFile.close();
 
     DPCTLSyclKernelBundleRef KBRef = nullptr;
-    EXPECT_NO_FATAL_FAILURE(
-        KBRef = DPCTLKernelBundle_CreateFromSpirv(
-            CRef, DRef, spirvBuffer.data(), spirvFileSize, nullptr));
+    EXPECT_NO_FATAL_FAILURE(KBRef = DPCTLKernelBundle_CreateFromSpirv(
+                                CRef, DRef, spirvBuffer.data(), spirvFileSize,
+                                nullptr, 0, nullptr));
     ASSERT_TRUE(KBRef == nullptr);
 }
diff --git a/libsyclinterface/tests/test_sycl_queue_submit.cpp b/libsyclinterface/tests/test_sycl_queue_submit.cpp
index ab5b6bef82..f2fc2b2140 100644
--- a/libsyclinterface/tests/test_sycl_queue_submit.cpp
+++ b/libsyclinterface/tests/test_sycl_queue_submit.cpp
@@ -242,7 +242,8 @@ struct TestQueueSubmit : public ::testing::Test
         auto CRef = DPCTLQueue_GetContext(QRef);
 
         KBRef = DPCTLKernelBundle_CreateFromSpirv(
-            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr);
+            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0,
+            nullptr);
         DPCTLDevice_Delete(DRef);
         DPCTLDeviceSelector_Delete(DSRef);
     }
@@ -282,7 +283,8 @@ struct TestQueueSubmitFP64 : public ::testing::Test
         auto CRef = DPCTLQueue_GetContext(QRef);
 
         KBRef = DPCTLKernelBundle_CreateFromSpirv(
-            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr);
+            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0,
+            nullptr);
         DPCTLDeviceSelector_Delete(DSRef);
     }
 
diff --git a/libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp b/libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp
index f6110375bb..8f1b97d1d4 100644
--- a/libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp
+++ b/libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp
@@ -237,7 +237,8 @@ struct TestQueueSubmitWithLocalAccessor : public ::testing::Test
         auto CRef = DPCTLQueue_GetContext(QRef);
 
         KBRef = DPCTLKernelBundle_CreateFromSpirv(
-            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr);
+            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0,
+            nullptr);
         DPCTLDevice_Delete(DRef);
         DPCTLDeviceSelector_Delete(DSRef);
     }
@@ -276,7 +277,8 @@ struct TestQueueSubmitWithLocalAccessorFP64 : public ::testing::Test
         auto CRef = DPCTLQueue_GetContext(QRef);
 
         KBRef = DPCTLKernelBundle_CreateFromSpirv(
-            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr);
+            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0,
+            nullptr);
         DPCTLDeviceSelector_Delete(DSRef);
     }
 
diff --git a/libsyclinterface/tests/test_sycl_queue_submit_raw_kernel_arg.cpp b/libsyclinterface/tests/test_sycl_queue_submit_raw_kernel_arg.cpp
index f40bc20066..04d3958d8d 100644
--- a/libsyclinterface/tests/test_sycl_queue_submit_raw_kernel_arg.cpp
+++ b/libsyclinterface/tests/test_sycl_queue_submit_raw_kernel_arg.cpp
@@ -262,7 +262,8 @@ struct TestQueueSubmitWithRawKernelArg : public ::testing::Test
         auto CRef = DPCTLQueue_GetContext(QRef);
 
         KBRef = DPCTLKernelBundle_CreateFromSpirv(
-            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr);
+            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0,
+            nullptr);
         DPCTLDevice_Delete(DRef);
         DPCTLDeviceSelector_Delete(DSRef);
     }
@@ -301,7 +302,8 @@ struct TestQueueSubmitWithRawKernelArgFP64 : public ::testing::Test
         auto CRef = DPCTLQueue_GetContext(QRef);
 
         KBRef = DPCTLKernelBundle_CreateFromSpirv(
-            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr);
+            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0,
+            nullptr);
         DPCTLDeviceSelector_Delete(DSRef);
     }
 
diff --git a/libsyclinterface/tests/test_sycl_queue_submit_work_group_memory_arg.cpp b/libsyclinterface/tests/test_sycl_queue_submit_work_group_memory_arg.cpp
index d0f44b7275..d1d1f69bfa 100644
--- a/libsyclinterface/tests/test_sycl_queue_submit_work_group_memory_arg.cpp
+++ b/libsyclinterface/tests/test_sycl_queue_submit_work_group_memory_arg.cpp
@@ -262,7 +262,8 @@ struct TestQueueSubmitWithWorkGroupMemory : public ::testing::Test
         auto CRef = DPCTLQueue_GetContext(QRef);
 
         KBRef = DPCTLKernelBundle_CreateFromSpirv(
-            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr);
+            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0,
+            nullptr);
         DPCTLDevice_Delete(DRef);
         DPCTLDeviceSelector_Delete(DSRef);
     }
@@ -301,7 +302,8 @@ struct TestQueueSubmitWithWorkGroupMemoryFP64 : public ::testing::Test
         auto CRef = DPCTLQueue_GetContext(QRef);
 
         KBRef = DPCTLKernelBundle_CreateFromSpirv(
-            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr);
+            CRef, DRef, spirvBuffer_.data(), spirvFileSize_, nullptr, 0,
+            nullptr);
         DPCTLDeviceSelector_Delete(DSRef);
     }
 

From d8b0afbeb07d819bee1fd8766e0fb0e43e6e4aaf Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Thu, 30 Apr 2026 11:35:16 -0700
Subject: [PATCH 21/24] add test for composite specialization constant

also removes "v" as a permitted specialization constant intermediate data type, as composite specialization constants are broken into multiple specialization constants, so structs end up passed as a single constant while the program expects multiple, and therefore, doesn't work as intended
---
 dpctl/program/_program.pyx                    |  12 ++---
 .../specialization_constant_composite.spv     | Bin 0 -> 3432 bytes
 dpctl/tests/test_sycl_program.py              |  41 ++++++++++++++++++
 3 files changed, 47 insertions(+), 6 deletions(-)
 create mode 100644 dpctl/tests/input_files/specialization_constant_composite.spv

diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx
index 50ad8aac02..166d46a3e3 100644
--- a/dpctl/program/_program.pyx
+++ b/dpctl/program/_program.pyx
@@ -282,8 +282,8 @@ cdef class SpecializationConstant:
       If the constructor is invoked with a single variadic argument, the
       argument is expected to either expose the Python buffer protocol or be
       coercible to a NumPy array. If the argument is coercible to a NumPy array
-      or is one, it must have a supported data type (bool, integral, floating
-      point, or void). The specialization constant will be constructed from the
+      or is one, it must have a supported data type (bool, integral, or
+      floating point). The specialization constant will be constructed from the
       data in the buffer
 
     - ``SpecializationConstant(spec_id, dtype, obj)``
@@ -291,7 +291,7 @@ cdef class SpecializationConstant:
     argument is a string, it is interpreted as a NumPy ``dtype`` string and the
     second argument will be coerced to a NumPy array with that data type.
     The data type specified by the first argument must be a supported data
-    type (bool, integral, floating point, or void).
+    type (bool, integral, or floating point).
 
     - ``SpecializationConstant(spec_id, nbytes, raw_ptr)``
       If the constructor is invoked with two variadic arguments where both are
@@ -360,13 +360,13 @@ cdef class SpecializationConstant:
             )
 
         if isinstance(target_obj, np.ndarray):
-            if target_obj.dtype.kind not in ("b", "i", "u", "f", "c", "V"):
+            if target_obj.dtype.kind not in ("b", "i", "u", "f", "c"):
                 raise TypeError(
                     "Coercion of input to buffer resulted in an unsupported "
                     f"data type '{target_obj.dtype}'. When coercing objects, "
                     "`SpecializationConstant` expects the data to coerce to a "
-                    "supported type: bool, integral, real or complex floating "
-                    "point, or void. To pass arbitrary data, use a "
+                    "supported type: bool, integral, or real or complex "
+                    "floating point. To pass arbitrary data, use a "
                     "`memoryview` or `bytes` object, or pass the pointer and "
                     "size directly."
                 )
diff --git a/dpctl/tests/input_files/specialization_constant_composite.spv b/dpctl/tests/input_files/specialization_constant_composite.spv
new file mode 100644
index 0000000000000000000000000000000000000000..6b69617a7dc599495bc6a43862832b1ef77eab0d
GIT binary patch
literal 3432
zcma)-ZBHCk6vt=FQYyt_ZE1aIk=oLNR4fnDx)numffb~brJzmPVb~pD?LOe{fPgir
zFZ6TxR$uu={5U3RQse(Ob5Ge2ns}4{p2vI6x#ylc+j8P+o0m#?XS|c%4_>m)dMCWJ
zSZJ%GCmlWQ=ygYLIQqOR@&Dt3=e2vMy!_&dpIcbY&ifnN{zlM<{7MjQ`ITBRl)Wvs
z&&@9R;of#wi2R_@*r|kyby^uyqW8AL>eT#TBPx2Ttwoi1F*?0&RqA=K3;x69{K9Bs
zzfd0WSBA!lVHA|g!=p>JooX>GE(O)ia4xJ@!}4k=+M1iqE#z<cqopXU<mO95+4=RM
zY&HuwJd}^>JB28(Ql@IvMwINzW|f2Uldn^iL4BaSFWD3Ox2RtF_sh>I$(%@g{Kj^v
zzUxo!l*&=AI#aH#2jyIKw^j(EQmvZPG3ED4lH)kveCdGxw5qqKL<?oLY>=~OGCF)F
zQ|GwY+ZE*)d(xduCaBkgef&DSS6}$~{$^MW>p>J2{bPF(=Y3QDjkr^T%RALlb#pPO
z2bC}i>kV&4(Ku(n@PE$A9nYvw?9s@GPscp-bt#sUe!n!H>xaeRZY&Nwao&mJq`h;Z
z*9B&XJjgc%`iOj9wAqf_Y`^Se?iBx7QS!2obE3#eEotuy$6LNd$$dhH&Wr!*5qV8A
z`7C~LL@peWzmeRW|A~_=f7!{FzahCf|3Q=7rL6c(2rT3?qV$G^{FZ1Bg4Td;3IBG%
zYB1e39NZdk`+ASa1CrTi_K(FA&thc7Ydzy<*u|jVEM#h<zbs-PFU8WKKmB2uKmA}`
zlz)pT`Jg?L={fZvpOd9c^pdp9kDaqyR^AjkczT1({ti+0phNByWu~dCTQYA99l-vY
zDE{cL$n&g|iG%He;=e80BXrBoo1_2eh)E3I2QjcAcRQJS$$3Q-9$<f06o2%{eWKJ$
z%=7XZa{4<?Z*en{*_YgZ+8cMW)$ys5lX}zMBgvc{nX^51cJ$z`a?@{oQcW`N6}ew$
zr<UOu={OE=i#4JzBX!LQvjTeNfEgg)1A)bzBR`3y!+jOW%msSpVN6wUHvUIg+{10<
z=N|Yzz;hpb3;t4%^!FCmPwy`a?0u?S^q%hwd&X5`((^<IPoMGQ4z3Eso0cE@_~xMF
z-Si5!|Au(`9<hH<`};)u1#A}YJ@L))Zc0W+yaB=Dy)Pb{)noAnrNf3F^Z$Xs-PjrK
zOM<`UY$D#T?u;YixhHJovlye|xzD39K9mj{%VLa)CkD25;dkZZ{%;G!?-BNOcINUU
zfxgo_=JKwSt#|0hrTe5w&$k6Vb2`zar#FAALjEtfTj-y<QEHwP$jyC!CR`M-QP<}J
zwPBkQu>GU@(8IqW5D(tYGA*8%S>fNV6RIb<KWd#3IFpy+e}Vr8>TiwDIwzib(!#WK
z-2c453=ofgGv{9l=*<^B?{z_-N2aeS68`s`?I*=VUKEIdp1r*7`vQHmy}!A8$vfMU
zsO4W4KOvYecT3zA0sjXAK0msAtK!Lt-OlxucxG`@;QYw=Sxn}getj*_uax=)PrshY
zAHRnJ`^N<9J$Hn@N3faWPEF4}H_rk0_gr71JA!>5eCb$^1-=Q7MAw@5Ks<N8F3R_f
z??XXo6X;QkXi+%sU4_mj@}j;C@w^Fa*91Gurg&^NN67Xzw#1Xu`c!f<a|Lg6g<nOm
zzS4irQx(<(;@iE}#G^m-P2SzMbmXG8UIBmR?OOqx`P;nJoz3R0A(`0d?afBwEhcZ4
zn9SjhfX!ms9PT<Bb7=4TJMk97zJ1K$o`B6_*nI9wM;zu8ADhnu@zjXj`t!Yb?wS5D
WpUC)GOzZno>8OGEgs1N|PyYdNu{1IO

literal 0
HcmV?d00001

diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py
index 7b5f8db20e..61b9fee843 100644
--- a/dpctl/tests/test_sycl_program.py
+++ b/dpctl/tests/test_sycl_program.py
@@ -300,3 +300,44 @@ def test_create_kernel_bundle_with_spec_const():
     ht_e.wait()
 
     assert np.all(y == 43)
+
+
+def test_create_kernel_bundle_with_composite_spec_const():
+    try:
+        q = dpctl.SyclQueue()
+    except dpctl.SyclQueueCreationError:
+        pytest.skip("Could not create default queue")
+
+    # composite specialization constants are separated into individual
+    # specialization constants with unique spec_ids
+    sp1 = dpctl_prog.SpecializationConstant(0, "i4", 10)
+    sp2 = dpctl_prog.SpecializationConstant(1, "f4", 2.5)
+    sp3 = dpctl_prog.SpecializationConstant(2, "?", 1)
+
+    spirv_file = get_spirv_abspath("specialization_constant_composite.spv")
+    with open(spirv_file, "br") as spv:
+        spv_bytes = spv.read()
+
+    kb = dpctl_prog.create_kernel_bundle_from_spirv(
+        q, spv_bytes, specializations=[sp1, sp2, sp3]
+    )
+    kernel = kb.get_sycl_kernel("_ZTS21StructSpecConstKernel")
+
+    n = 128
+    x = np.ones(n, dtype="f4")
+    y = np.zeros_like(x)
+
+    x_usm = dpctl.memory.MemoryUSMDevice(x.nbytes, queue=q)
+    y_usm = dpctl.memory.MemoryUSMDevice(y.nbytes, queue=q)
+
+    e1 = q.memcpy_async(x_usm, x, x.nbytes)
+    e2 = q.submit(kernel, [x_usm, y_usm], [n], dEvents=[e1])
+    e3 = q.memcpy_async(y, y_usm, y.nbytes, [e2])
+
+    ht_e = q._submit_keep_args_alive([x_usm], [e3])
+
+    e3.wait()
+    ht_e.wait()
+
+    # 1.0 * 10 + 2.5 = 12.5
+    assert np.all(y == 12.5)

From 020ec2b9205b702826b61d1df12946df3dfcdc07 Mon Sep 17 00:00:00 2001
From: "dependabot[bot]" <49699333+dependabot[bot]@users.noreply.github.com>
Date: Mon, 4 May 2026 17:39:22 +0000
Subject: [PATCH 22/24] Bump github/codeql-action from 4.35.2 to 4.35.3

Bumps [github/codeql-action](https://github.com/github/codeql-action) from 4.35.2 to 4.35.3.
- [Release notes](https://github.com/github/codeql-action/releases)
- [Changelog](https://github.com/github/codeql-action/blob/main/CHANGELOG.md)
- [Commits](https://github.com/github/codeql-action/compare/95e58e9a2cdfd71adc6e0353d5c52f41a045d225...e46ed2cbd01164d986452f91f178727624ae40d7)

---
updated-dependencies:
- dependency-name: github/codeql-action
  dependency-version: 4.35.3
  dependency-type: direct:production
  update-type: version-update:semver-patch
...

Signed-off-by: dependabot[bot] <support@github.com>
---
 .github/workflows/openssf-scorecard.yml | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/.github/workflows/openssf-scorecard.yml b/.github/workflows/openssf-scorecard.yml
index 5e6b9ec8b0..cc512f59be 100644
--- a/.github/workflows/openssf-scorecard.yml
+++ b/.github/workflows/openssf-scorecard.yml
@@ -69,6 +69,6 @@ jobs:
 
       # Upload the results to GitHub's code scanning dashboard.
       - name: "Upload to code-scanning"
-        uses: github/codeql-action/upload-sarif@95e58e9a2cdfd71adc6e0353d5c52f41a045d225 # v4.35.2
+        uses: github/codeql-action/upload-sarif@e46ed2cbd01164d986452f91f178727624ae40d7 # v4.35.3
         with:
           sarif_file: results.sarif

From d2c6bb6707ff954fade21c0897e74cf14d68c009 Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Mon, 4 May 2026 20:27:13 -0700
Subject: [PATCH 23/24] add program.utils namespace with SPIRV parser

---
 dpctl/program/__init__.py       |   5 ++
 dpctl/program/utils/__init__.py |  25 ++++++++
 dpctl/program/utils/_utils.py   | 106 ++++++++++++++++++++++++++++++++
 3 files changed, 136 insertions(+)
 create mode 100644 dpctl/program/utils/__init__.py
 create mode 100644 dpctl/program/utils/_utils.py

diff --git a/dpctl/program/__init__.py b/dpctl/program/__init__.py
index bb16f9611b..e1e625ff7a 100644
--- a/dpctl/program/__init__.py
+++ b/dpctl/program/__init__.py
@@ -45,6 +45,11 @@
     "SpecializationConstant",
 ]
 
+# add submodules
+__all__ += [
+    "utils",
+]
+
 
 def __getattr__(name):
     if name == "SyclProgram":
diff --git a/dpctl/program/utils/__init__.py b/dpctl/program/utils/__init__.py
new file mode 100644
index 0000000000..474f154f95
--- /dev/null
+++ b/dpctl/program/utils/__init__.py
@@ -0,0 +1,25 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2025 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""
+A collection of utility functions for dpctl.program module.
+"""
+
+from ._utils import parse_spirv_specializations
+
+__all__ = [
+    "parse_spirv_specializations",
+]
diff --git a/dpctl/program/utils/_utils.py b/dpctl/program/utils/_utils.py
new file mode 100644
index 0000000000..2a16802dea
--- /dev/null
+++ b/dpctl/program/utils/_utils.py
@@ -0,0 +1,106 @@
+#                      Data Parallel Control (dpctl)
+#
+# Copyright 2020-2025 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#    http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""Implements various utilities for the dpctl.program module."""
+
+from enum import IntEnum
+
+import numpy as np
+
+
+class SpirvOpCode(IntEnum):
+    OpName = 5
+    OpTypeBool = 20
+    OpTypeInt = 21
+    OpTypeFloat = 22
+    OpSpecConstantTrue = 48
+    OpSpecConstantFalse = 49
+    OpSpecConstant = 50
+    OpDecorate = 71
+
+
+class SpirvDecoration(IntEnum):
+    SpecId = 1
+
+
+def parse_spirv_specializations(
+    spv_bytes: bytes | bytearray | memoryview,
+) -> dict[int, dict[str, str]]:
+    words = np.frombuffer(spv_bytes, dtype=np.uint32)
+
+    # verify magic number
+    if len(words) < 5 or words[0] != 0x07230203:
+        raise ValueError("Invalid SPIR-V binary")
+
+    types = {}
+    ids = {}
+    names = {}
+    constants = {}
+
+    i = 5  # skip 5 word header
+    while i < len(words):
+        word = words[i]
+        opcode = word & 0xFFFF
+        word_count = word >> 16
+
+        if word_count == 0:
+            raise ValueError(f"Invalid SPIR-V instruction at word index {i}")
+
+        if opcode == SpirvOpCode.OpTypeBool:
+            result_id = int(words[i + 1])
+            types[result_id] = "?"
+        elif opcode == SpirvOpCode.OpTypeInt:
+            result_id = int(words[i + 1])
+            width = int(words[i + 2])
+            signed = int(words[i + 3])
+            prefix = "i" if signed else "u"
+            types[result_id] = f"{prefix}{width // 8}"
+        elif opcode == SpirvOpCode.OpTypeFloat:
+            result_id = int(words[i + 1])
+            width = int(words[i + 2])
+            types[result_id] = f"f{width // 8}"
+        elif opcode in (
+            SpirvOpCode.OpSpecConstant,
+            SpirvOpCode.OpSpecConstantTrue,
+            SpirvOpCode.OpSpecConstantFalse,
+        ):
+            type_id = int(words[i + 1])
+            result_id = int(words[i + 2])
+            constants[result_id] = type_id
+        elif opcode == SpirvOpCode.OpDecorate:
+            target_id = int(words[i + 1])
+            decoration = int(words[i + 2])
+            if decoration == SpirvDecoration.SpecId:
+                ids[target_id] = int(words[i + 3])
+        elif opcode == SpirvOpCode.OpName:
+            target_id = int(words[i + 1])
+            name_bytes = words[i + 2 : i + word_count].tobytes()
+            names[target_id] = name_bytes.split(b"\x00", 1)[0].decode("utf-8")
+
+        i += word_count
+
+    result = {}
+    for target_id, spec_id in ids.items():
+        type_id = constants.get(target_id)
+        dtype_str = types.get(type_id, "unknown_type")
+        name = names.get(target_id, f"unnamed_spec_const_{spec_id}")
+
+        result[spec_id] = {
+            "name": name,
+            "dtype": dtype_str,
+        }
+
+    return result

From 043aa83af7e37e87334270d5963b8b78403bdf80 Mon Sep 17 00:00:00 2001
From: Nikita Grigorian <nikita.grigorian@intel.com>
Date: Mon, 4 May 2026 23:51:48 -0700
Subject: [PATCH 24/24] Refactor specialiazation constants to use dataclass

also adds spec_id, itemsize, and default_value fields
---
 dpctl/program/utils/_utils.py                 | 122 +++++++++++++++---
 .../specialization_constant_composite.spv     | Bin 3432 -> 3504 bytes
 dpctl/tests/test_sycl_program.py              |  47 +++++++
 setup.py                                      |   1 +
 4 files changed, 152 insertions(+), 18 deletions(-)

diff --git a/dpctl/program/utils/_utils.py b/dpctl/program/utils/_utils.py
index 2a16802dea..86df0855ad 100644
--- a/dpctl/program/utils/_utils.py
+++ b/dpctl/program/utils/_utils.py
@@ -16,6 +16,7 @@
 
 """Implements various utilities for the dpctl.program module."""
 
+from dataclasses import dataclass
 from enum import IntEnum
 
 import numpy as np
@@ -29,6 +30,7 @@ class SpirvOpCode(IntEnum):
     OpSpecConstantTrue = 48
     OpSpecConstantFalse = 49
     OpSpecConstant = 50
+    OpFunction = 54
     OpDecorate = 71
 
 
@@ -36,9 +38,52 @@ class SpirvDecoration(IntEnum):
     SpecId = 1
 
 
+@dataclass(frozen=True)
+class SpecializationConstantInfo:
+    """Data class representing specialization constant information."""
+
+    spec_id: int
+    dtype: str
+    name: str
+    itemsize: int
+    default_value: int | float | bool | None
+
+
 def parse_spirv_specializations(
     spv_bytes: bytes | bytearray | memoryview,
-) -> dict[int, dict[str, str]]:
+) -> tuple[SpecializationConstantInfo]:
+    """
+    Parses SPIR-V byte stream to extract information about specializations,
+    including the specialization IDs, types, names, and default values.
+
+    Note that the dtype information may be imprecise, as the compiler may
+    choose to, for example, represent a bool as char, or may represent both
+    signed and unsigned integers as unsigned integer bit buckets of the same
+    length.
+
+    Args:
+        spv_bytes (bytes | bytearray | memoryview):
+            the SPIR-V byte stream.
+
+    Returns:
+        tuple[SpecializationConstantInfo]:
+            a tuple of parsed constants and their information represented by
+            `SpecializationConstantInfo` objects, sorted by their
+            specialization IDs. The length of the tuple is equal to the number
+            of specialization constants found. Each
+            `SpecializationConstantInfo` object contains the following
+            attributes:
+
+            - `spec_id` (int): The specialization ID.
+            - `dtype` (str): A NumPy style string representing the data type.
+            - `itemsize` (int): The size of the specialization constant in
+                bytes.
+            - `name` (str): The variable name. If not preserved in the binary,
+                a default name in the format `unnamed_spec_const_{spec_id}` is
+                used.
+            - `default_value` (int | float | bool | None): The default value of
+                the specialization constant. If not specified, `None` is used.
+    """
     words = np.frombuffer(spv_bytes, dtype=np.uint32)
 
     # verify magic number
@@ -49,6 +94,7 @@ def parse_spirv_specializations(
     ids = {}
     names = {}
     constants = {}
+    defaults = {}
 
     i = 5  # skip 5 word header
     while i < len(words):
@@ -59,27 +105,45 @@ def parse_spirv_specializations(
         if word_count == 0:
             raise ValueError(f"Invalid SPIR-V instruction at word index {i}")
 
-        if opcode == SpirvOpCode.OpTypeBool:
+        if opcode == SpirvOpCode.OpFunction:
+            # everything following is not relevant to specialization constant
+            # parsing, so we can stop parsing at this point
+            break
+        elif opcode == SpirvOpCode.OpTypeBool:
             result_id = int(words[i + 1])
-            types[result_id] = "?"
+            types[result_id] = {"dtype": "?", "itemsize": 1}
         elif opcode == SpirvOpCode.OpTypeInt:
             result_id = int(words[i + 1])
             width = int(words[i + 2])
             signed = int(words[i + 3])
             prefix = "i" if signed else "u"
-            types[result_id] = f"{prefix}{width // 8}"
+            types[result_id] = {
+                "dtype": f"{prefix}{width // 8}",
+                "itemsize": width // 8,
+            }
         elif opcode == SpirvOpCode.OpTypeFloat:
             result_id = int(words[i + 1])
             width = int(words[i + 2])
-            types[result_id] = f"f{width // 8}"
-        elif opcode in (
-            SpirvOpCode.OpSpecConstant,
-            SpirvOpCode.OpSpecConstantTrue,
-            SpirvOpCode.OpSpecConstantFalse,
-        ):
+            types[result_id] = {
+                "dtype": f"f{width // 8}",
+                "itemsize": width // 8,
+            }
+        elif opcode == SpirvOpCode.OpSpecConstant:
+            type_id = int(words[i + 1])
+            result_id = int(words[i + 2])
+            constants[result_id] = type_id
+            literal_words = words[i + 3 : i + word_count]
+            defaults[result_id] = literal_words.tobytes()
+        elif opcode == SpirvOpCode.OpSpecConstantTrue:
             type_id = int(words[i + 1])
             result_id = int(words[i + 2])
             constants[result_id] = type_id
+            defaults[result_id] = True
+        elif opcode == SpirvOpCode.OpSpecConstantFalse:
+            type_id = int(words[i + 1])
+            result_id = int(words[i + 2])
+            constants[result_id] = type_id
+            defaults[result_id] = False
         elif opcode == SpirvOpCode.OpDecorate:
             target_id = int(words[i + 1])
             decoration = int(words[i + 2])
@@ -92,15 +156,37 @@ def parse_spirv_specializations(
 
         i += word_count
 
-    result = {}
+    # a spec ID may appear multiple times in the same binary with different
+    # target IDs. We only need to keep one, so skip duplicates
+    unique_ids = set()
+    result = []
     for target_id, spec_id in ids.items():
+        if spec_id in unique_ids:
+            continue
+        unique_ids.add(spec_id)
         type_id = constants.get(target_id)
-        dtype_str = types.get(type_id, "unknown_type")
+        type_info = types.get(type_id, {"dtype": "unknown_type", "itemsize": 0})
         name = names.get(target_id, f"unnamed_spec_const_{spec_id}")
 
-        result[spec_id] = {
-            "name": name,
-            "dtype": dtype_str,
-        }
-
-    return result
+        dtype_str = type_info["dtype"]
+        raw_default = defaults.get(target_id)
+        default_value = None
+        if isinstance(raw_default, bytes):
+            try:
+                default_value = np.frombuffer(raw_default, dtype=dtype_str)[
+                    0
+                ].item()
+            except Exception:
+                default_value = None
+
+        result.append(
+            SpecializationConstantInfo(
+                spec_id=spec_id,
+                dtype=dtype_str,
+                name=name,
+                itemsize=type_info["itemsize"],
+                default_value=default_value,
+            )
+        )
+
+    return tuple(sorted(result, key=lambda x: x.spec_id))
diff --git a/dpctl/tests/input_files/specialization_constant_composite.spv b/dpctl/tests/input_files/specialization_constant_composite.spv
index 6b69617a7dc599495bc6a43862832b1ef77eab0d..c262f97ff6dfe48767c2ce52a31aea146a0d6b3e 100644
GIT binary patch
literal 3504
zcmZ{lYjaao6oz+5C|D3gZeFm4q7(!wy$Bct3pJG3T1rzuQRnnDIc=lKNlZ>FMNs&l
zzr(Nk$v@)HaYmiN@qJGAqM1(FGjG;)uf6tKYoBCyoY>_u8Mn{%xhJl>4!9jICl<QX
z(>_o4czV*)Ay1zj?)v|E#JPUA+m+|;2Bq1B;!LpI3YNo85;VeOHE1-ek?gxt`}EEE
zAbQY>DoGG_I%|zcvG(Z1jOeXa6i?0!cao}8ZdqmGQtWpF%G9}6h2Z`|d3L<BQK^pw
zcSkN%qa>`=N5|)zYjHKI&WG_zG#$0$sD7`OtWMu7&6dvx<FzDelxAup#hIm%VzCG}
zI#N#BYn7y|R3@8oC+Y4g7Ig;q=e+&O47wz6?hE#${#oTq{eJx=Bbgg%Pta-A+Uvpf
zwOTzX#Z&d>Qdlp=>&;4-)S7Wg*OcE&=^5MlmTM2G&ra2Lw`irVk_~hBLP3`=6xv)D
zd%vP=WA9d{Pzc-Ya09=+?&arxL9i0VQ9DecYOrlB(sN&ze<!Wf@WNVLi&y5tcG!rL
zsNHc>ipD)p3)|+u<*xL2N_Aq7Mn-%t<(aQzsobsi3*-5ESR8s|ao~ybb{Z$=4vL->
zm?6$VJ}Xd1<U^vpcH~Wa)^YEC?@LT<^q0M#ijo5hd0G^?n^(@wc;3z~OFkv+)!p&G
zzeTP{?v2k}^vcUy<c8$l^J6dD`Asj|`RkH<&wtz_A67p2O$aRHeWKKfh5UwS9zpAj
zI)*>Bu^dcC53>Txbkx|^86S)92g~Gh;)!MQRq>o*F{Z?8J>h56$FLqClP@)A5d-;7
zDsB2(&CH(~v5v~0oH-AgmrS3@2l=2ZyF`yk%Y5102V~{Vp@XMh$n4)M${uvc`$d^~
z@*0rL8$<`Nzaol1`s4CE;AP@qJEHh+isl9KcuSO6)SE9Ovr3QfL64od1ER#mj-2;0
zcOZ`wqVNFwJEHibM}AL~I}rDfyvDq~;Pn>wBgyRR&PmRF>}7Hx&NVN0?~rqglBowW
z_YFnazvyvUd*}f^8NuW}FP~OD$YorhpTyw}vNE!9|IdUmMJ2~sfqB7)Il;!9u<vt$
zWpU!J_<V0mazjAR%&_-@G9zF9k685UqRypXd}rY48{d?_RX6=jrgxwpd`sB-NayE8
zj|=R%qCV1xL9c_S_W044cLn0zkY8T(lz@)+H6+-6zMa@?{{!u(k7opI7VoV1-gv{3
z(Gl;QVDa7;k8Nwb^U`6%FDp7C&@cS*!fi?LqXIV8Z~l%Si0}5Ud!EI(AU!?a8snmL
z*jN_hl6YcZ>lgmedCbIz0`Uih4c(nty)007YR9Zz@v_wpJu^Geqvsoh{+j6Z9zC`B
zM;Y?JLvNvfsiWlliNM+Pdr~+mU?Z;^0=Z!;3fTTte(2#}6Nm?IceyE^m{Y=khj%z+
z`bVxMfjhZOHszVUtEvHaofXI@BixXV{@)Uq0pd|_W`9mVk8e&u&l|p-@|*gmtoYA+
z+s}%LTo#Cdp1u1-7laAHV*c*?a@X7Lh}!x0#9#06rMJZWLcpK5iO*Ai-hJ_$iQVq?
zrFdqMdvSkc{Oo+I*H_X}uZ-#iPraVVAHS~!_A^sfdwPU^P_UV!r>3XRy>meSo~h!O
z1p7WLNyiEWz6pz>l^(t-o}NddeB<~&EDL-es8LpQMc65jXI`*(wJIJP%kry<=S^T^
zw(Tz8h{tAggluo)Tk)J}b*g)rxq`R3!mlNeht1WZ`0oVnM11RO+w1Yo3HI(f;yH`l
zh6MbXw?x2Z{x)xG-e&W*E}7Wq?ah8K-eU4*iOC#35U^QHo5KxnV-D?oKNN2<?Aym2
z{vcqp7&f1eq$3XViI2_aWAWsOo%vi8{ZXLL)Q9;*#?N9}-G7pf9GFjd>TdJ&KQWv+
A5C8xG

literal 3432
zcma)-ZBHCk6vt=FQYyt_ZE1aIk=oLNR4fnDx)numffb~brJzmPVb~pD?LOe{fPgir
zFZ6TxR$uu={5U3RQse(Ob5Ge2ns}4{p2vI6x#ylc+j8P+o0m#?XS|c%4_>m)dMCWJ
zSZJ%GCmlWQ=ygYLIQqOR@&Dt3=e2vMy!_&dpIcbY&ifnN{zlM<{7MjQ`ITBRl)Wvs
z&&@9R;of#wi2R_@*r|kyby^uyqW8AL>eT#TBPx2Ttwoi1F*?0&RqA=K3;x69{K9Bs
zzfd0WSBA!lVHA|g!=p>JooX>GE(O)ia4xJ@!}4k=+M1iqE#z<cqopXU<mO95+4=RM
zY&HuwJd}^>JB28(Ql@IvMwINzW|f2Uldn^iL4BaSFWD3Ox2RtF_sh>I$(%@g{Kj^v
zzUxo!l*&=AI#aH#2jyIKw^j(EQmvZPG3ED4lH)kveCdGxw5qqKL<?oLY>=~OGCF)F
zQ|GwY+ZE*)d(xduCaBkgef&DSS6}$~{$^MW>p>J2{bPF(=Y3QDjkr^T%RALlb#pPO
z2bC}i>kV&4(Ku(n@PE$A9nYvw?9s@GPscp-bt#sUe!n!H>xaeRZY&Nwao&mJq`h;Z
z*9B&XJjgc%`iOj9wAqf_Y`^Se?iBx7QS!2obE3#eEotuy$6LNd$$dhH&Wr!*5qV8A
z`7C~LL@peWzmeRW|A~_=f7!{FzahCf|3Q=7rL6c(2rT3?qV$G^{FZ1Bg4Td;3IBG%
zYB1e39NZdk`+ASa1CrTi_K(FA&thc7Ydzy<*u|jVEM#h<zbs-PFU8WKKmB2uKmA}`
zlz)pT`Jg?L={fZvpOd9c^pdp9kDaqyR^AjkczT1({ti+0phNByWu~dCTQYA99l-vY
zDE{cL$n&g|iG%He;=e80BXrBoo1_2eh)E3I2QjcAcRQJS$$3Q-9$<f06o2%{eWKJ$
z%=7XZa{4<?Z*en{*_YgZ+8cMW)$ys5lX}zMBgvc{nX^51cJ$z`a?@{oQcW`N6}ew$
zr<UOu={OE=i#4JzBX!LQvjTeNfEgg)1A)bzBR`3y!+jOW%msSpVN6wUHvUIg+{10<
z=N|Yzz;hpb3;t4%^!FCmPwy`a?0u?S^q%hwd&X5`((^<IPoMGQ4z3Eso0cE@_~xMF
z-Si5!|Au(`9<hH<`};)u1#A}YJ@L))Zc0W+yaB=Dy)Pb{)noAnrNf3F^Z$Xs-PjrK
zOM<`UY$D#T?u;YixhHJovlye|xzD39K9mj{%VLa)CkD25;dkZZ{%;G!?-BNOcINUU
zfxgo_=JKwSt#|0hrTe5w&$k6Vb2`zar#FAALjEtfTj-y<QEHwP$jyC!CR`M-QP<}J
zwPBkQu>GU@(8IqW5D(tYGA*8%S>fNV6RIb<KWd#3IFpy+e}Vr8>TiwDIwzib(!#WK
z-2c453=ofgGv{9l=*<^B?{z_-N2aeS68`s`?I*=VUKEIdp1r*7`vQHmy}!A8$vfMU
zsO4W4KOvYecT3zA0sjXAK0msAtK!Lt-OlxucxG`@;QYw=Sxn}getj*_uax=)PrshY
zAHRnJ`^N<9J$Hn@N3faWPEF4}H_rk0_gr71JA!>5eCb$^1-=Q7MAw@5Ks<N8F3R_f
z??XXo6X;QkXi+%sU4_mj@}j;C@w^Fa*91Gurg&^NN67Xzw#1Xu`c!f<a|Lg6g<nOm
zzS4irQx(<(;@iE}#G^m-P2SzMbmXG8UIBmR?OOqx`P;nJoz3R0A(`0d?afBwEhcZ4
zn9SjhfX!ms9PT<Bb7=4TJMk97zJ1K$o`B6_*nI9wM;zu8ADhnu@zjXj`t!Yb?wS5D
WpUC)GOzZno>8OGEgs1N|PyYdNu{1IO

diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py
index 61b9fee843..564f40bed9 100644
--- a/dpctl/tests/test_sycl_program.py
+++ b/dpctl/tests/test_sycl_program.py
@@ -23,6 +23,7 @@
 
 import dpctl
 import dpctl.program as dpctl_prog
+from dpctl.program.utils import parse_spirv_specializations
 
 
 def get_spirv_abspath(fn):
@@ -341,3 +342,49 @@ def test_create_kernel_bundle_with_composite_spec_const():
 
     # 1.0 * 10 + 2.5 = 12.5
     assert np.all(y == 12.5)
+
+
+def test_spirv_specializations_parser():
+    spirv_file = get_spirv_abspath("specialization_constant_kernel.spv")
+    with open(spirv_file, "rb") as spv:
+        spv_bytes = spv.read()
+    spec_consts = parse_spirv_specializations(spv_bytes)
+    assert len(spec_consts) == 1
+    assert spec_consts[0].dtype == "u4"
+
+    spirv_file = get_spirv_abspath("specialization_constant_composite.spv")
+    with open(spirv_file, "rb") as spv:
+        spv_bytes = spv.read()
+
+    spec_consts = parse_spirv_specializations(spv_bytes)
+    assert len(spec_consts) == 3
+    spec_const0, spec_const1, spec_const2 = spec_consts
+    assert spec_const0.dtype == "u4"
+    assert spec_const0.itemsize == 4
+    assert spec_const0.name == "unnamed_spec_const_0"
+    assert spec_const0.default_value == 1
+
+    assert spec_const1.dtype == "f4"
+    assert spec_const1.itemsize == 4
+    assert spec_const1.name == "unnamed_spec_const_1"
+    assert spec_const1.default_value == 0
+
+    # compiler translates bool to char
+    assert spec_const2.dtype == "u1"
+    assert spec_const2.itemsize == 1
+    assert spec_const2.name == "unnamed_spec_const_2"
+    assert spec_const2.default_value == 0
+
+
+def test_spirv_specializations_parser_no_spec_consts():
+    spirv_file = get_spirv_abspath("multi_kernel.spv")
+    with open(spirv_file, "rb") as spv:
+        spv_bytes = spv.read()
+    spec_consts = parse_spirv_specializations(spv_bytes)
+    assert not spec_consts
+
+
+def test_spirv_specializations_parser_invalid_spirv():
+    invalid_spv = b"\x00\x01\x02\x03\x04\x05"
+    with pytest.raises(ValueError):
+        parse_spirv_specializations(invalid_spv)
diff --git a/setup.py b/setup.py
index 1b10322ef8..2c44bd4a11 100644
--- a/setup.py
+++ b/setup.py
@@ -27,6 +27,7 @@
         "dpctl",
         "dpctl.memory",
         "dpctl.program",
+        "dpctl.program.utils",
         "dpctl.utils",
     ],
     package_data={