From 5832d33709b61685c82799dab6007e1b1d75b257 Mon Sep 17 00:00:00 2001 From: Rafal Rudnicki Date: Wed, 22 Apr 2026 19:52:48 +0000 Subject: [PATCH 1/4] $ git rebase --continue sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_common.hpp: needs merge You must edit all merge conflicts and then mark them as resolved using git addRevert "Revert "Modify CI to enable tests on CUDA 13.1 and add Dockerfiles for CUDA 13.1" (#21816)" This reverts commit 91080cf6b4dbe12ae1e714aafcd1216fd384692d. --- .github/workflows/sycl-containers.yaml | 7 ++ .github/workflows/sycl-linux-precommit.yml | 14 ++++ .github/workflows/ur-build-hw.yml | 1 + .github/workflows/ur-precommit.yml | 8 +++ .../ubuntu2404_build_cuda131.Dockerfile | 71 +++++++++++++++++++ ...buntu2404_intel_drivers_cuda131.Dockerfile | 25 +++++++ 6 files changed, 126 insertions(+) create mode 100644 devops/containers/ubuntu2404_build_cuda131.Dockerfile create mode 100644 devops/containers/ubuntu2404_intel_drivers_cuda131.Dockerfile 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"] + From 5d6566c7f0f5a145515a355e0427c398c0e8f2e2 Mon Sep 17 00:00:00 2001 From: Rafal Rudnicki Date: Fri, 24 Apr 2026 09:01:28 +0000 Subject: [PATCH 2/4] [SYCL][E2E] Split vulkan_sycl_image_interop tests --- ...an_sycl_image_interop_read_1d_channels.cpp | 39 ++++ ...kan_sycl_image_interop_read_1d_common.hpp} | 171 +++++------------- ...kan_sycl_image_interop_read_1d_sampled.cpp | 39 ++++ ..._sycl_image_interop_read_1d_semaphores.cpp | 27 +++ ...an_sycl_image_interop_read_2d_channels.cpp | 43 +++++ ...kan_sycl_image_interop_read_2d_common.hpp} | 97 ++-------- ...kan_sycl_image_interop_read_2d_sampled.cpp | 42 +++++ ..._sycl_image_interop_read_2d_semaphores.cpp | 26 +++ ...ge_interop_write_1d_unsampled_channels.cpp | 42 +++++ ...age_interop_write_1d_unsampled_common.hpp} | 77 ++------ ..._interop_write_1d_unsampled_semaphores.cpp | 26 +++ ...ge_interop_write_2d_unsampled_channels.cpp | 42 +++++ ...age_interop_write_2d_unsampled_common.hpp} | 70 ++----- ..._interop_write_2d_unsampled_semaphores.cpp | 25 +++ 14 files changed, 442 insertions(+), 324 deletions(-) create mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_channels.cpp rename sycl/test-e2e/bindless_images/vulkan_interop/{vulkan_sycl_image_interop_read_1d.cpp => vulkan_sycl_image_interop_read_1d_common.hpp} (71%) create mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_sampled.cpp create mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_semaphores.cpp create mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_channels.cpp rename sycl/test-e2e/bindless_images/vulkan_interop/{vulkan_sycl_image_interop_read_2d.cpp => vulkan_sycl_image_interop_read_2d_common.hpp} (77%) create mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_sampled.cpp create mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d_semaphores.cpp create mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled_channels.cpp rename sycl/test-e2e/bindless_images/vulkan_interop/{vulkan_sycl_image_interop_write_1d_unsampled.cpp => vulkan_sycl_image_interop_write_1d_unsampled_common.hpp} (85%) create mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled_semaphores.cpp create mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled_channels.cpp rename sycl/test-e2e/bindless_images/vulkan_interop/{vulkan_sycl_image_interop_write_2d_unsampled.cpp => vulkan_sycl_image_interop_write_2d_unsampled_common.hpp} (85%) create mode 100644 sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled_semaphores.cpp 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 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 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 @@ -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 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" From 8d93092ff17ab38f1f588cab8836405942516980 Mon Sep 17 00:00:00 2001 From: Rafal Rudnicki Date: Fri, 24 Apr 2026 09:02:02 +0000 Subject: [PATCH 3/4] [SYCL][E2E] Split bindless image read_sampled test --- .../bindless_images/read_sampled_1d.cpp | 31 ++++++++++++++++ .../bindless_images/read_sampled_2d.cpp | 34 +++++++++++++++++ ...ad_sampled.cpp => read_sampled_common.hpp} | 37 +++++-------------- 3 files changed, 75 insertions(+), 27 deletions(-) create mode 100644 sycl/test-e2e/bindless_images/read_sampled_1d.cpp create mode 100644 sycl/test-e2e/bindless_images/read_sampled_2d.cpp rename sycl/test-e2e/bindless_images/{read_sampled.cpp => read_sampled_common.hpp} (97%) 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 #include #include +#include +#include #include @@ -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 dims, sycl::range 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); } From 294e39f4909cbc6a79ef3d504d2ca17c3bd6b2dd Mon Sep 17 00:00:00 2001 From: Rafal Rudnicki Date: Fri, 24 Apr 2026 09:02:27 +0000 Subject: [PATCH 4/4] [SYCL][E2E] Split WG Memory basic usage test --- ...basic_usage.cpp => basic_usage_common.hpp} | 28 +++---------------- .../WorkGroupMemory/basic_usage_test.cpp | 26 +++++++++++++++++ .../WorkGroupMemory/basic_usage_test_ptr.cpp | 26 +++++++++++++++++ 3 files changed, 56 insertions(+), 24 deletions(-) rename sycl/test-e2e/WorkGroupMemory/{basic_usage.cpp => basic_usage_common.hpp} (95%) create mode 100644 sycl/test-e2e/WorkGroupMemory/basic_usage_test.cpp create mode 100644 sycl/test-e2e/WorkGroupMemory/basic_usage_test_ptr.cpp 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 #include #include #include #include +#include #include #include 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 void test_ptr() { } swap_array_2d(arr1, arr2, 8); } - -int main() { - test(); - test(); - test(); - if (q.get_device().has(sycl::aspect::fp16)) - test(); - test_ptr(); - test_ptr(); - test_ptr(); - test_ptr(); - if (q.get_device().has(sycl::aspect::fp16)) - test_ptr(); - test_ptr(); - 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 q; + +#include "./basic_usage_common.hpp" + +int main() { + test(); + test(); + test(); + test(); + if (q.get_device().has(sycl::aspect::fp16)) + test(); + 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 q; + +#include "./basic_usage_common.hpp" + +int main() { + test_ptr(); + test_ptr(); + test_ptr(); + test_ptr(); + if (q.get_device().has(sycl::aspect::fp16)) + test_ptr(); + return 0; +}