[GPU] SYCL infra for GPU plugin and GHA pipeline (openvinotoolkit#26067)

Co-authored-by: Ilya Lavrenov <[email protected]>

.github/dockerfiles/docker_tag

@@ -1 +1 @@
-pr-25992
+pr-26067

.github/dockerfiles/ov_build/ubuntu_22_04_x64_dpcpp/Dockerfile

@@ -0,0 +1,95 @@
+FROM openvinogithubactions.azurecr.io/dockerhub/ubuntu:22.04
+
+USER root
+
+# APT configuration
+RUN echo 'Acquire::Retries "10";' > /etc/apt/apt.conf && \
+    echo 'APT::Get::Assume-Yes "true";' >> /etc/apt/apt.conf && \
+    echo 'APT::Get::Fix-Broken "true";' >> /etc/apt/apt.conf && \
+    echo 'APT::Get::no-install-recommends "true";' >> /etc/apt/apt.conf
+
+ENV DEBIAN_FRONTEND="noninteractive" \
+    TZ="Europe/London"
+
+RUN apt-get update && \
+    apt-get install software-properties-common && \
+    add-apt-repository --yes --no-update ppa:git-core/ppa && \
+    add-apt-repository --yes --no-update ppa:deadsnakes/ppa && \
+    apt-get update && \
+    apt-get install \
+        wget \
+        curl \
+        git \
+        ca-certificates \
+        gpg-agent \
+        tzdata \
+        libtbb2 \
+        # Pythons \
+        python3.11-dev \
+        python3.11-venv \
+        python3.11-distutils \
+        default-jdk \
+        && \
+    rm -rf /var/lib/apt/lists/*
+
+
+# Install OneAPI Toolkit
+RUN wget -O- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB | gpg --dearmor | tee /usr/share/keyrings/intel-oneapi-archive-keyring.gpg > /dev/null
+RUN echo "deb [signed-by=/usr/share/keyrings/intel-oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main " > /etc/apt/sources.list.d/oneAPI.list
+
+RUN apt-get update && \
+    apt-get install \
+        intel-oneapi-compiler-dpcpp-cpp=2024.2.1-1079 \
+        && \
+    rm -rf /var/lib/apt/lists/*
+
+# Install build dependencies
+ADD install_build_dependencies.sh /install_build_dependencies.sh
+RUN chmod +x /install_build_dependencies.sh && \
+    /install_build_dependencies.sh && \
+    rm -rf /var/lib/apt/lists/*
+
+# Install sscache
+ARG SCCACHE_VERSION="v0.7.5"
+ENV SCCACHE_HOME="/opt/sccache" \
+    SCCACHE_PATH="/opt/sccache/sccache"
+
+RUN mkdir ${SCCACHE_HOME} && cd ${SCCACHE_HOME} && \
+    SCCACHE_ARCHIVE="sccache-${SCCACHE_VERSION}-x86_64-unknown-linux-musl.tar.gz" && \
+    curl -SLO https://github.com/mozilla/sccache/releases/download/${SCCACHE_VERSION}/${SCCACHE_ARCHIVE} && \
+    tar -xzf ${SCCACHE_ARCHIVE} --strip-components=1 && rm ${SCCACHE_ARCHIVE}
+
+# Setup pip
+ENV PIP_VERSION="24.0"
+RUN curl https://bootstrap.pypa.io/get-pip.py -o get-pip.py && \
+    python3 get-pip.py --no-cache-dir pip==${PIP_VERSION} && \
+    python3.11 get-pip.py --no-cache-dir pip==${PIP_VERSION} && \
+    rm -f get-pip.py
+
+# Use Python 3.11 as default
+# Using venv here 'cause other methods to switch the default Python on Ubuntu break both system and wheels build
+RUN python3.11 -m venv venv
+ENV PATH="/venv/bin:$SCCACHE_HOME:$PATH"
+
+ENV PIP_CACHE_DIR=/mount/caches/pip/linux/${PIP_VERSION}
+
+# OneAPI env
+ENV ONEAPI_ROOT=/opt/intel/oneapi
+ENV PKG_CONFIG_PATH=/opt/intel/oneapi/compiler/2024.2/lib/pkgconfig
+ENV DIAGUTIL_PATH=/opt/intel/oneapi/dpcpp-ct/2024.2/etc/dpct/sys_check/sys_check.sh:/opt/intel/oneapi/debugger/2024.2/etc/debugger/sys_check/sys_check.py:/opt/intel/oneapi/compiler/2024.2/etc/compiler/sys_check/sys_check.sh
+ENV MANPATH=/opt/intel/oneapi/debugger/2024.2/share/man:/opt/intel/oneapi/compiler/2024.2/share/man:
+ENV GDB_INFO=/opt/intel/oneapi/debugger/2024.2/share/info/
+ENV CMAKE_PREFIX_PATH=/opt/intel/oneapi/compiler/2024.2
+ENV CMPLR_ROOT=/opt/intel/oneapi/compiler/2024.2
+ENV INFOPATH=/opt/intel/oneapi/debugger/2024.2/share/info
+ENV LIBRARY_PATH=/opt/intel/oneapi/compiler/2024.2/lib
+ENV OCL_ICD_FILENAMES=/opt/intel/oneapi/compiler/2024.2/lib/libintelocl.so
+ENV LD_LIBRARY_PATH=/opt/intel/oneapi/debugger/2024.2/opt/debugger/lib:/opt/intel/oneapi/compiler/2024.2/opt/compiler/lib:/opt/intel/oneapi/compiler/2024.2/lib
+ENV NLSPATH=/opt/intel/oneapi/mkl/2024.2/share/locale/%l_%t/%N:/opt/intel/oneapi/compiler/2024.2/lib/compiler/locale/%l_%t/%N
+ENV PATH=$PATH:/opt/intel/oneapi/dev-utilities/2024.2/bin:/opt/intel/oneapi/debugger/2024.2/opt/debugger/bin:/opt/intel/oneapi/compiler/2024.2/bin
+ENV INTEL_PYTHONHOME=/opt/intel/oneapi/debugger/2024.2/opt/debugger
+ENV CPATH=/opt/intel/oneapi/dpl/2022.6/include:/opt/intel/oneapi/dev-utilities/2024.2/include:
+
+# Set Intel DPC++ as a default compiler
+ENV CC=icx
+ENV CXX=icpx

.github/workflows/ubuntu_22_dpcpp.yml

@@ -0,0 +1,90 @@
+name: Linux (Ubuntu 22.04, Python 3.11, Intel DPC++ Compiler)
+on:
+  workflow_dispatch:
+  pull_request:
+  merge_group:
+
+concurrency:
+  # github.ref is not unique in post-commit
+  group: ${{ github.event_name == 'push' && github.run_id || github.ref }}-ubuntu-22-dpcpp
+  cancel-in-progress: true
+
+permissions: read-all
+
+jobs:
+  Smart_CI:
+    runs-on: ubuntu-latest
+    outputs:
+      affected_components: "${{ steps.smart_ci.outputs.affected_components }}"
+      changed_components: "${{ steps.smart_ci.outputs.changed_components }}"
+      skip_workflow: "${{ steps.smart_ci.outputs.skip_workflow }}"
+    steps:
+      - name: checkout action
+        uses: actions/checkout@692973e3d937129bcbf40652eb9f2f61becf3332 # v4.1.7
+        with:
+          sparse-checkout: .github/actions/smart-ci
+
+      - name: Get affected components
+        id: smart_ci
+        uses: ./.github/actions/smart-ci
+        with:
+          repository: ${{ github.repository }}
+          pr: ${{ github.event.number }}
+          commit_sha: ${{ github.sha }}
+          ref_name: ${{ github.ref_name }}
+          component_pattern: "category: (.*)"
+          repo_token: ${{ secrets.GITHUB_TOKEN }}
+          skip_when_only_listed_labels_set: 'docs'
+          skip_when_only_listed_files_changed: '*.md,*.rst,*.png,*.jpg,*.svg'
+
+      - name: Show affected components
+        run: |
+          echo "${{ toJSON(steps.smart_ci.outputs.affected_components) }}"
+        shell: bash
+
+  Docker:
+    needs: Smart_CI
+    runs-on: aks-linux-4-cores-16gb-docker-build
+    container:
+      image: openvinogithubactions.azurecr.io/docker_build:0.2
+      volumes:
+        - /mount:/mount
+    outputs:
+      images: "${{ steps.handle_docker.outputs.images }}"
+    steps:
+      - name: Checkout
+        uses: actions/checkout@692973e3d937129bcbf40652eb9f2f61becf3332 # v4.1.7
+
+      - uses: ./.github/actions/handle_docker
+        id: handle_docker
+        with:
+          images: |
+            ov_build/ubuntu_22_04_x64_dpcpp
+          registry: 'openvinogithubactions.azurecr.io'
+          dockerfiles_root_dir: '.github/dockerfiles'
+          changed_components: ${{ needs.smart_ci.outputs.changed_components }}
+
+  Build:
+    needs: [Docker, Smart_CI]
+    if: "!needs.smart_ci.outputs.skip_workflow"
+    uses: ./.github/workflows/job_build_linux.yml
+    with:
+      runner: 'aks-linux-16-cores-32gb'
+      container: '{"image": "${{ fromJSON(needs.docker.outputs.images).ov_build.ubuntu_22_04_x64_dpcpp }}", "volumes": ["/mount:/mount"], "options": "-e SCCACHE_AZURE_BLOB_CONTAINER -e SCCACHE_AZURE_CONNECTION_STRING"}'
+      affected-components: ${{ needs.smart_ci.outputs.affected_components }}
+      event-name: ${{ github.event_name }}
+      os: 'ubuntu_22_04_dpcpp'
+
+  Overall_Status:
+    name: ci/gha_overall_status_ubuntu_22.04_dpcpp
+    needs: [Smart_CI, Build]
+    if: ${{ always() }}
+    runs-on: ubuntu-latest
+    steps:
+      - name: Check status of all jobs
+        if: >-
+          ${{
+            contains(needs.*.result, 'failure') ||
+            contains(needs.*.result, 'cancelled')
+          }}
+        run: exit 1

cmake/developer_package/compile_flags/os_flags.cmake

@@ -289,6 +289,8 @@ endif()
 if(NOT DEFINED CMAKE_CXX_STANDARD)
     if(CMAKE_CXX_COMPILER_ID STREQUAL "MSVC")
         set(CMAKE_CXX_STANDARD 14)
+    elseif(OV_COMPILER_IS_INTEL_LLVM)
+        set(CMAKE_CXX_STANDARD 17)
     else()
         set(CMAKE_CXX_STANDARD 11)
     endif()

src/plugins/intel_cpu/src/nodes/kernels/scaled_attn/executor_pa_common.hpp

@@ -34,6 +34,7 @@ struct PagedAttentionExecutor {
     static const size_t ID_ALIBI_SLOPES = 11;               // [H|0], float
     static const size_t ID_MAX_CONTEXT_LEN = 12;            // []
     virtual void execute(const std::vector<ov::intel_cpu::MemoryPtr>& inputs, const std::vector<ov::intel_cpu::MemoryPtr> outputs) = 0;
+    virtual ~PagedAttentionExecutor() = default;
 };
 
 #ifdef OPENVINO_ARCH_X86_64

src/plugins/intel_gpu/CMakeLists.txt

@@ -8,6 +8,10 @@ endif()
 
 set (TARGET_NAME "openvino_intel_gpu_plugin")
 
+if(OV_COMPILER_IS_INTEL_LLVM)
+    find_package(IntelSYCL REQUIRED)
+endif()
+
 if((CMAKE_COMPILER_IS_GNUCXX OR OV_COMPILER_IS_CLANG) AND CMAKE_CXX_STANDARD GREATER_EQUAL 20)
     set(CMAKE_CXX_FLAGS "-Wno-error=deprecated ${CMAKE_CXX_FLAGS}")
 endif()

src/plugins/intel_gpu/include/intel_gpu/graph/kernel_impl_params.hpp

@@ -90,7 +90,7 @@ struct kernel_impl_params final {
                        , primary_input_idx(0) {
     }
 
-    virtual ~kernel_impl_params() = default;
+    ~kernel_impl_params() = default;
 
     const layout& get_input_layout(size_t idx = 0) const {
         OPENVINO_ASSERT(input_layouts.size() > idx,

src/plugins/intel_gpu/include/intel_gpu/primitives/implementation_desc.hpp

@@ -18,6 +18,7 @@ enum class impl_types : uint8_t {
     common = 1 << 1,
     ocl = 1 << 2,
     onednn = 1 << 3,
+    sycl = 1 << 4,
     any = 0xFF,
 };
 

src/plugins/intel_gpu/include/intel_gpu/runtime/engine_configuration.hpp

@@ -11,11 +11,13 @@ namespace cldnn {
 /// @brief Defines available engine types
 enum class engine_types : int32_t {
     ocl,
+    sycl
 };
 
 inline std::ostream& operator<<(std::ostream& os, const engine_types& type) {
     switch (type) {
     case engine_types::ocl: os << "ocl"; break;
+    case engine_types::sycl: os << "sycl"; break;
     default: os << "unknown"; break;
     }
 

src/plugins/intel_gpu/src/graph/CMakeLists.txt

@@ -4,6 +4,7 @@
 
 set(TARGET_NAME "openvino_intel_gpu_graph")
 
+file(GLOB_RECURSE SYCL_SOURCES "${CMAKE_CURRENT_SOURCE_DIR}/impls/sycl/*")
 file(GLOB_RECURSE LIBRARY_SRC
     "${INCLUDE_DIR}/*.h"
     "${INCLUDE_DIR}/*.hpp"
@@ -22,6 +23,10 @@ if (NOT ENABLE_ONEDNN_FOR_GPU)
     endforeach(SOURCE_FILE)
 endif()
 
+if(NOT OV_COMPILER_IS_INTEL_LLVM)
+    list(REMOVE_ITEM LIBRARY_SRC ${SYCL_SOURCES})
+endif()
+
 add_library(${TARGET_NAME} STATIC ${LIBRARY_SRC})
 
 target_include_directories(${TARGET_NAME} PUBLIC
@@ -33,6 +38,11 @@ target_include_directories(${TARGET_NAME} PUBLIC
 target_compile_options(${TARGET_NAME} PRIVATE
   $<$<CONFIG:Release>:$<IF:$<CXX_COMPILER_ID:MSVC>,/Os,-Os>>)
 
+if(OV_COMPILER_IS_INTEL_LLVM)
+    add_sycl_to_target(TARGET ${TARGET_NAME} SOURCES ${SYCL_SOURCES})
+    target_compile_definitions(${TARGET_NAME} PUBLIC OV_GPU_WITH_SYCL)
+endif()
+
 target_link_libraries(${TARGET_NAME} PUBLIC OpenCL::OpenCL openvino::shape_inference)
 target_link_libraries(${TARGET_NAME} PRIVATE openvino_intel_gpu_kernels
                                              openvino_intel_gpu_runtime

src/plugins/intel_gpu/src/graph/impls/sycl/README.md

@@ -0,0 +1,16 @@
+# How to build with DPC++ support
+
+1. Install OneAPI base toolkit. Guide: https://www.intel.com/content/www/us/en/docs/oneapi/installation-guide-linux/2024-0/installation.html
+2. Export environment:
+    $ source /opt/intel/oneapi/setvars.sh
+
+3. configure cmake with the following additional options:
+    - `-DCMAKE_CXX_FLAGS:STRING=--gcc-install-dir=/lib/gcc/x86_64-linux-gnu/12/ -DCMAKE_C_FLAGS:STRING=--gcc-install-dir=/lib/gcc/x86_64-linux-gnu/12/`
+        - This WA is needed if multiple GCC version available in the system
+    - `-DCMAKE_CXX_STANDARD:STRING=17`
+        - Sycl requires c++17
+    - `-DENABLE_SYSTEM_OPENCL=OFF`
+        - May help to avoid opencl icd/header conflicts as sycl package may have no clhpp headers
+    - `-DCMAKE_C_COMPILER:FILEPATH=icx -DCMAKE_CXX_COMPILER:FILEPATH=icpx`
+        - For now find_package(IntelSYCL) doesn't work if compiler is not icpx
+4. make -j$(nproc)

src/plugins/intel_gpu/src/graph/impls/sycl/primitive_sycl_base.h

@@ -0,0 +1,90 @@
+// Copyright (C) 2024 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma once
+
+#include "primitive_inst.h"
+#include "intel_gpu/runtime/memory.hpp"
+#include "register.hpp"
+#include "utils.hpp"
+#include "runtime/ocl/ocl_event.hpp"
+
+#include <vector>
+
+#include "sycl/sycl.hpp"
+
+namespace cldnn {
+namespace sycl {
+
+static std::mutex cacheAccessMutex;
+
+template <class PType>
+struct typed_primitive_sycl_impl : public typed_primitive_impl<PType> {
+    const engine* _engine;
+
+    typed_primitive_sycl_impl(const engine& engine, const ExecutionConfig& config, std::shared_ptr<WeightsReorderParams> weights_reorder = nullptr)
+        : typed_primitive_impl<PType>(weights_reorder, "sycl_kernel"),
+        _engine(&engine) { }
+
+    typed_primitive_sycl_impl()
+        : typed_primitive_impl<PType>({}, "undef"),
+          _engine(nullptr) {
+    }
+
+    bool is_cpu() const override { return false; }
+    bool is_onednn() const override { return false; }
+
+    void save(BinaryOutputBuffer& ob) const override {
+    }
+
+    void load(BinaryInputBuffer& ib) override {
+    }
+
+protected:
+    void init_kernels(const kernels_cache&, const kernel_impl_params&) override { }
+
+    void set_arguments_impl(typed_primitive_inst<PType>& instance) override {
+        if (instance.can_be_optimized())
+            return;
+    }
+
+    void update_dispatch_data(const kernel_impl_params& impl_params) override {}
+
+    void set_arguments_impl(typed_primitive_inst<PType>& instance, kernel_arguments_data& args) override {
+        if (instance.can_be_optimized()) {
+            return;
+        }
+    }
+
+    event::ptr execute_impl(const std::vector<event::ptr>& /* events */,
+                            typed_primitive_inst<PType>& instance) override {
+        auto& network = instance.get_network();
+        auto& stream = network.get_stream();
+        auto net_id = network.get_id();
+        event::ptr event;
+
+
+        return event;
+    }
+
+    static event::ptr to_ocl_event(stream& stream, ::sycl::event e) {
+        if (stream.get_queue_type() == QueueTypes::out_of_order) {
+            auto native_events = get_native<::sycl::backend::opencl, ::sycl::event>(e);
+            std::vector<event::ptr> events;
+            for (auto& e : native_events) {
+                events.push_back(std::make_shared<ocl::ocl_event>(cl::Event(e, true)));
+            }
+            return events.empty() ? stream.create_user_event(true) : stream.group_events(events);
+        } else {
+            return stream.create_user_event(true);
+        }
+    }
+
+    std::vector<layout> get_internal_buffer_layouts_impl() const override {
+        return {};
+    }
+};
+
+}  // namespace sycl
+}  // namespace cldnn