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..78a46f5ee4019 100644
--- a/.github/workflows/ur-precommit.yml
+++ b/.github/workflows/ur-precommit.yml
@@ -73,6 +73,14 @@ jobs:
           - 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: cuda13
+            docker_image: "ghcr.io/intel/llvm/ubuntu2404_intel_drivers_cuda131:latest"
+            image_options: -u 1001 --privileged --cap-add SYS_ADMIN --gpus all
+            
+            # 12.6.3
           - name: CUDA
             runner: UR_CUDA
             image_options: -u 1001 --privileged --cap-add SYS_ADMIN --gpus all
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 71%
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..301fbcb3ee774 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,6 +85,8 @@ VK_FORMAT_R8G8B8A8_UNORM
 */
 // 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_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"