diff --git a/.github/workflows/sycl-containers.yaml b/.github/workflows/sycl-containers.yaml
index 3fe84c663e1d8..2b3be35570127 100644
--- a/.github/workflows/sycl-containers.yaml
+++ b/.github/workflows/sycl-containers.yaml
@@ -51,6 +51,10 @@ jobs:
             file: ubuntu2404_build
             tag: latest
             build_args: ""
+          - name: Build Ubuntu 24.04 Docker image with CUDA 13.1
+            file: ubuntu2404_build_cuda131
+            tag: latest
+            build_args: ""
     steps:
       - name: Checkout
         uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
@@ -87,6 +91,9 @@ jobs:
           - name: Build + Intel Drivers Ubuntu 24.04 Docker image
             file: ubuntu2404_intel_drivers
             tag: latest
+          - name: Build + Intel Drivers Ubuntu 24.04 Docker image with CUDA 13.1
+            file: ubuntu2404_intel_drivers_cuda131
+            tag: latest
     steps:
       - name: Checkout
         uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
diff --git a/.github/workflows/sycl-linux-precommit.yml b/.github/workflows/sycl-linux-precommit.yml
index 0443dd4032268..7883f090eb359 100644
--- a/.github/workflows/sycl-linux-precommit.yml
+++ b/.github/workflows/sycl-linux-precommit.yml
@@ -159,6 +159,11 @@ jobs:
             runner: '["Linux", "cuda"]'
             image_options: -u 1001 --gpus all --cap-add SYS_ADMIN
             target_devices: cuda:gpu
+          - name: NVIDIA/CUDA 13.1
+            runner: '["Linux", "cuda13"]'
+            image: "ghcr.io/intel/llvm/ubuntu2404_intel_drivers_cuda131:latest"
+            image_options: -u 1001 --gpus all --cap-add SYS_ADMIN
+            target_devices: cuda:gpu
           - name: Intel / Arc A-Series Graphics
             runner: '["Linux", "arc"]'
             target_devices: level_zero:gpu;opencl:gpu;level_zero_v2:gpu         
@@ -247,6 +252,11 @@ jobs:
             runner: '["Linux", "cuda"]'
             image_options: -u 1001 --gpus all --cap-add SYS_ADMIN
             target_devices: cuda:gpu
+          - name: NVIDIA/CUDA 13.1
+            runner: '["Linux", "cuda13"]'
+            image: "ghcr.io/intel/llvm/ubuntu2404_intel_drivers_cuda131:latest"
+            image_options: -u 1001 --gpus all --cap-add SYS_ADMIN
+            target_devices: cuda:gpu
 
     uses: ./.github/workflows/sycl-linux-run-tests.yml
     with:
@@ -287,6 +297,10 @@ jobs:
           - name: CUDA system
             runner: '["Linux", "cuda"]'
             image_extra_opts: --gpus all
+          - name: CUDA system with CUDA 13.1
+            runner: '["Linux", "cuda13"]'
+            image: "ghcr.io/intel/llvm/ubuntu2404_intel_drivers_cuda131:latest"
+            image_extra_opts: --gpus all
     uses: ./.github/workflows/sycl-linux-run-tests.yml
     with:
       name: Perf tests on ${{ matrix.name }}
diff --git a/.github/workflows/ur-build-hw.yml b/.github/workflows/ur-build-hw.yml
index f8d029ae537a2..4723d4954402f 100644
--- a/.github/workflows/ur-build-hw.yml
+++ b/.github/workflows/ur-build-hw.yml
@@ -129,6 +129,7 @@ jobs:
         ${{ matrix.adapter.other_name != '' && format('-DUR_BUILD_ADAPTER_{0}=ON', matrix.adapter.other_name) || '' }}
         -DUR_STATIC_LOADER=${{matrix.adapter.static_Loader}}
         -DUR_STATIC_ADAPTER_${{matrix.adapter.name}}=${{matrix.adapter.static_adapter}}
+        ${{ matrix.adapter.name == 'CUDA' && '-DUR_CONFORMANCE_NVIDIA_ARCH="sm_75"' || '' }}
         -DUR_DPCXX=./dpcpp_compiler/bin/clang++
         -DUR_SYCL_LIBRARY_DIR=./dpcpp_compiler/lib
         -DCMAKE_INSTALL_PREFIX=./install
diff --git a/.github/workflows/ur-precommit.yml b/.github/workflows/ur-precommit.yml
index 403ac4d285826..b5f6e986ced99 100644
--- a/.github/workflows/ur-precommit.yml
+++ b/.github/workflows/ur-precommit.yml
@@ -56,39 +56,47 @@ jobs:
       # Extra native CPU jobs are here to force the loader to be used.
       # UR will not use the loader if there is only one target.
         include:
-          - name: L0
-            runner: UR_L0
-            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
-          - name: L0_V2
-            runner: UR_L0
-            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
-          - name: L0
-            runner: UR_L0
-            static: ON
-            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
-          - name: L0
-            runner: UR_L0
-            other_adapter: NATIVE_CPU
-            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
-          - name: L0_V2
-            runner: UR_L0_BMG
-            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
+          # - name: L0
+          #   runner: UR_L0
+          #   image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
+          # - name: L0_V2
+          #   runner: UR_L0
+          #   image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
+          # - name: L0
+          #   runner: UR_L0
+          #   static: ON
+          #   image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
+          # - name: L0
+          #   runner: UR_L0
+          #   other_adapter: NATIVE_CPU
+          #   image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
+          # - name: L0_V2
+          #   runner: UR_L0_BMG
+          #   image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
+
+          # 13.1
           - name: CUDA
-            runner: UR_CUDA
+            runner: cuda13
+            docker_image: "ghcr.io/intel/llvm/ubuntu2404_intel_drivers_cuda131:latest"
             image_options: -u 1001 --privileged --cap-add SYS_ADMIN --gpus all
-          - name: OPENCL
-            runner: UR_OPENCL
-            docker_image: "ghcr.io/intel/llvm/ubuntu2204_build:latest"
-            image_options: -u 1001 --device=/dev/dri --device=/dev/kfd --privileged --cap-add SYS_ADMIN
-          - name: OPENCL
-            runner: UR_OPENCL
-            other_adapter: NATIVE_CPU
-            docker_image: "ghcr.io/intel/llvm/ubuntu2204_build:latest"
-            image_options: -u 1001 --device=/dev/dri --device=/dev/kfd --privileged --cap-add SYS_ADMIN
-          - name: NATIVE_CPU
-            runner: UR_NATIVE_CPU
-            docker_image: "ghcr.io/intel/llvm/ubuntu2204_build:latest"
-            image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
+
+          # 12.6.3
+          # - name: CUDA
+          #   runner: UR_CUDA
+          #   image_options: -u 1001 --privileged --cap-add SYS_ADMIN --gpus all
+          # - name: OPENCL
+          #   runner: UR_OPENCL
+          #   docker_image: "ghcr.io/intel/llvm/ubuntu2204_build:latest"
+          #   image_options: -u 1001 --device=/dev/dri --device=/dev/kfd --privileged --cap-add SYS_ADMIN
+          # - name: OPENCL
+          #   runner: UR_OPENCL
+          #   other_adapter: NATIVE_CPU
+          #   docker_image: "ghcr.io/intel/llvm/ubuntu2204_build:latest"
+          #   image_options: -u 1001 --device=/dev/dri --device=/dev/kfd --privileged --cap-add SYS_ADMIN
+          # - name: NATIVE_CPU
+          #   runner: UR_NATIVE_CPU
+          #   docker_image: "ghcr.io/intel/llvm/ubuntu2204_build:latest"
+          #   image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
     uses: ./.github/workflows/ur-build-hw.yml
     with:
       adapter_name: ${{ matrix.name }}
diff --git a/devops/containers/ubuntu2404_build_cuda131.Dockerfile b/devops/containers/ubuntu2404_build_cuda131.Dockerfile
new file mode 100644
index 0000000000000..883781a24365b
--- /dev/null
+++ b/devops/containers/ubuntu2404_build_cuda131.Dockerfile
@@ -0,0 +1,71 @@
+FROM nvidia/cuda:13.1.0-devel-ubuntu24.04
+
+ENV DEBIAN_FRONTEND=noninteractive
+
+USER root
+
+# Configure LLVM nightly repo
+RUN apt-get update -qq && apt-get install --no-install-recommends -yqq curl ca-certificates
+RUN curl -sSL https://apt.llvm.org/llvm-snapshot.gpg.key -o /etc/apt/trusted.gpg.d/apt.llvm.org.asc
+RUN echo 'deb http://apt.llvm.org/noble/ llvm-toolchain-noble main' > /etc/apt/sources.list.d/llvm.list
+RUN echo 'deb http://apt.llvm.org/noble/ llvm-toolchain-noble-22 main' > /etc/apt/sources.list.d/llvm.list
+
+# Install SYCL prerequisites
+COPY scripts/install_build_tools.sh /install.sh
+RUN /install.sh
+
+COPY scripts/install_vulkan.sh /install_vulkan.sh
+RUN /install_vulkan.sh
+
+# libzstd-dev installed by default on Ubuntu 24.04 is not compiled with -fPIC flag.
+# This causes linking errors when building SYCL runtime.
+# Bug: https://github.com/intel/llvm/issues/15935
+# Workaround: build zstd from sources with -fPIC flag.
+COPY scripts/build_zstd.sh /build_zstd.sh
+RUN /build_zstd.sh
+
+SHELL ["/bin/bash", "-ec"]
+
+# Make the directory if it doesn't exist yet.
+# This location is recommended by the distribution maintainers.
+RUN mkdir --parents --mode=0755 /etc/apt/keyrings
+# Download the key, convert the signing-key to a full
+# keyring required by apt and store in the keyring directory
+RUN curl -sSL https://repo.radeon.com/rocm/rocm.gpg.key | \
+gpg --dearmor | tee /etc/apt/keyrings/rocm.gpg > /dev/null && \
+# Add rocm repo
+echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/rocm.gpg] https://repo.radeon.com/amdgpu/6.3/ubuntu noble main" \
+    | tee /etc/apt/sources.list.d/amdgpu.list && \
+echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/rocm.gpg] https://repo.radeon.com/rocm/apt/6.3 noble main" \
+    |  tee --append /etc/apt/sources.list.d/rocm.list && \
+echo -e 'Package: *\nPin: release o=repo.radeon.com\nPin-Priority: 600' \
+    | tee /etc/apt/preferences.d/rocm-pin-600 && \
+echo -e 'Package: *\nPin: release o=repo.radeon.com\nPin-Priority: 600' \
+    | tee /etc/apt/preferences.d/rocm-pin-600
+# Install the ROCM kernel driver
+RUN apt update && apt install -yqq rocm-dev && \
+  apt-get clean && \
+  rm -rf /var/lib/apt/lists/*
+
+# Fix Vulkan install inside container
+# https://stackoverflow.com/questions/74965945/vulkan-is-unable-to-detect-nvidia-gpu-from-within-a-docker-container-when-using
+RUN apt-get update && \
+    apt-get install -y libegl1 && \
+    apt-get install -y --no-install-recommends --download-only libnvidia-gl-565 && \
+    dpkg-deb --extract /var/cache/apt/archives/libnvidia-gl-565_*.deb extracted && \
+    cp -R ./extracted/usr/* /usr/ && \
+    rm -rf /var/lib/apt/lists/* /var/cache/apt/archives/*.deb ./extracted
+
+COPY scripts/create-sycl-user.sh /user-setup.sh
+RUN /user-setup.sh
+
+COPY scripts/docker_entrypoint.sh /docker_entrypoint.sh
+
+COPY actions/cached_checkout /actions/cached_checkout
+COPY actions/cleanup /actions/cleanup
+COPY scripts/install_drivers.sh /opt/install_drivers.sh
+
+USER sycl
+
+ENTRYPOINT ["/docker_entrypoint.sh"]
+
diff --git a/devops/containers/ubuntu2404_intel_drivers_cuda131.Dockerfile b/devops/containers/ubuntu2404_intel_drivers_cuda131.Dockerfile
new file mode 100644
index 0000000000000..c9771bd305364
--- /dev/null
+++ b/devops/containers/ubuntu2404_intel_drivers_cuda131.Dockerfile
@@ -0,0 +1,25 @@
+ARG base_tag=latest
+ARG base_image=ghcr.io/intel/llvm/ubuntu2404_build_cuda131
+
+FROM $base_image:$base_tag
+
+ENV DEBIAN_FRONTEND=noninteractive
+
+USER root
+
+RUN apt update && apt install -yqq wget
+
+COPY scripts/install_drivers.sh /
+COPY dependencies.json /
+
+RUN mkdir /runtimes
+ENV INSTALL_LOCATION=/runtimes
+RUN --mount=type=secret,id=github_token \
+    GITHUB_TOKEN=$(cat /run/secrets/github_token) /install_drivers.sh dependencies.json --all
+
+COPY scripts/drivers_entrypoint.sh /drivers_entrypoint.sh
+
+USER sycl
+
+ENTRYPOINT ["/bin/bash", "/drivers_entrypoint.sh"]
+
diff --git a/sycl/test-e2e/WorkGroupMemory/basic_usage.cpp b/sycl/test-e2e/WorkGroupMemory/basic_usage_common.hpp
similarity index 95%
rename from sycl/test-e2e/WorkGroupMemory/basic_usage.cpp
rename to sycl/test-e2e/WorkGroupMemory/basic_usage_common.hpp
index c63f16733b289..1d63e58c0c4dc 100644
--- a/sycl/test-e2e/WorkGroupMemory/basic_usage.cpp
+++ b/sycl/test-e2e/WorkGroupMemory/basic_usage_common.hpp
@@ -1,23 +1,19 @@
-// UNSUPPORTED: hip
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17339
-// RUN: %{build} -o %t.out
-// RUN: %{run} %t.out
-// XFAIL: spirv-backend
-// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18230
+// Shared implementation for the work-group memory basic usage tests.
+
+#pragma once
 
 #include <sycl/builtins.hpp>
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/experimental/work_group_memory.hpp>
 #include <sycl/group_barrier.hpp>
 #include <sycl/half_type.hpp>
+#include <sycl/queue.hpp>
 
 #include <cassert>
 #include <cstring>
 
 namespace syclexp = sycl::ext::oneapi::experimental;
 
-sycl::queue q;
-
 // This test performs a swap of two scalars/arrays inside a kernel using a
 // work_group_memory object as a temporary buffer. The test is done for scalar
 // types and bounded arrays. After the kernel finishes, it is verified on the
@@ -407,19 +403,3 @@ template <typename T> void test_ptr() {
   }
   swap_array_2d(arr1, arr2, 8);
 }
-
-int main() {
-  test<int>();
-  test<char>();
-  test<uint16_t>();
-  if (q.get_device().has(sycl::aspect::fp16))
-    test<sycl::half>();
-  test_ptr<float *>();
-  test_ptr<int *>();
-  test_ptr<char *>();
-  test_ptr<uint16_t *>();
-  if (q.get_device().has(sycl::aspect::fp16))
-    test_ptr<sycl::half *>();
-  test_ptr<float *>();
-  return 0;
-}
diff --git a/sycl/test-e2e/WorkGroupMemory/basic_usage_test.cpp b/sycl/test-e2e/WorkGroupMemory/basic_usage_test.cpp
new file mode 100644
index 0000000000000..5589cfcc7770d
--- /dev/null
+++ b/sycl/test-e2e/WorkGroupMemory/basic_usage_test.cpp
@@ -0,0 +1,26 @@
+// Non-pointer types version of the basic usage test.
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17339
+
+// XFAIL: spirv-backend
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18230
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/queue.hpp>
+
+sycl::queue q;
+
+#include "./basic_usage_common.hpp"
+
+int main() {
+  test<float>();
+  test<int>();
+  test<char>();
+  test<uint16_t>();
+  if (q.get_device().has(sycl::aspect::fp16))
+    test<sycl::half>();
+  return 0;
+}
diff --git a/sycl/test-e2e/WorkGroupMemory/basic_usage_test_ptr.cpp b/sycl/test-e2e/WorkGroupMemory/basic_usage_test_ptr.cpp
new file mode 100644
index 0000000000000..b675e1182165c
--- /dev/null
+++ b/sycl/test-e2e/WorkGroupMemory/basic_usage_test_ptr.cpp
@@ -0,0 +1,26 @@
+// Pointer-types version of the basic usage test.
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17339
+
+// XFAIL: spirv-backend
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18230
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/queue.hpp>
+
+sycl::queue q;
+
+#include "./basic_usage_common.hpp"
+
+int main() {
+  test_ptr<float *>();
+  test_ptr<int *>();
+  test_ptr<char *>();
+  test_ptr<uint16_t *>();
+  if (q.get_device().has(sycl::aspect::fp16))
+    test_ptr<sycl::half *>();
+  return 0;
+}
diff --git a/sycl/test-e2e/bindless_images/read_sampled_1d.cpp b/sycl/test-e2e/bindless_images/read_sampled_1d.cpp
new file mode 100644
index 0000000000000..dadb9ae81efcf
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/read_sampled_1d.cpp
@@ -0,0 +1,31 @@
+// 1D version of sampled image read test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Returning non-FP values from sampling fails on HIP.
+
+// UNSUPPORTED: linux && arch-intel_gpu_bmg_g21 && level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20223
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include "./read_sampled_common.hpp"
+
+int main() {
+  const unsigned int seed = 0;
+  const float offset = 20.0;
+
+  std::cout << "Running 1D Sampled Image Tests!\n";
+  bool result1D = runAll1D(offset, seed);
+
+  if (result1D) {
+    std::cout << "All tests passed!\n";
+  } else {
+    std::cerr << "An error has occurred!\n";
+    return 1;
+  }
+
+  return 0;
+}
diff --git a/sycl/test-e2e/bindless_images/read_sampled_2d.cpp b/sycl/test-e2e/bindless_images/read_sampled_2d.cpp
new file mode 100644
index 0000000000000..7e14ab05a2130
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/read_sampled_2d.cpp
@@ -0,0 +1,34 @@
+// 2D version of sampled image read test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Returning non-FP values from sampling fails on HIP.
+
+// UNSUPPORTED: linux && arch-intel_gpu_bmg_g21 && level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20223
+
+// XFAIL: cuda && cuda-major-ge-13
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/21807
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include "./read_sampled_common.hpp"
+
+int main() {
+  const unsigned int seed = 0;
+  const float offset = 20.0;
+
+  std::cout << "Running 2D Sampled Image Tests!\n";
+  bool result2D = runAll2D(offset, seed);
+
+  if (result2D) {
+    std::cout << "All tests passed!\n";
+  } else {
+    std::cerr << "An error has occurred!\n";
+    return 1;
+  }
+
+  return 0;
+}
diff --git a/sycl/test-e2e/bindless_images/read_sampled.cpp b/sycl/test-e2e/bindless_images/read_sampled_common.hpp
similarity index 97%
rename from sycl/test-e2e/bindless_images/read_sampled.cpp
rename to sycl/test-e2e/bindless_images/read_sampled_common.hpp
index c30f0a05b2c65..3166a2c9b275f 100644
--- a/sycl/test-e2e/bindless_images/read_sampled.cpp
+++ b/sycl/test-e2e/bindless_images/read_sampled_common.hpp
@@ -1,13 +1,6 @@
-// REQUIRES: aspect-ext_oneapi_bindless_images
+// Shared implementation for the sampled image read tests.
 
-// UNSUPPORTED: hip
-// UNSUPPORTED-INTENDED: Returning non-FP values from sampling fails on HIP.
-
-// UNSUPPORTED: linux && arch-intel_gpu_bmg_g21 && level_zero_v2_adapter
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20223
-
-// RUN: %{build} -o %t.out
-// RUN: %{run} %t.out
+#pragma once
 
 // Print test names and pass status
 // #define VERBOSE_LV1
@@ -26,6 +19,8 @@
 #include <iostream>
 #include <sycl/accessor_image.hpp>
 #include <sycl/detail/core.hpp>
+#include <sycl/image.hpp>
+#include <type_traits>
 
 #include <sycl/ext/oneapi/bindless_images.hpp>
 
@@ -488,7 +483,7 @@ bool runTests(sycl::range<2> dims, sycl::range<2> localSize, float offset,
               int seed, sycl::coordinate_normalization_mode normMode) {
 
   // addressing_mode::none currently removed due to
-  // inconsistent behavour when switching between
+  // inconsistent behaviour when switching between
   // normalized and unnormalized coords.
   sycl::addressing_mode addrModes[4] = {
       sycl::addressing_mode::repeat, sycl::addressing_mode::mirrored_repeat,
@@ -718,22 +713,10 @@ bool runAll(sycl::range<NDims> dims, sycl::range<NDims> localSize, float offset,
   return offsetPassed && noOffsetPassed;
 }
 
-int main() {
-
-  const unsigned int seed = 0;
-  const float offset = 20.0;
-
-  std::cout << "Running 1D Sampled Image Tests!\n";
-  bool result1D = runAll<1>({128}, {32}, offset, seed);
-  std::cout << "Running 2D Sampled Image Tests!\n";
-  bool result2D = runAll<2>({16, 16}, {8, 8}, offset, seed);
-
-  if (result1D && result2D) {
-    std::cout << "All tests passed!\n";
-  } else {
-    std::cerr << "An error has occurred!\n";
-    return 1;
-  }
+inline bool runAll1D(float offset, int seed) {
+  return runAll<1>(sycl::range<1>{128}, sycl::range<1>{32}, offset, seed);
+}
 
-  return 0;
+inline bool runAll2D(float offset, int seed) {
+  return runAll<2>(sycl::range<2>{16, 16}, sycl::range<2>{8, 8}, offset, seed);
 }
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_channels.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_channels.cpp
new file mode 100644
index 0000000000000..1f5a63e5abb90
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_channels.cpp
@@ -0,0 +1,39 @@
+// Unsampled version of the Vulkan/SYCL 1D image read interop test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
+// REQUIRES: vulkan
+
+// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+
+// clang-format off
+// RUN: %{run} %t.out --type float --channels 1 32
+// RUN: %{run} %t.out --type float --channels 2 32
+// RUN: %{run} %t.out --type float --channels 4 32
+// RUN: %{run} %t.out --type half --channels 1 32
+// RUN: %{run} %t.out --type half --channels 2 32
+// RUN: %{run} %t.out --type half --channels 4 32
+// RUN: %{run} %t.out --type int32 --channels 1 32
+// RUN: %{run} %t.out --type int32 --channels 2 32
+// RUN: %{run} %t.out --type int32 --channels 4 32
+// RUN: %{run} %t.out --type uint32 --channels 1 32
+// RUN: %{run} %t.out --type uint32 --channels 2 32
+// RUN: %{run} %t.out --type uint32 --channels 4 32
+// RUN: %{run} %t.out --type int16 --channels 1 32
+// RUN: %{run} %t.out --type int16 --channels 2 32
+// RUN: %{run} %t.out --type int16 --channels 4 32
+// RUN: %{run} %t.out --type uint16 --channels 1 32
+// RUN: %{run} %t.out --type uint16 --channels 2 32
+// RUN: %{run} %t.out --type uint16 --channels 4 32
+// RUN: %{run} %t.out --type uint8 --channels 1 32
+// RUN: %{run} %t.out --type uint8 --channels 2 32
+// RUN: %{run} %t.out --type uint8 --channels 4 32
+// RUN: %{run} %t.out --type int8 --channels 1 32
+// RUN: %{run} %t.out --type int8 --channels 2 32
+// RUN: %{run} %t.out --type int8 --channels 4 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 1 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 2 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 4 32
+// clang-format on
+
+#include "./vulkan_sycl_image_interop_read_1d_common.hpp"
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_common.hpp
similarity index 65%
rename from sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp
rename to sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_common.hpp
index bbbc5b7c1e97b..2d854067305a2 100644
--- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_common.hpp
@@ -1,6 +1,3 @@
-// REQUIRES: aspect-ext_oneapi_bindless_images
-// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
-// REQUIRES: vulkan
 
 // XFAIL: windows && gpu-intel-dg2
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/21985
@@ -8,135 +5,64 @@
 // UNSUPPORTED: windows && arch-intel_gpu_bmg_g21
 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22084
 
-// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+// Shared implementation for the Vulkan/SYCL 1D image read interop tests.
 
 /*
-    Run ALL the vulkan formats through the gauntlet. sampled and unsampled.
-    This entire test takes less than 30 seconds on a slow machine.  MUCH faster
-   (and more complete coveraage) than SFINAE based approach.
+  Run ALL the vulkan formats through the gauntlet. sampled and unsampled.
+  This entire test takes less than 30 seconds on a slow machine.  MUCH faster
+  (and more complete coveraage) than SFINAE based approach.
 
-    IF a particular variant is having problems on some platform, please do NOT
-   just disable the whole test, instead use   RUN~IF (SOMETHING) yadda-yadda
-    to enable/disable that variant.
-
-    For semaphore testing, we run just a sampling. Note, that on Linux if there
-   is a failure in the first section, then likely ALL semaphore tests afterwards
-   will fail. This is being tracked as a separate issue.
+  IF a particular variant is having problems on some platform, please do NOT
+  just disable the whole test, instead use   RUN~IF (SOMETHING) yadda-yadda
+  to enable/disable that variant.
 
+  For semaphore testing, we run just a sampling. Note, that on Linux if there
+  is a failure in the first section, then likely ALL semaphore tests afterwards
+  will fail. This is being tracked as a separate issue.
 */
-// clang-format off
-// RUN: %{run} %t.out --type float --channels 1 32
-// RUN: %{run} %t.out --type float --channels 2 32
-// RUN: %{run} %t.out --type float --channels 4 32
-// RUN: %{run} %t.out --type half --channels 1 32
-// RUN: %{run} %t.out --type half --channels 2 32
-// RUN: %{run} %t.out --type half --channels 4 32
-// RUN: %{run} %t.out --type int32 --channels 1 32
-// RUN: %{run} %t.out --type int32 --channels 2 32
-// RUN: %{run} %t.out --type int32 --channels 4 32
-// RUN: %{run} %t.out --type uint32 --channels 1 32
-// RUN: %{run} %t.out --type uint32 --channels 2 32
-// RUN: %{run} %t.out --type uint32 --channels 4 32
-// RUN: %{run} %t.out --type int16 --channels 1 32
-// RUN: %{run} %t.out --type int16 --channels 2 32
-// RUN: %{run} %t.out --type int16 --channels 4 32
-// RUN: %{run} %t.out --type uint16 --channels 1 32
-// RUN: %{run} %t.out --type uint16 --channels 2 32
-// RUN: %{run} %t.out --type uint16 --channels 4 32
-// RUN: %{run} %t.out --type uint8 --channels 1 32
-// RUN: %{run} %t.out --type uint8 --channels 2 32
-// RUN: %{run} %t.out --type uint8 --channels 4 32
-// RUN: %{run} %t.out --type int8 --channels 1 32
-// RUN: %{run} %t.out --type int8 --channels 2 32
-// RUN: %{run} %t.out --type int8 --channels 4 32
-// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 1 32
-// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 2 32
-// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 4 32
-// RUN: %{run} %t.out --type float --channels 1 --sampled 32
-// RUN: %{run} %t.out --type float --channels 2 --sampled 32
-// RUN: %{run} %t.out --type float --channels 4 --sampled 32
-// RUN: %{run} %t.out --type half --channels 1 --sampled 32
-// RUN: %{run} %t.out --type half --channels 2 --sampled 32
-// RUN: %{run} %t.out --type half --channels 4 --sampled 32
-// RUN: %{run} %t.out --type int32 --channels 1 --sampled 32
-// RUN: %{run} %t.out --type int32 --channels 2 --sampled 32
-// RUN: %{run} %t.out --type int32 --channels 4 --sampled 32
-// RUN: %{run} %t.out --type uint32 --channels 1 --sampled 32
-// RUN: %{run} %t.out --type uint32 --channels 2 --sampled 32
-// RUN: %{run} %t.out --type uint32 --channels 4 --sampled 32
-// RUN: %{run} %t.out --type int16 --channels 1 --sampled 32
-// RUN: %{run} %t.out --type int16 --channels 2 --sampled 32
-// RUN: %{run} %t.out --type int16 --channels 4 --sampled 32
-// RUN: %{run} %t.out --type uint16 --channels 1 --sampled 32
-// RUN: %{run} %t.out --type uint16 --channels 2 --sampled 32
-// RUN: %{run} %t.out --type uint16 --channels 4 --sampled 32
-// RUN: %{run} %t.out --type uint8 --channels 1 --sampled 32
-// RUN: %{run} %t.out --type uint8 --channels 2 --sampled 32
-// RUN: %{run} %t.out --type uint8 --channels 4 --sampled 32
-// RUN: %{run} %t.out --type int8 --channels 1 --sampled 32
-// RUN: %{run} %t.out --type int8 --channels 2 --sampled 32
-// RUN: %{run} %t.out --type int8 --channels 4 --sampled 32
-// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 1 --sampled 32
-// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 2 --sampled 32
-// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 4 --sampled 32
-
-// RUN: %{run} %t.out --type float --channels 1 32 --semaphores
-// RUN: %{run} %t.out --type float --channels 2 32 --semaphores
-// RUN: %{run} %t.out --type float --channels 4 32 --semaphores
-// RUN: %{run} %t.out --type half --channels 1 32 --semaphores
-// RUN: %{run} %t.out --type int32 --channels 2 32 --semaphores
-// RUN: %{run} %t.out --type uint32 --channels 4 32 --semaphores
-// RUN: %{run} %t.out --type int16 --channels 1 32 --semaphores
-// RUN: %{run} %t.out --type uint16 --channels 2 32 --semaphores
-// RUN: %{run} %t.out --type uint8 --channels 4 32 --semaphores
-// RUN: %{run} %t.out --type int8 --channels 1 32 --semaphores
-// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 2 32 --semaphores
-// RUN: %{run} %t.out --type float --channels 4 --sampled 32 --semaphores
-// RUN: %{run} %t.out --type int16 --channels 4 --sampled 32 --semaphores
-// RUN: %{run} %t.out --type int8 --channels 4 --sampled 32 --semaphores
-// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 4 --sampled 32 --semaphores
 
 /*
-
-The block above tests these formats, sampled and unsampled, with and without
-semaphores
-
-VK_FORMAT_R32_SFLOAT
-VK_FORMAT_R32G32_SFLOAT
-VK_FORMAT_R32G32B32A32_SFLOAT
-VK_FORMAT_R16_SFLOAT
-VK_FORMAT_R16G16_SFLOAT
-VK_FORMAT_R16G16B16A16_SFLOAT
-VK_FORMAT_R32_SINT
-VK_FORMAT_R32G32_SINT
-VK_FORMAT_R32G32B32A32_SINT
-VK_FORMAT_R32_UINT
-VK_FORMAT_R32G32_UINT
-VK_FORMAT_R32G32B32A32_UINT
-VK_FORMAT_R16_SINT
-VK_FORMAT_R16G16_SINT
-VK_FORMAT_R16G16B16A16_SINT
-VK_FORMAT_R16_UINT
-VK_FORMAT_R16G16_UINT
-VK_FORMAT_R16G16B16A16_UINT
-VK_FORMAT_R8_UINT
-VK_FORMAT_R8G8_UINT
-VK_FORMAT_R8G8B8A8_UINT
-VK_FORMAT_R8_SINT
-VK_FORMAT_R8G8_SINT
-VK_FORMAT_R8G8B8A8_SINT
-VK_FORMAT_R8_UNORM
-VK_FORMAT_R8G8_UNORM
-VK_FORMAT_R8G8B8A8_UNORM
-
+  The block above tests these formats, sampled and unsampled, with and without
+  semaphores
+
+  VK_FORMAT_R32_SFLOAT
+  VK_FORMAT_R32G32_SFLOAT
+  VK_FORMAT_R32G32B32A32_SFLOAT
+  VK_FORMAT_R16_SFLOAT
+  VK_FORMAT_R16G16_SFLOAT
+  VK_FORMAT_R16G16B16A16_SFLOAT
+  VK_FORMAT_R32_SINT
+  VK_FORMAT_R32G32_SINT
+  VK_FORMAT_R32G32B32A32_SINT
+  VK_FORMAT_R32_UINT
+  VK_FORMAT_R32G32_UINT
+  VK_FORMAT_R32G32B32A32_UINT
+  VK_FORMAT_R16_SINT
+  VK_FORMAT_R16G16_SINT
+  VK_FORMAT_R16G16B16A16_SINT
+  VK_FORMAT_R16_UINT
+  VK_FORMAT_R16G16_UINT
+  VK_FORMAT_R16G16B16A16_UINT
+  VK_FORMAT_R8_UINT
+  VK_FORMAT_R8G8_UINT
+  VK_FORMAT_R8G8B8A8_UINT
+  VK_FORMAT_R8_SINT
+  VK_FORMAT_R8G8_SINT
+  VK_FORMAT_R8G8B8A8_SINT
+  VK_FORMAT_R8_UNORM
+  VK_FORMAT_R8G8_UNORM
+  VK_FORMAT_R8G8B8A8_UNORM
 */
 
+// clang-format off
 /*
   Vulkan/SYCL 1D Image Read Test (Sampled + Unsampled)
 
-  clang++ -fsycl -o vsr_1d_test.bin vulkan_sycl_image_interop_read_1d.cpp -lvulkan -I$VULKAN_SDK/include -L$VULKAN_SDK/lib
+  clang++ -fsycl -o vsr_1d_test.bin vulkan_sycl_image_interop_read_1d.cpp
+  -lvulkan -I$VULKAN_SDK/include -L$VULKAN_SDK/lib
 
-  clang++ -fsycl -o vsr_1d_test.exe vulkan_sycl_image_interop_read_1d.cpp -Wno-ignored-attributes  -lvulkan-1 -I$VULKAN_SDK/Include -L$VULKAN_SDK/Lib
+  clang++ -fsycl -o vsr_1d_test.exe vulkan_sycl_image_interop_read_1d.cpp
+  -Wno-ignored-attributes  -lvulkan-1 -I$VULKAN_SDK/Include -L$VULKAN_SDK/Lib
 
   USAGE:
     ./vsr_1d_test.bin [FLAGS] [Wx]
@@ -146,8 +72,9 @@ VK_FORMAT_R8G8B8A8_UNORM
     --semaphores   Use Vulkan Semaphores for SYCL Interop Sync
     --linear       Use LINEAR tiling for the Vulkan Image (default is OPTIMAL)
     --channels X   Set number of channels (1, 2, or 4). Default is 4 (RGBA)
-    --type XXX     Set data type (float, half, uint32, int32, uint16, int16, uint8, int8, unorm8). Default is float 
-    Wx             Set custom Width (e.g. 64x)
+    --type XXX     Set data type (float, half, uint32, int32, uint16, int16,
+  uint8, int8, unorm8). Default is float Wx             Set custom Width (e.g.
+  64x)
 
   EXAMPLES:
     ./vsr_1d_test.bin
@@ -158,8 +85,12 @@ VK_FORMAT_R8G8B8A8_UNORM
 */
 // clang-format on
 
+#pragma once
+
 #include "vulkan_setup.hpp"
 
+#include <chrono>
+#include <cstdlib>
 #include <optional>
 #include <string>
 #include <sycl/builtins.hpp>
@@ -167,6 +98,65 @@ VK_FORMAT_R8G8B8A8_UNORM
 #include <sycl/ext/oneapi/bindless_images.hpp>
 #include <sycl/ext/oneapi/bindless_images_interop.hpp>
 #include <sycl/image.hpp>
+#include <sycl/platform.hpp>
+
+inline std::string formatVulkanVersion(uint32_t version) {
+  return std::to_string(VK_API_VERSION_MAJOR(version)) + "." +
+         std::to_string(VK_API_VERSION_MINOR(version)) + "." +
+         std::to_string(VK_API_VERSION_PATCH(version));
+}
+
+inline void printVulkanDependencyVersions(const VulkanContext &vkCtx) {
+  uint32_t loaderVersion = VK_API_VERSION_1_0;
+  auto enumerateInstanceVersion =
+      reinterpret_cast<PFN_vkEnumerateInstanceVersion>(
+          vkGetInstanceProcAddr(VK_NULL_HANDLE, "vkEnumerateInstanceVersion"));
+  if (enumerateInstanceVersion != nullptr) {
+    VkResult result = enumerateInstanceVersion(&loaderVersion);
+    if (result != VK_SUCCESS)
+      loaderVersion = VK_API_VERSION_1_0;
+  }
+
+  VkPhysicalDeviceProperties props{};
+  vkGetPhysicalDeviceProperties(vkCtx.physicalDevice, &props);
+
+  std::cout << "[DEPS] Vulkan loader API version: "
+            << formatVulkanVersion(loaderVersion) << std::endl;
+  std::cout << "[DEPS] Vulkan device: " << props.deviceName << std::endl;
+  std::cout << "[DEPS] Vulkan device API version: "
+            << formatVulkanVersion(props.apiVersion) << std::endl;
+  std::cout << "[DEPS] Vulkan driver version (raw): " << props.driverVersion
+            << std::endl;
+  std::cout << "[DEPS] Vulkan vendor/device ID: 0x" << std::hex
+            << props.vendorID << "/0x" << props.deviceID << std::dec
+            << std::endl;
+}
+
+inline void printSyclDependencyVersions(const sycl::queue &q) {
+  const sycl::device dev = q.get_device();
+  const sycl::platform platform = dev.get_platform();
+
+#ifdef SYCL_LANGUAGE_VERSION
+  std::cout << "[DEPS] SYCL language version: " << SYCL_LANGUAGE_VERSION
+            << std::endl;
+#endif
+#ifdef __SYCL_COMPILER_VERSION
+  std::cout << "[DEPS] SYCL compiler version macro: "
+            << __SYCL_COMPILER_VERSION << std::endl;
+#endif
+
+  std::cout << "[DEPS] SYCL platform: "
+            << platform.get_info<sycl::info::platform::name>() << " | vendor: "
+            << platform.get_info<sycl::info::platform::vendor>()
+            << " | version: "
+            << platform.get_info<sycl::info::platform::version>() << std::endl;
+
+  std::cout << "[DEPS] SYCL device: "
+            << dev.get_info<sycl::info::device::name>() << " | vendor: "
+            << dev.get_info<sycl::info::device::vendor>() << " | version: "
+            << dev.get_info<sycl::info::device::version>() << " | driver: "
+            << dev.get_info<sycl::info::device::driver_version>() << std::endl;
+}
 
 // ---------------------------------------------------------
 // SYCL TYPE MAPPING HELPERS
@@ -232,6 +222,25 @@ int runTest(
     bool useSampled, VkFormat fmtOverride = VK_FORMAT_UNDEFINED,
     std::optional<sycl::image_channel_type> syclOverride = std::nullopt) {
 
+  const bool profileEnabled = std::getenv("VULKAN_SYCL_PROFILE") != nullptr;
+  using Clock = std::chrono::steady_clock;
+  auto profileStart = Clock::now();
+  auto profileLast = profileStart;
+  auto logProfile = [&](const char *label) {
+    if (!profileEnabled)
+      return;
+    auto now = Clock::now();
+    auto stepMs =
+        std::chrono::duration_cast<std::chrono::milliseconds>(now - profileLast)
+            .count();
+    auto totalMs =
+        std::chrono::duration_cast<std::chrono::milliseconds>(now - profileStart)
+            .count();
+    std::cout << "[PROFILE] " << label << " step_ms=" << stepMs
+              << " total_ms=" << totalMs << std::endl;
+    profileLast = now;
+  };
+
   VkImageTiling tiling =
       useLinear ? VK_IMAGE_TILING_LINEAR : VK_IMAGE_TILING_OPTIMAL;
   VkFormat vkFormat = (fmtOverride != VK_FORMAT_UNDEFINED)
@@ -242,25 +251,32 @@ int runTest(
 
   // Setup Vulkan
   VulkanContext vkCtx = createVulkanContext();
+  printVulkanDependencyVersions(vkCtx);
+  logProfile("createVulkanContext");
   VkExtent3D extent = {(uint32_t)width, 1, 1};
   ImageResources imgRes =
       createExportableImage(vkCtx, extent, vkFormat, VK_IMAGE_TYPE_1D, tiling);
+  logProfile("createExportableImage");
 
   // Semaphores
   VkSemaphore vkSem = VK_NULL_HANDLE;
   if (useSemaphores)
     vkSem = createExportableSemaphore(vkCtx);
+  logProfile("createExportableSemaphore");
 
   // Upload test data
   if (!uploadAndVerify<T>(vkCtx, imgRes, vkSem, channels)) {
     std::cerr << "Vulkan Upload Failed!" << std::endl;
     return 1;
   }
+  logProfile("uploadAndVerify");
 
   // SYCL Import and Verification
   namespace syclexp = sycl::ext::oneapi::experimental;
   try {
     sycl::queue q;
+    printSyclDependencyVersions(q);
+    logProfile("create_sycl_queue");
 
     // Import Memory (Platform Specific)
 #ifdef _WIN32
@@ -277,6 +293,7 @@ int runTest(
 
     syclexp::external_mem extMem = syclexp::import_external_memory(
         extMemDesc, q.get_device(), q.get_context());
+    logProfile("import_external_memory");
 
     // Import Semaphore (Platform Specific)
     syclexp::external_semaphore extSem;
@@ -294,6 +311,7 @@ int runTest(
       extSem = syclexp::import_external_semaphore(extSemDesc, q.get_device(),
                                                   q.get_context());
     }
+    logProfile("import_external_semaphore");
 
     // Create Image Descriptor
     sycl::image_channel_type syclType = syclOverride.has_value()
@@ -306,6 +324,7 @@ int runTest(
     // Map external memory
     syclexp::image_mem_handle devHandle = syclexp::map_external_image_memory(
         extMem, imgDesc, q.get_device(), q.get_context());
+    logProfile("map_external_image_memory");
 
     // Branch: Sampled vs Unsampled
     syclexp::sampled_image_handle sampledHandle;
@@ -324,6 +343,7 @@ int runTest(
       unsampledHandle = syclexp::create_image(devHandle, imgDesc,
                                               q.get_device(), q.get_context());
     }
+    logProfile("create_image_handle");
 
     // Output Buffer
     size_t totalValues = width * channels;
@@ -336,6 +356,7 @@ int runTest(
         h.ext_oneapi_wait_external_semaphore(extSem);
       });
     }
+    logProfile("submit_external_semaphore_wait");
 
     // Kernel: Read image data
     q.submit([&](sycl::handler &h) {
@@ -410,6 +431,7 @@ int runTest(
          }
        });
      }).wait();
+    logProfile("submit_and_wait_kernel");
 
     std::cout << "SYCL Kernel Executed." << std::endl;
 
@@ -430,6 +452,7 @@ int runTest(
         errorCount++;
       }
     }
+    logProfile("host_verify");
 
     if (passed) {
       std::cout << "SUCCESS! All " << totalValues << " values match."
@@ -455,8 +478,10 @@ int runTest(
                                           q.get_context());
       vkDestroySemaphore(vkCtx.device, vkSem, nullptr);
     }
+    logProfile("release_external_resources");
 
     cleanupVulkan(vkCtx, imgRes);
+    logProfile("cleanupVulkan");
     return passed ? 0 : 1;
 
   } catch (std::exception &e) {
@@ -470,6 +495,22 @@ int runTest(
 // MAIN
 // ---------------------------------------------------------
 int main(int argc, char **argv) {
+  // Enable profiling by default
+  ::setenv("VULKAN_SYCL_PROFILE", "1", 0);
+  
+  const bool profileEnabled = std::getenv("VULKAN_SYCL_PROFILE") != nullptr;
+  using Clock = std::chrono::steady_clock;
+  auto processStart = Clock::now();
+  auto finish = [&](int rc) {
+    if (profileEnabled) {
+      auto totalMs = std::chrono::duration_cast<std::chrono::milliseconds>(
+                         Clock::now() - processStart)
+                         .count();
+      std::cout << "[PROFILE] process_total_ms=" << totalMs << std::endl;
+    }
+    return rc;
+  };
+
   int width = 16;
   int channels = 4;
   bool useLinear = false;
@@ -519,37 +560,37 @@ int main(int argc, char **argv) {
 
   // Dispatch to appropriate type
   if (type == "float")
-    return runTest<float>(width, channels, useLinear, useSemaphores,
-                          useSampled);
+    return finish(runTest<float>(width, channels, useLinear, useSemaphores,
+                                 useSampled));
   if (type == "half")
-    return runTest<sycl::half>(width, channels, useLinear, useSemaphores,
-                               useSampled);
+    return finish(runTest<sycl::half>(width, channels, useLinear,
+                                      useSemaphores, useSampled));
 
   if (type == "int32")
-    return runTest<int32_t>(width, channels, useLinear, useSemaphores,
-                            useSampled);
+    return finish(runTest<int32_t>(width, channels, useLinear, useSemaphores,
+                                   useSampled));
   if (type == "uint32")
-    return runTest<uint32_t>(width, channels, useLinear, useSemaphores,
-                             useSampled);
+    return finish(runTest<uint32_t>(width, channels, useLinear,
+                                    useSemaphores, useSampled));
 
   if (type == "int16")
-    return runTest<int16_t>(width, channels, useLinear, useSemaphores,
-                            useSampled);
+    return finish(runTest<int16_t>(width, channels, useLinear, useSemaphores,
+                                   useSampled));
   if (type == "uint16")
-    return runTest<uint16_t>(width, channels, useLinear, useSemaphores,
-                             useSampled);
+    return finish(runTest<uint16_t>(width, channels, useLinear,
+                                    useSemaphores, useSampled));
 
   if (type == "uint8")
-    return runTest<uint8_t>(width, channels, useLinear, useSemaphores,
-                            useSampled);
+    return finish(runTest<uint8_t>(width, channels, useLinear, useSemaphores,
+                                   useSampled));
   if (type == "int8")
-    return runTest<int8_t>(width, channels, useLinear, useSemaphores,
-                           useSampled);
+    return finish(runTest<int8_t>(width, channels, useLinear, useSemaphores,
+                                  useSampled));
 
   if (type == "unorm8") {
-    return runTest<uint8_t>(width, channels, useLinear, useSemaphores,
-                            useSampled, getUnorm8Format(channels),
-                            sycl::image_channel_type::unorm_int8);
+    return finish(runTest<uint8_t>(width, channels, useLinear, useSemaphores,
+                                   useSampled, getUnorm8Format(channels),
+                                   sycl::image_channel_type::unorm_int8));
   }
 
   std::cerr << "Unknown type: " << type << std::endl;
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_sampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_sampled.cpp
new file mode 100644
index 0000000000000..9cee6db306dbf
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_sampled.cpp
@@ -0,0 +1,39 @@
+// Sampled-only version of the Vulkan/SYCL 1D image read interop test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
+// REQUIRES: vulkan
+
+// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+
+// clang-format off
+// RUN: %{run} %t.out --type float --channels 1 --sampled 32
+// RUN: %{run} %t.out --type float --channels 2 --sampled 32
+// RUN: %{run} %t.out --type float --channels 4 --sampled 32
+// RUN: %{run} %t.out --type half --channels 1 --sampled 32
+// RUN: %{run} %t.out --type half --channels 2 --sampled 32
+// RUN: %{run} %t.out --type half --channels 4 --sampled 32
+// RUN: %{run} %t.out --type int32 --channels 1 --sampled 32
+// RUN: %{run} %t.out --type int32 --channels 2 --sampled 32
+// RUN: %{run} %t.out --type int32 --channels 4 --sampled 32
+// RUN: %{run} %t.out --type uint32 --channels 1 --sampled 32
+// RUN: %{run} %t.out --type uint32 --channels 2 --sampled 32
+// RUN: %{run} %t.out --type uint32 --channels 4 --sampled 32
+// RUN: %{run} %t.out --type int16 --channels 1 --sampled 32
+// RUN: %{run} %t.out --type int16 --channels 2 --sampled 32
+// RUN: %{run} %t.out --type int16 --channels 4 --sampled 32
+// RUN: %{run} %t.out --type uint16 --channels 1 --sampled 32
+// RUN: %{run} %t.out --type uint16 --channels 2 --sampled 32
+// RUN: %{run} %t.out --type uint16 --channels 4 --sampled 32
+// RUN: %{run} %t.out --type uint8 --channels 1 --sampled 32
+// RUN: %{run} %t.out --type uint8 --channels 2 --sampled 32
+// RUN: %{run} %t.out --type uint8 --channels 4 --sampled 32
+// RUN: %{run} %t.out --type int8 --channels 1 --sampled 32
+// RUN: %{run} %t.out --type int8 --channels 2 --sampled 32
+// RUN: %{run} %t.out --type int8 --channels 4 --sampled 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 1 --sampled 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 2 --sampled 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 4 --sampled 32
+// clang-format on
+
+#include "./vulkan_sycl_image_interop_read_1d_common.hpp"
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_semaphores.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_semaphores.cpp
new file mode 100644
index 0000000000000..75811685f7c57
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_semaphores.cpp
@@ -0,0 +1,27 @@
+// Semaphore version of the Vulkan/SYCL 1D image read interop test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
+// REQUIRES: vulkan
+
+// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+
+// clang-format off
+// RUN: %{run} %t.out --type float --channels 1 32 --semaphores
+// RUN: %{run} %t.out --type float --channels 2 32 --semaphores
+// RUN: %{run} %t.out --type float --channels 4 32 --semaphores
+// RUN: %{run} %t.out --type half --channels 1 32 --semaphores
+// RUN: %{run} %t.out --type int32 --channels 2 32 --semaphores
+// RUN: %{run} %t.out --type uint32 --channels 4 32 --semaphores
+// RUN: %{run} %t.out --type int16 --channels 1 32 --semaphores
+// RUN: %{run} %t.out --type uint16 --channels 2 32 --semaphores
+// RUN: %{run} %t.out --type uint8 --channels 4 32 --semaphores
+// RUN: %{run} %t.out --type int8 --channels 1 32 --semaphores
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 2 32 --semaphores
+// RUN: %{run} %t.out --type float --channels 4 --sampled 32 --semaphores
+// RUN: %{run} %t.out --type int16 --channels 4 --sampled 32 --semaphores
+// RUN: %{run} %t.out --type int8 --channels 4 --sampled 32 --semaphores
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 4 --sampled 32 --semaphores
+// clang-format on
+
+#include "./vulkan_sycl_image_interop_read_1d_common.hpp"
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_channels.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_channels.cpp
new file mode 100644
index 0000000000000..fa2017aa25683
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_channels.cpp
@@ -0,0 +1,43 @@
+// Unsampled channel-coverage version of the Vulkan/SYCL 2D image read interop
+// test
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
+// REQUIRES: vulkan
+
+// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+
+// UNSUPPORTED: linux
+// UNSUPPORTED-TRACKER: GSD-12357
+
+// clang-format off
+// RUN: %{run} %t.out --type float --channels 1 32x33
+// RUN: %{run} %t.out --type float --channels 2 32x33
+// RUN: %{run} %t.out --type float --channels 4 32x33
+// RUN: %{run} %t.out --type half --channels 1 32x33
+// RUN: %{run} %t.out --type half --channels 2 32x33
+// RUN: %{run} %t.out --type half --channels 4 32x33
+// RUN: %{run} %t.out --type int32 --channels 1 32x33
+// RUN: %{run} %t.out --type int32 --channels 2 32x33
+// RUN: %{run} %t.out --type int32 --channels 4 32x33
+// RUN: %{run} %t.out --type uint32 --channels 1 32x33
+// RUN: %{run} %t.out --type uint32 --channels 2 32x33
+// RUN: %{run} %t.out --type uint32 --channels 4 32x33
+// RUN: %{run} %t.out --type int16 --channels 1 32x33
+// RUN: %{run} %t.out --type int16 --channels 2 32x33
+// RUN: %{run} %t.out --type int16 --channels 4 32x33
+// RUN: %{run} %t.out --type uint16 --channels 1 32x33
+// RUN: %{run} %t.out --type uint16 --channels 2 32x33
+// RUN: %{run} %t.out --type uint16 --channels 4 32x33
+// RUN: %{run} %t.out --type uint8 --channels 1 32x33
+// RUN: %{run} %t.out --type uint8 --channels 2 32x33
+// RUN: %{run} %t.out --type uint8 --channels 4 32x33
+// RUN: %{run} %t.out --type int8 --channels 1 32x33
+// RUN: %{run} %t.out --type int8 --channels 2 32x33
+// RUN: %{run} %t.out --type int8 --channels 4 32x33
+// RUN: %{run} %t.out --type unorm8 --channels 1 32x33
+// RUN: %{run} %t.out --type unorm8 --channels 2 32x33
+// RUN: %{run} %t.out --type unorm8 --channels 4 32x33
+// clang-format on
+
+#include "./vulkan_sycl_image_interop_read_2d_common.hpp"
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_common.hpp
similarity index 77%
rename from sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp
rename to sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_common.hpp
index ac59681b4071b..bc213f7b487b3 100644
--- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_common.hpp
@@ -1,11 +1,4 @@
-// REQUIRES: aspect-ext_oneapi_bindless_images
-// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
-// REQUIRES: vulkan
-
-// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
-
-// UNSUPPORTED: linux
-// UNSUPPORTED-TRACKER: GSD-12357
+// Shared implementation for the Vulkan/SYCL 2D image read interop tests.
 
 // XFAIL: windows && gpu-intel-dg2
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/21985
@@ -14,87 +7,19 @@
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/21986
 
 /*
-    Run ALL the vulkan formats through the gauntlet. sampled and unsampled.
-    This entire test takes less than 30 seconds on a slow machine.  MUCH faster
-   (and more complete coveraage) than SFINAE based approach.
-
-    IF a particular variant is having problems on some platform, please do NOT
-   just disable the whole test, instead use   RUN~IF (SOMETHING) yadda-yadda
-    to enable/disable that variant.
+  Run ALL the vulkan formats through the gauntlet. sampled and unsampled.
+  This entire test takes less than 30 seconds on a slow machine.  MUCH faster
+  (and more complete coveraage) than SFINAE based approach.
 
-    For semaphore testing, we run just a sampling. Note, that on Linux if there
-   is a failure in the first section, then likely ALL semaphore tests afterwards
-   will fail. This is being tracked as a separate issue.
+  IF a particular variant is having problems on some platform, please do NOT
+  just disable the whole test, instead use   RUN~IF (SOMETHING) yadda-yadda
+  to enable/disable that variant.
 
+  For semaphore testing, we run just a sampling. Note, that on Linux if there
+  is a failure in the first section, then likely ALL semaphore tests afterwards
+  will fail. This is being tracked as a separate issue.
 */
 
-// RUN: %{run} %t.out --type float --channels 1 32x33
-// RUN: %{run} %t.out --type float --channels 2 32x33
-// RUN: %{run} %t.out --type float --channels 4 32x33
-// RUN: %{run} %t.out --type half --channels 1 32x33
-// RUN: %{run} %t.out --type half --channels 2 32x33
-// RUN: %{run} %t.out --type half --channels 4 32x33
-// RUN: %{run} %t.out --type int32 --channels 1 32x33
-// RUN: %{run} %t.out --type int32 --channels 2 32x33
-// RUN: %{run} %t.out --type int32 --channels 4 32x33
-// RUN: %{run} %t.out --type uint32 --channels 1 32x33
-// RUN: %{run} %t.out --type uint32 --channels 2 32x33
-// RUN: %{run} %t.out --type uint32 --channels 4 32x33
-// RUN: %{run} %t.out --type int16 --channels 1 32x33
-// RUN: %{run} %t.out --type int16 --channels 2 32x33
-// RUN: %{run} %t.out --type int16 --channels 4 32x33
-// RUN: %{run} %t.out --type uint16 --channels 1 32x33
-// RUN: %{run} %t.out --type uint16 --channels 2 32x33
-// RUN: %{run} %t.out --type uint16 --channels 4 32x33
-// RUN: %{run} %t.out --type uint8 --channels 1 32x33
-// RUN: %{run} %t.out --type uint8 --channels 2 32x33
-// RUN: %{run} %t.out --type uint8 --channels 4 32x33
-// RUN: %{run} %t.out --type int8 --channels 1 32x33
-// RUN: %{run} %t.out --type int8 --channels 2 32x33
-// RUN: %{run} %t.out --type int8 --channels 4 32x33
-// RUN: %{run} %t.out --type unorm8 --channels 1 32x33
-// RUN: %{run} %t.out --type unorm8 --channels 2 32x33
-// RUN: %{run} %t.out --type unorm8 --channels 4 32x33
-// RUN: %{run} %t.out --type float --channels 1 --sampled 32x33
-// RUN: %{run} %t.out --type float --channels 2 --sampled 32x33
-// RUN: %{run} %t.out --type float --channels 4 --sampled 32x33
-// RUN: %{run} %t.out --type half --channels 1 --sampled 32x33
-// RUN: %{run} %t.out --type half --channels 2 --sampled 32x33
-// RUN: %{run} %t.out --type half --channels 4 --sampled 32x33
-// RUN: %{run} %t.out --type int32 --channels 1 --sampled 32x33
-// RUN: %{run} %t.out --type int32 --channels 2 --sampled 32x33
-// RUN: %{run} %t.out --type int32 --channels 4 --sampled 32x33
-// RUN: %{run} %t.out --type uint32 --channels 1 --sampled 32x33
-// RUN: %{run} %t.out --type uint32 --channels 2 --sampled 32x33
-// RUN: %{run} %t.out --type uint32 --channels 4 --sampled 32x33
-// RUN: %{run} %t.out --type int16 --channels 1 --sampled 32x33
-// RUN: %{run} %t.out --type int16 --channels 2 --sampled 32x33
-// RUN: %{run} %t.out --type int16 --channels 4 --sampled 32x33
-// RUN: %{run} %t.out --type uint16 --channels 1 --sampled 32x33
-// RUN: %{run} %t.out --type uint16 --channels 2 --sampled 32x33
-// RUN: %{run} %t.out --type uint16 --channels 4 --sampled 32x33
-// RUN: %{run} %t.out --type uint8 --channels 1 --sampled 32x33
-// RUN: %{run} %t.out --type uint8 --channels 2 --sampled 32x33
-// RUN: %{run} %t.out --type uint8 --channels 4 --sampled 32x33
-// RUN: %{run} %t.out --type int8 --channels 1 --sampled 32x33
-// RUN: %{run} %t.out --type int8 --channels 2 --sampled 32x33
-// RUN: %{run} %t.out --type int8 --channels 4 --sampled 32x33
-// RUN: %{run} %t.out --type unorm8 --channels 1 --sampled 32x33
-// RUN: %{run} %t.out --type unorm8 --channels 2 --sampled 32x33
-// RUN: %{run} %t.out --type unorm8 --channels 4 --sampled 32x33
-
-// RUN: %{run} %t.out --type float --channels 1 32x33 --semaphores
-// RUN: %{run} %t.out --type float --channels 4 32x33 --semaphores
-// RUN: %{run} %t.out --type half --channels 1 32x33 --semaphores
-// RUN: %{run} %t.out --type uint32 --channels 4 32x33 --semaphores
-// RUN: %{run} %t.out --type uint16 --channels 2 32x33 --semaphores
-// RUN: %{run} %t.out --type int8 --channels 1 32x33 --semaphores
-// RUN: %{run} %t.out --type float --channels 4 --sampled 32x33 --semaphores
-// RUN: %{run} %t.out --type int32 --channels 4 --sampled 32x33 --semaphores
-// RUN: %{run} %t.out --type int16 --channels 4 --sampled 32x33 --semaphores
-// RUN: %{run} %t.out --type uint8 --channels 2 --sampled 32x33 --semaphores
-// RUN: %{run} %t.out --type unorm8 --channels 4 --sampled 32x33 --semaphores
-
 // clang-format off
 /*
   Vulkan/SYCL 2D Image Read Test (Sampled + Unsampled)
@@ -126,6 +51,8 @@
 */
 // clang-format on
 
+#pragma once
+
 #include "vulkan_setup.hpp"
 
 #include <optional>
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_sampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_sampled.cpp
new file mode 100644
index 0000000000000..868e79fa228d1
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_sampled.cpp
@@ -0,0 +1,42 @@
+// Sampled-only version of the Vulkan/SYCL 2D image read interop test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
+// REQUIRES: vulkan
+
+// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+
+// UNSUPPORTED: linux
+// UNSUPPORTED-TRACKER: GSD-12357
+
+// clang-format off
+// RUN: %{run} %t.out --type float --channels 1 --sampled 32x33
+// RUN: %{run} %t.out --type float --channels 2 --sampled 32x33
+// RUN: %{run} %t.out --type float --channels 4 --sampled 32x33
+// RUN: %{run} %t.out --type half --channels 1 --sampled 32x33
+// RUN: %{run} %t.out --type half --channels 2 --sampled 32x33
+// RUN: %{run} %t.out --type half --channels 4 --sampled 32x33
+// RUN: %{run} %t.out --type int32 --channels 1 --sampled 32x33
+// RUN: %{run} %t.out --type int32 --channels 2 --sampled 32x33
+// RUN: %{run} %t.out --type int32 --channels 4 --sampled 32x33
+// RUN: %{run} %t.out --type uint32 --channels 1 --sampled 32x33
+// RUN: %{run} %t.out --type uint32 --channels 2 --sampled 32x33
+// RUN: %{run} %t.out --type uint32 --channels 4 --sampled 32x33
+// RUN: %{run} %t.out --type int16 --channels 1 --sampled 32x33
+// RUN: %{run} %t.out --type int16 --channels 2 --sampled 32x33
+// RUN: %{run} %t.out --type int16 --channels 4 --sampled 32x33
+// RUN: %{run} %t.out --type uint16 --channels 1 --sampled 32x33
+// RUN: %{run} %t.out --type uint16 --channels 2 --sampled 32x33
+// RUN: %{run} %t.out --type uint16 --channels 4 --sampled 32x33
+// RUN: %{run} %t.out --type uint8 --channels 1 --sampled 32x33
+// RUN: %{run} %t.out --type uint8 --channels 2 --sampled 32x33
+// RUN: %{run} %t.out --type uint8 --channels 4 --sampled 32x33
+// RUN: %{run} %t.out --type int8 --channels 1 --sampled 32x33
+// RUN: %{run} %t.out --type int8 --channels 2 --sampled 32x33
+// RUN: %{run} %t.out --type int8 --channels 4 --sampled 32x33
+// RUN: %{run} %t.out --type unorm8 --channels 1 --sampled 32x33
+// RUN: %{run} %t.out --type unorm8 --channels 2 --sampled 32x33
+// RUN: %{run} %t.out --type unorm8 --channels 4 --sampled 32x33
+// clang-format on
+
+#include "./vulkan_sycl_image_interop_read_2d_common.hpp"
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_semaphores.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_semaphores.cpp
new file mode 100644
index 0000000000000..6068fa1fc8c94
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_semaphores.cpp
@@ -0,0 +1,26 @@
+// Semaphore version of the Vulkan/SYCL 2D image read interop test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
+// REQUIRES: vulkan
+
+// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+
+// UNSUPPORTED: linux
+// UNSUPPORTED-TRACKER: GSD-12357
+
+// clang-format off
+// RUN: %{run} %t.out --type float --channels 1 32x33 --semaphores
+// RUN: %{run} %t.out --type float --channels 4 32x33 --semaphores
+// RUN: %{run} %t.out --type half --channels 1 32x33 --semaphores
+// RUN: %{run} %t.out --type uint32 --channels 4 32x33 --semaphores
+// RUN: %{run} %t.out --type uint16 --channels 2 32x33 --semaphores
+// RUN: %{run} %t.out --type int8 --channels 1 32x33 --semaphores
+// RUN: %{run} %t.out --type float --channels 4 --sampled 32x33 --semaphores
+// RUN: %{run} %t.out --type int32 --channels 4 --sampled 32x33 --semaphores
+// RUN: %{run} %t.out --type int16 --channels 4 --sampled 32x33 --semaphores
+// RUN: %{run} %t.out --type uint8 --channels 2 --sampled 32x33 --semaphores
+// RUN: %{run} %t.out --type unorm8 --channels 4 --sampled 32x33 --semaphores
+// clang-format on
+
+#include "./vulkan_sycl_image_interop_read_2d_common.hpp"
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled_channels.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled_channels.cpp
new file mode 100644
index 0000000000000..dccfa78eb2dc2
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled_channels.cpp
@@ -0,0 +1,42 @@
+// Channel coverage version of the Vulkan/SYCL 1D unsampled write interop test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
+// REQUIRES: vulkan
+
+// UNSUPPORTED: windows
+// UNSUPPORTED-TRACKER: CMPLRLLVM-73525
+
+// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+
+// clang-format off
+// RUN: %{run} %t.out --type float --channels 1 32
+// RUN: %{run} %t.out --type float --channels 2 32
+// RUN: %{run} %t.out --type float --channels 4 32
+// RUN: %{run} %t.out --type half --channels 1 32
+// RUN: %{run} %t.out --type half --channels 2 32
+// RUN: %{run} %t.out --type half --channels 4 32
+// RUN: %{run} %t.out --type int32 --channels 1 32
+// RUN: %{run} %t.out --type int32 --channels 2 32
+// RUN: %{run} %t.out --type int32 --channels 4 32
+// RUN: %{run} %t.out --type uint32 --channels 1 32
+// RUN: %{run} %t.out --type uint32 --channels 2 32
+// RUN: %{run} %t.out --type uint32 --channels 4 32
+// RUN: %{run} %t.out --type int16 --channels 1 32
+// RUN: %{run} %t.out --type int16 --channels 2 32
+// RUN: %{run} %t.out --type int16 --channels 4 32
+// RUN: %{run} %t.out --type uint16 --channels 1 32
+// RUN: %{run} %t.out --type uint16 --channels 2 32
+// RUN: %{run} %t.out --type uint16 --channels 4 32
+// RUN: %{run} %t.out --type uint8 --channels 1 32
+// RUN: %{run} %t.out --type uint8 --channels 2 32
+// RUN: %{run} %t.out --type uint8 --channels 4 32
+// RUN: %{run} %t.out --type int8 --channels 1 32
+// RUN: %{run} %t.out --type int8 --channels 2 32
+// RUN: %{run} %t.out --type int8 --channels 4 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 1 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 2 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 4 32
+// clang-format on
+
+#include "./vulkan_sycl_image_interop_write_1d_unsampled_common.hpp"
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled_common.hpp
similarity index 85%
rename from sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled.cpp
rename to sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled_common.hpp
index aa8532eaf61a5..2b790b9a91bed 100644
--- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled.cpp
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled_common.hpp
@@ -1,12 +1,4 @@
-
-// REQUIRES: aspect-ext_oneapi_bindless_images
-// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
-// REQUIRES: vulkan
-
-// UNSUPPORTED: windows
-// UNSUPPORTED-TRACKER: CMPLRLLVM-73525
-
-// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+// Shared implementation for the Vulkan/SYCL 1D unsampled write interop tests.
 
 /*
   Run all the vulkan formats through a write test. Note this is unsampled only,
@@ -14,82 +6,41 @@
 
   IF a particular variant is having problems on some platform, please do NOT
   just disable the whole test, instead use   RUN~IF: (SOMETHING) yadda-yadda
-    to enable/disable that variant.
+  to enable/disable that variant.
 
-    For semaphore testing, we run just a sampling. Note, that on Linux if there
+  For semaphore testing, we run just a sampling. Note, that on Linux if there
   is a failure in the first section, then likely ALL semaphore tests afterwards
   will fail. This is being tracked as a separate issue.
-
 */
-// clang-format off
-
-// RUN: %{run} %t.out --type float --channels 1 32
-// RUN: %{run} %t.out --type float --channels 2 32
-// RUN: %{run} %t.out --type float --channels 4 32
-// RUN: %{run} %t.out --type half --channels 1 32
-// RUN: %{run} %t.out --type half --channels 2 32
-// RUN: %{run} %t.out --type half --channels 4 32
-// RUN: %{run} %t.out --type int32 --channels 1 32
-// RUN: %{run} %t.out --type int32 --channels 2 32
-// RUN: %{run} %t.out --type int32 --channels 4 32
-// RUN: %{run} %t.out --type uint32 --channels 1 32
-// RUN: %{run} %t.out --type uint32 --channels 2 32
-// RUN: %{run} %t.out --type uint32 --channels 4 32
-// RUN: %{run} %t.out --type int16 --channels 1 32
-// RUN: %{run} %t.out --type int16 --channels 2 32
-// RUN: %{run} %t.out --type int16 --channels 4 32
-// RUN: %{run} %t.out --type uint16 --channels 1 32
-// RUN: %{run} %t.out --type uint16 --channels 2 32
-// RUN: %{run} %t.out --type uint16 --channels 4 32
-// RUN: %{run} %t.out --type uint8 --channels 1 32
-// RUN: %{run} %t.out --type uint8 --channels 2 32
-// RUN: %{run} %t.out --type uint8 --channels 4 32
-// RUN: %{run} %t.out --type int8 --channels 1 32
-// RUN: %{run} %t.out --type int8 --channels 2 32
-// RUN: %{run} %t.out --type int8 --channels 4 32
-// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 1 32
-// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 2 32
-// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 4 32
-
-// On Linux L0, there are problem with semaphores and latest drivers.
-// GSD-12371 GSD-12339
-
-// RUN-IF: !level_zero, %{run} %t.out --type float --channels 1 32 --semaphores
-// RUN-IF: !level_zero, %{run} %t.out --type half --channels 2 32 --semaphores
-// RUN-IF: !level_zero, %{run} %t.out --type int32 --channels 4 32 --semaphores
-// RUN-IF: !level_zero, %{run} %t.out --type uint32 --channels 1 32 --semaphores
-// RUN-IF: !level_zero, %{run} %t.out --type int16 --channels 2 32 --semaphores
-// RUN-IF: !level_zero, %{run} %t.out --type uint16 --channels 4 32 --semaphores
-// RUN-IF: !level_zero, %{run} %t.out --type uint8 --channels 1 32 --semaphores
-// RUN-IF: !level_zero, %{run} %t.out --type int8 --channels 4 32 --semaphores
-// CUDA doesn't support unorm8, level_zero has issues with semaphores
-// XXX-IF: !cuda, %{run} %t.out --type unorm8 --channels 2 32 --semaphores
-
 
+// clang-format off
 /*
-  Vulkan/SYCL  1D Unsampled Write Image
-
+  Vulkan/SYCL 1D Unsampled Write Image
 
   clang++ -fsycl -o vsw_1d_test.bin vulkan_sycl_image_interop_write_1d_unsampled.cpp -lvulkan -I$VULKAN_SDK/include -L$VULKAN_SDK/lib
 
   clang++ -fsycl -o vsw_1d_test.exe vulkan_sycl_image_interop_write_1d_unsampled.cpp -Wno-ignored-attributes -lvulkan-1 -I$VULKAN_SDK/Include -L$VULKAN_SDK/Lib
 
-
+  USAGE:
     ./vsw_1d_test.bin
     ./vsw_1d_test.bin --semaphores
-    FLAGS
+  
+  FLAGS:
     --sampled      ERROR: Sampled image writes are not supported
     --semaphores   Use Vulkan Semaphores for SYCL Interop Sync
     --linear       Use LINEAR tiling for the Vulkan Image (default is OPTIMAL)
     --channels  X  Set number of channels (1, 2, or 4). Default is 4 (RGBA)
     --type  XXX    Set data type (float, half, uint32, int32, uint16, int16, uint8, int8, unorm8). Default is float 
     Wx             Set custom Width .  Put "x" after
-
+  
+  EXAMPLES:
     ./vsw_1d_test.bin --semaphores --channels 2 --linear 64x
 
- */
+*/
 // clang-format on
 
+#pragma once
+
 #include "vulkan_setup.hpp"
 
 #include <optional>
@@ -464,4 +415,4 @@ int main(int argc, char **argv) {
                             getUnorm8Format(channels),
                             sycl::image_channel_type::unorm_int8);
   return 1;
-}
\ No newline at end of file
+}
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled_semaphores.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled_semaphores.cpp
new file mode 100644
index 0000000000000..35a1384c1aeca
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled_semaphores.cpp
@@ -0,0 +1,26 @@
+// Semaphore coverage version of the Vulkan/SYCL 1D unsampled write interop
+// test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
+// REQUIRES: vulkan
+
+// UNSUPPORTED: windows
+// UNSUPPORTED-TRACKER: CMPLRLLVM-73525
+
+// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+
+// clang-format off
+// RUN-IF: !level_zero, %{run} %t.out --type float --channels 1 32 --semaphores
+// RUN-IF: !level_zero, %{run} %t.out --type half --channels 2 32 --semaphores
+// RUN-IF: !level_zero, %{run} %t.out --type int32 --channels 4 32 --semaphores
+// RUN-IF: !level_zero, %{run} %t.out --type uint32 --channels 1 32 --semaphores
+// RUN-IF: !level_zero, %{run} %t.out --type int16 --channels 2 32 --semaphores
+// RUN-IF: !level_zero, %{run} %t.out --type uint16 --channels 4 32 --semaphores
+// RUN-IF: !level_zero, %{run} %t.out --type uint8 --channels 1 32 --semaphores
+// RUN-IF: !level_zero, %{run} %t.out --type int8 --channels 4 32 --semaphores
+// CUDA doesn't support unorm8, level_zero has issues with semaphores
+// XXX-IF: !cuda, %{run} %t.out --type unorm8 --channels 2 32 --semaphores
+// clang-format on
+
+#include "./vulkan_sycl_image_interop_write_1d_unsampled_common.hpp"
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled_channels.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled_channels.cpp
new file mode 100644
index 0000000000000..21c8e9a52260d
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled_channels.cpp
@@ -0,0 +1,42 @@
+// Channel coverage version of the Vulkan/SYCL 2D unsampled write interop test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
+// REQUIRES: vulkan
+
+// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+
+// UNSUPPORTED: linux
+// UNSUPPORTED-TRACKER: GSD-12357
+
+// clang-format off
+// RUN: %{run} %t.out --type float --channels 1 32x33
+// RUN: %{run} %t.out --type float --channels 2 32x33
+// RUN: %{run} %t.out --type float --channels 4 32x33
+// RUN: %{run} %t.out --type half --channels 1 32x33
+// RUN: %{run} %t.out --type half --channels 2 32x33
+// RUN: %{run} %t.out --type half --channels 4 32x33
+// RUN: %{run} %t.out --type int32 --channels 1 32x33
+// RUN: %{run} %t.out --type int32 --channels 2 32x33
+// RUN: %{run} %t.out --type int32 --channels 4 32x33
+// RUN: %{run} %t.out --type uint32 --channels 1 32x33
+// RUN: %{run} %t.out --type uint32 --channels 2 32x33
+// RUN: %{run} %t.out --type uint32 --channels 4 32x33
+// RUN: %{run} %t.out --type int16 --channels 1 32x33
+// RUN: %{run} %t.out --type int16 --channels 2 32x33
+// RUN: %{run} %t.out --type int16 --channels 4 32x33
+// RUN: %{run} %t.out --type uint16 --channels 1 32x33
+// RUN: %{run} %t.out --type uint16 --channels 2 32x33
+// RUN: %{run} %t.out --type uint16 --channels 4 32x33
+// RUN: %{run} %t.out --type uint8 --channels 1 32x33
+// RUN: %{run} %t.out --type uint8 --channels 2 32x33
+// RUN: %{run} %t.out --type uint8 --channels 4 32x33
+// RUN: %{run} %t.out --type int8 --channels 1 32x33
+// RUN: %{run} %t.out --type int8 --channels 2 32x33
+// RUN: %{run} %t.out --type int8 --channels 4 32x33
+// RUN: %{run} %t.out --type unorm8 --channels 1 32x33
+// RUN: %{run} %t.out --type unorm8 --channels 2 32x33
+// RUN: %{run} %t.out --type unorm8 --channels 4 32x33
+// clang-format on
+
+#include "./vulkan_sycl_image_interop_write_2d_unsampled_common.hpp"
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled_common.hpp
similarity index 85%
rename from sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp
rename to sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled_common.hpp
index 3644fd78b7fd0..03a0b7929b3a3 100644
--- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled_common.hpp
@@ -1,11 +1,4 @@
-// REQUIRES: aspect-ext_oneapi_bindless_images
-// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
-// REQUIRES: vulkan
-
-// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
-
-// UNSUPPORTED: linux
-// UNSUPPORTED-TRACKER: GSD-12357
+// Shared implementation for the Vulkan/SYCL 2D unsampled write interop tests.
 
 // XFAIL: windows && gpu-intel-dg2
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/21985
@@ -14,56 +7,19 @@
 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/21986
 
 /*
-    Run all the vulkan formats through a write test. Note this is unsampled
-   only, you can't "write" with the image sampler.
 
-    IF a particular variant is having problems on some platform, please do NOT
-   just disable the whole test, instead use   RUN~IF: (SOMETHING) yadda-yadda
-    to enable/disable that variant.
+  Run all the vulkan formats through a write test. Note this is unsampled
+  only, you can't "write" with the image sampler.
 
-    For semaphore testing, we run just a sampling. Note, that on Linux if there
-   is a failure in the first section, then likely ALL semaphore tests afterwards
-   will fail. This is being tracked as a separate issue.
+  IF a particular variant is having problems on some platform, please do NOT
+  just disable the whole test, instead use   RUN~IF: (SOMETHING) yadda-yadda
+  to enable/disable that variant.
 
-*/
+  For semaphore testing, we run just a sampling. Note, that on Linux if there
+  is a failure in the first section, then likely ALL semaphore tests afterwards
+  will fail. This is being tracked as a separate issue.
 
-// RUN: %{run} %t.out --type float --channels 1 32x33
-// RUN: %{run} %t.out --type float --channels 2 32x33
-// RUN: %{run} %t.out --type float --channels 4 32x33
-// RUN: %{run} %t.out --type half --channels 1 32x33
-// RUN: %{run} %t.out --type half --channels 2 32x33
-// RUN: %{run} %t.out --type half --channels 4 32x33
-// RUN: %{run} %t.out --type int32 --channels 1 32x33
-// RUN: %{run} %t.out --type int32 --channels 2 32x33
-// RUN: %{run} %t.out --type int32 --channels 4 32x33
-// RUN: %{run} %t.out --type uint32 --channels 1 32x33
-// RUN: %{run} %t.out --type uint32 --channels 2 32x33
-// RUN: %{run} %t.out --type uint32 --channels 4 32x33
-// RUN: %{run} %t.out --type int16 --channels 1 32x33
-// RUN: %{run} %t.out --type int16 --channels 2 32x33
-// RUN: %{run} %t.out --type int16 --channels 4 32x33
-// RUN: %{run} %t.out --type uint16 --channels 1 32x33
-// RUN: %{run} %t.out --type uint16 --channels 2 32x33
-// RUN: %{run} %t.out --type uint16 --channels 4 32x33
-// RUN: %{run} %t.out --type uint8 --channels 1 32x33
-// RUN: %{run} %t.out --type uint8 --channels 2 32x33
-// RUN: %{run} %t.out --type uint8 --channels 4 32x33
-// RUN: %{run} %t.out --type int8 --channels 1 32x33
-// RUN: %{run} %t.out --type int8 --channels 2 32x33
-// RUN: %{run} %t.out --type int8 --channels 4 32x33
-// RUN: %{run} %t.out --type unorm8 --channels 1 32x33
-// RUN: %{run} %t.out --type unorm8 --channels 2 32x33
-// RUN: %{run} %t.out --type unorm8 --channels 4 32x33
-
-// RUN: %{run} %t.out --type float --channels 1 32x33 --semaphores
-// RUN: %{run} %t.out --type half --channels 2 32x33 --semaphores
-// RUN: %{run} %t.out --type int32 --channels 4 32x33 --semaphores
-// RUN: %{run} %t.out --type uint32 --channels 1 32x33 --semaphores
-// RUN: %{run} %t.out --type int16 --channels 2 32x33 --semaphores
-// RUN: %{run} %t.out --type uint16 --channels 4 32x33 --semaphores
-// RUN: %{run} %t.out --type uint8 --channels 1 32x33 --semaphores
-// RUN: %{run} %t.out --type int8 --channels 2 32x33 --semaphores
-// RUN: %{run} %t.out --type unorm8 --channels 4 32x33 --semaphores
+*/
 
 // clang-format off
 /*
@@ -71,9 +27,10 @@
 
   clang++ -fsycl -o vsw_2d_test.exe vulkan_sycl_image_interop_write_2d_unsampled.cpp -Wno-ignored-attributes -lvulkan-1 -I$VULKAN_SDK/Include -L$VULKAN_SDK/Lib
 
+  USAGE:
     ./vsw_2d_test.bin
 
-    FLAGS
+  FLAGS:
     --sampled      ERROR: Sampled image writes are not supported
     --semaphores   Use Vulkan Semaphores for SYCL Interop Sync
     --linear       Use LINEAR tiling for the Vulkan Image (default is OPTIMAL)
@@ -82,11 +39,14 @@
                    Default is float 
     WxH            Set custom Width x Height (e.g. 8x4)
 
+  EXAMPLES:
     ./vsw_2d_test.bin --semaphores --channels 2 --linear 8x4
 
 */
 // clang-format on
 
+#pragma once
+
 #include "vulkan_setup.hpp"
 
 #include <optional>
diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled_semaphores.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled_semaphores.cpp
new file mode 100644
index 0000000000000..b587871577c98
--- /dev/null
+++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled_semaphores.cpp
@@ -0,0 +1,25 @@
+// Semaphore coverage version of the Vulkan/SYCL 2D unsampled write interop
+// test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
+// REQUIRES: vulkan
+
+// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+
+// UNSUPPORTED: linux
+// UNSUPPORTED-TRACKER: GSD-12357
+
+// clang-format off
+// RUN: %{run} %t.out --type float --channels 1 32x33 --semaphores
+// RUN: %{run} %t.out --type half --channels 2 32x33 --semaphores
+// RUN: %{run} %t.out --type int32 --channels 4 32x33 --semaphores
+// RUN: %{run} %t.out --type uint32 --channels 1 32x33 --semaphores
+// RUN: %{run} %t.out --type int16 --channels 2 32x33 --semaphores
+// RUN: %{run} %t.out --type uint16 --channels 4 32x33 --semaphores
+// RUN: %{run} %t.out --type uint8 --channels 1 32x33 --semaphores
+// RUN: %{run} %t.out --type int8 --channels 2 32x33 --semaphores
+// RUN: %{run} %t.out --type unorm8 --channels 4 32x33 --semaphores
+// clang-format on
+
+#include "./vulkan_sycl_image_interop_write_2d_unsampled_common.hpp"