[GPU] Sycl op example, base class cleanup, docs update (#26304)

Vladimir Paramuzov · web-flow · commit fd25ea4a2ec3 · 2024-08-30T07:46:11.000Z
### Details:
- Added naive FC sycl impl as an example and it will also be used to
verify base class compilation in CI util we have any production sycl op.
 - Moved sycl readme to common gpu docs folder with some minor updates
diff --git a/src/plugins/intel_gpu/docs/gpu_plugin_with_sycl.md b/src/plugins/intel_gpu/docs/gpu_plugin_with_sycl.md
@@ -0,0 +1,21 @@
+# 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:
+    - [Linux] `-DCMAKE_C_COMPILER:FILEPATH=icx -DCMAKE_CXX_COMPILER:FILEPATH=icpx`
+      [Windows] `-DCMAKE_C_COMPILER:FILEPATH=icx -DCMAKE_CXX_COMPILER:FILEPATH=icx`
+        - For now find_package(IntelSYCL) doesn't work if compiler is not icpx, so we need to update compilers globally for the whole project
+    - `-DENABLE_INTEL_CPU=OFF`
+        - OneAPI toolkit with OneDNN installed may cause CPU plugin build issue due to weird include files resolver which prefer system onednn intead of
+          CPU fork which causes build issue. Alternatively, OneDNN can be removed from OneAPI toolkit installation.
+    - [Linux] `-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
+    - `-DENABLE_SYSTEM_OPENCL=OFF`
+        - May help to avoid opencl icd/header conflicts as sycl package may have no clhpp headers
+    - `-DCMAKE_CXX_COMPILER_LAUNCHER=ccache`
+        - For some reason with latest OneAPI package versions each `make` call causes full project recompilation, so the recommendation is to enable caching
+
+4. cmake --build . --config Release --parallel
diff --git a/src/plugins/intel_gpu/src/graph/impls/sycl/README.md b/src/plugins/intel_gpu/src/graph/impls/sycl/README.md
diff --git a/src/plugins/intel_gpu/src/graph/impls/sycl/impl_example.cpp b/src/plugins/intel_gpu/src/graph/impls/sycl/impl_example.cpp
@@ -0,0 +1,262 @@
+// Copyright (C) 2024 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include "fully_connected_inst.h"
+#include "intel_gpu/primitives/reorder.hpp"
+#include "ocl/ocl_event.hpp"
+#include "ocl/sycl_engine.hpp"
+#include "ocl/sycl_stream.hpp"
+#include "openvino/core/type/element_type.hpp"
+#include "primitive_sycl_base.h"
+#include "impls/registry/implementation_map.hpp"
+
+#include "impls/ocl/kernel_selector_helper.h"
+
+#include "sycl/sycl.hpp"
+#include "sycl/ext/oneapi/experimental/builtins.hpp"
+
+#include <memory>
+
+#ifdef __SYCL_DEVICE_ONLY__
+          #define CONSTANT __attribute__((opencl_constant))
+#else
+          #define CONSTANT
+#endif
+
+namespace cldnn {
+namespace sycl {
+
+template <typename A, typename B>
+struct AccumulatorType {
+    using type = float;
+};
+
+template<> struct AccumulatorType<::sycl::half, ::sycl::half> {
+    using type = ::sycl::half;
+};
+
+template<> struct AccumulatorType<::sycl::half, uint8_t> {
+    using type = ::sycl::half;
+};
+
+
+template<> struct AccumulatorType<::sycl::half, int8_t> {
+    using type = ::sycl::half;
+};
+
+template<typename AType, typename WType, typename ZPType, typename ScaleType, typename DType>
+::sycl::event run_fc_int4_woq(::sycl::queue& queue, bool enqueue_barrier, const AType* a, const WType* w, const ZPType* zp, const ScaleType* s, DType* dst,
+                              size_t M, size_t N, size_t K, size_t group_size, size_t groups_num, const ov::Shape& out_shape, optional_value<float> dzp_s) {
+    if (enqueue_barrier) {
+        queue.submit([=](::sycl::handler& cgh) {
+            cgh.ext_oneapi_barrier();
+        });
+    }
+
+    bool has_value = dzp_s.has_value();
+    float dzp_value = dzp_s.value_or(0.0f);
+    return queue.submit([=](::sycl::handler& cgh) {
+        cgh.parallel_for(::sycl::range<3>(out_shape[0], out_shape[1], out_shape[2]), [=](::sycl::id<3> index) {
+            const uint b = index[0];
+            const uint m = index[1];
+            const uint n = index[2];
+            using accum_t = typename AccumulatorType<AType, WType>::type;
+            accum_t accumulator = 0.0f;
+
+            const uint dst_index = n + m*N + b*N*M;
+            for (uint y = 0; y < K; ++y) {
+                const uint input0_offset = y + m*K + b*M*K;
+                const uint decomp_offset = (y / group_size % groups_num)*N + n % N;
+                const uint filter_offset = y + n*K;
+                const uint zp_offset = 0;
+
+
+                accum_t zp_val = has_value ? static_cast<accum_t>(dzp_value) : static_cast<accum_t>(zp[zp_offset]);
+                accum_t scale = s[decomp_offset];
+                const WType packed = w[filter_offset / 2];
+
+                const WType v0 = packed & 0x0F;
+                const WType v1 = (packed & 0xF0) >> 4;
+                accum_t unpacked = filter_offset % 2 == 0 ? v0 : v1;
+
+                accum_t filter_val = (unpacked - zp_val) * scale;
+                accumulator += a[input0_offset] * filter_val;
+            }
+            dst[dst_index] = accumulator;
+        });
+    });
+}
+
+template<typename AType, typename WType, typename ZPType, typename ScaleType, typename DType>
+::sycl::event run_fc_int8_woq(::sycl::queue& queue, bool enqueue_barrier, const AType* a, const WType* w, const ZPType* zp, const ScaleType* s, DType* dst,
+                     size_t M, size_t N, size_t K, size_t group_size, size_t groups_num, const ov::Shape& out_shape, optional_value<float> dzp_s) {
+    if (enqueue_barrier) {
+        queue.submit([=](::sycl::handler& cgh) {
+            cgh.ext_oneapi_barrier();
+        });
+    }
+
+    bool has_value = dzp_s.has_value();
+    float dzp_value = dzp_s.value_or(0.0f);
+
+    return queue.submit([=](::sycl::handler& cgh) {
+        cgh.parallel_for(::sycl::range<3>(out_shape[0], out_shape[1], out_shape[2]), [=](::sycl::id<3> index) {
+            const uint b = index[0];
+            const uint m = index[1];
+            const uint n = index[2];
+            using accum_t = typename AccumulatorType<AType, WType>::type;
+            accum_t accumulator = 0.0f;
+
+            for (uint y = 0; y < K; ++y) {
+                const uint input0_offset = y + m*K + b*M*K;
+                const uint zp_offset = (y / group_size % groups_num)*N + n % N;
+                const uint decomp_offset = (y / group_size % groups_num)*N + n % N;
+                const uint filter_offset = y + n*K;
+
+                accum_t zp_val = has_value ? static_cast<accum_t>(dzp_value) : static_cast<accum_t>(zp[zp_offset]);
+                accum_t scale = s[decomp_offset];
+                accum_t filter_compressed = static_cast<accum_t>(w[filter_offset]);
+                accum_t filter_val = (filter_compressed - zp_val) * scale;
+                accumulator += a[input0_offset] * filter_val;
+            }
+            const uint dst_index = n + m*N + b*N*M;
+            dst[dst_index] = accumulator;
+        });
+    });
+}
+
+struct fully_connected_sycl_example : typed_primitive_sycl_impl<fully_connected> {
+    using parent = typed_primitive_sycl_impl<fully_connected>;
+    using parent::parent;
+
+    DECLARE_OBJECT_TYPE_SERIALIZATION(cldnn::sycl::fully_connected_sycl_example)
+
+    std::unique_ptr<primitive_impl> clone() const override {
+        return make_unique<fully_connected_sycl_example>(*this);
+    }
+
+    event::ptr execute_impl(const std::vector<event::ptr>& /* events */, typed_primitive_inst<fully_connected>& instance) override {
+        auto& network = instance.get_network();
+        const auto& desc = instance.get_typed_desc<fully_connected>();
+
+        auto& stream = downcast<ocl::sycl_stream>(network.get_stream());
+        auto& engine = downcast<ocl::sycl_engine>(network.get_engine());
+        ::sycl::context sycl_context = engine.get_sycl_context();
+        ::sycl::queue& sycl_queue = stream.get_sycl_queue();
+
+        const auto& params = instance.get_impl_params();
+        auto out_shape = params->output_layouts[0].get_shape();
+
+        auto output = instance.output_memory_ptr(0);
+        auto weights = instance.weights_memory();
+        auto bias = instance.bias_term() ? instance.bias_memory() : nullptr;
+
+        std::vector<memory::ptr> inputs = { instance.input_memory_ptr(0) };
+        size_t in_id = instance.bias_term() ? 3 : 2;
+        if (!desc->decompression_scale.empty())
+            inputs.push_back(instance.dep_memory_ptr(in_id++));
+
+        if (!desc->decompression_zero_point.empty())
+            inputs.push_back(instance.dep_memory_ptr(in_id));
+
+        OPENVINO_ASSERT(!instance.bias_term() && !instance.get_node().has_fused_primitives());
+
+        ov::element::Type_t in_t = params->input_layouts[0].data_type;
+        ov::element::Type_t wei_t = params->weights_layout.value().data_type;
+        ov::element::Type_t out_t = params->output_layouts[0].data_type;
+        ov::element::Type_t ds_t = params->input_layouts[2].data_type;
+        ov::element::Type_t dzp_t = inputs.size() == 3 ? params->input_layouts[3].data_type : ov::element::Type_t::undefined;
+
+        OPENVINO_ASSERT(out_shape.size() == 3);
+        size_t M = out_shape[1];
+        size_t N = out_shape[2];
+        size_t K = params->weights_layout.value().get_partial_shape()[1].get_length();
+        size_t groups_num = params->input_layouts[2].get_shape()[1];
+        size_t group_size = K / groups_num;
+
+        OPENVINO_ASSERT(inputs.size() >= 2);
+
+        auto dzp_scalar = desc->decompression_zero_point_scalar;
+
+        bool barrier = stream.get_queue_type() == QueueTypes::out_of_order;
+
+        #define CASE(InputType, WeightsType, ZPType, ScaleType, DstType) \
+            in_t == ov::element::InputType && \
+            wei_t == ov::element::WeightsType && \
+            out_t == ov::element::DstType && \
+            ds_t == ov::element::ScaleType && \
+            dzp_t == ov::element::ZPType
+
+        if ((CASE(f32, u4, f32, f32, f32)) || (CASE(f32, u4, undefined, f32, f32))) {
+            const float* in = static_cast<const float*>(inputs[0]->buffer_ptr());
+            const uint8_t* wei = static_cast<const uint8_t*>(weights->buffer_ptr());
+            float* out = static_cast<float*>(output->buffer_ptr());
+            const float* ds = static_cast<const float*>(inputs[1]->buffer_ptr());
+            const float* dzp = inputs.size() == 3 ? static_cast<const float*>(inputs[2]->buffer_ptr()) : nullptr;
+
+            return to_ocl_event(stream, run_fc_int4_woq(sycl_queue, barrier, in, wei, dzp, ds, out, M, N, K, group_size, groups_num, out_shape, dzp_scalar));
+        } else if ((CASE(f16, u4, f16, f16, f16)) || (CASE(f16, u4, undefined, f16, f16))) {
+            const ::sycl::half* in = static_cast<const ::sycl::half*>(inputs[0]->buffer_ptr());
+            const uint8_t* wei = static_cast<const uint8_t*>(weights->buffer_ptr());
+            ::sycl::half* out = static_cast<::sycl::half*>(output->buffer_ptr());
+            const ::sycl::half* ds = static_cast<const ::sycl::half*>(inputs[1]->buffer_ptr());
+            const ::sycl::half* dzp = inputs.size() == 3 ? static_cast<const ::sycl::half*>(inputs[2]->buffer_ptr()) : nullptr;
+
+
+            return to_ocl_event(stream, run_fc_int4_woq(sycl_queue, barrier, in, wei, dzp, ds, out, M, N, K, group_size, groups_num, out_shape, dzp_scalar));
+        } else if ((CASE(f16, u4, f16, f16, f32)) || (CASE(f16, u4, undefined, f16, f32))) {
+            const ::sycl::half* in = static_cast<const ::sycl::half*>(inputs[0]->buffer_ptr());
+            const uint8_t* wei = static_cast<const uint8_t*>(weights->buffer_ptr());
+            float* out = static_cast<float*>(output->buffer_ptr());
+            const ::sycl::half* ds = static_cast<const ::sycl::half*>(inputs[1]->buffer_ptr());
+            const ::sycl::half* dzp = inputs.size() == 3 ? static_cast<const ::sycl::half*>(inputs[2]->buffer_ptr()) : nullptr;
+
+
+            return to_ocl_event(stream, run_fc_int4_woq(sycl_queue, barrier, in, wei, dzp, ds, out, M, N, K, group_size, groups_num, out_shape, dzp_scalar));
+        } else if ((CASE(f32, u8, f32, f32, f32)) || (CASE(f32, u8, undefined, f32, f32))) {
+            const float* in = static_cast<const float*>(inputs[0]->buffer_ptr());
+            const uint8_t* wei = static_cast<const uint8_t*>(weights->buffer_ptr());
+            float* out = static_cast<float*>(output->buffer_ptr());
+            const float* ds = static_cast<const float*>(inputs[1]->buffer_ptr());
+            const float* dzp = inputs.size() == 3 ? static_cast<const float*>(inputs[2]->buffer_ptr()) : nullptr;
+
+            return to_ocl_event(stream, run_fc_int8_woq(sycl_queue, barrier, in, wei, dzp, ds, out, M, N, K, group_size, groups_num, out_shape, dzp_scalar));
+        } else if ((CASE(f16, u8, f16, f16, f16)) || (CASE(f16, u8, undefined, f16, f16))) {
+            const ::sycl::half* in = static_cast<const ::sycl::half*>(inputs[0]->buffer_ptr());
+            const uint8_t* wei = static_cast<const uint8_t*>(weights->buffer_ptr());
+            ::sycl::half* out = static_cast<::sycl::half*>(output->buffer_ptr());
+            const ::sycl::half* ds = static_cast<const ::sycl::half*>(inputs[1]->buffer_ptr());
+            const ::sycl::half* dzp = inputs.size() == 3 ? static_cast<const ::sycl::half*>(inputs[2]->buffer_ptr()) : nullptr;
+
+            return to_ocl_event(stream, run_fc_int8_woq(sycl_queue, barrier, in, wei, dzp, ds, out, M, N, K, group_size, groups_num, out_shape, dzp_scalar));
+        } else if ((CASE(f16, u8, f16, f16, f32)) || (CASE(f16, u8, undefined, f16, f32))) {
+            const ::sycl::half* in = static_cast<const ::sycl::half*>(inputs[0]->buffer_ptr());
+            const uint8_t* wei = static_cast<const uint8_t*>(weights->buffer_ptr());
+            float* out = static_cast<float*>(output->buffer_ptr());
+            const ::sycl::half* ds = static_cast<const ::sycl::half*>(inputs[1]->buffer_ptr());
+            const ::sycl::half* dzp = inputs.size() == 3 ? static_cast<const ::sycl::half*>(inputs[2]->buffer_ptr()) : nullptr;
+
+            return to_ocl_event(stream, run_fc_int8_woq(sycl_queue, barrier, in, wei, dzp, ds, out, M, N, K, group_size, groups_num, out_shape, dzp_scalar));
+        } else {
+            OPENVINO_THROW("No instance for given types found: ", in_t, " ", wei_t, " ", out_t, " ", ds_t, " ", dzp_t);
+        }
+    }
+
+    static std::shared_ptr<WeightsReorderParams> get_weights_reorder(const kernel_impl_params& impl_params) {
+        auto source_weights_layout = impl_params.get_input_layout(1);
+        auto target_weights_layout = source_weights_layout;
+        target_weights_layout.format = format::oiyx;
+
+        return std::make_shared<WeightsReorderParams>(source_weights_layout, target_weights_layout);
+    }
+
+    static std::unique_ptr<primitive_impl> create(const fully_connected_node& arg, const kernel_impl_params& impl_params) {
+        auto& engine = impl_params.prog->get_engine();
+        auto& config = impl_params.prog->get_config();
+        return cldnn::make_unique<fully_connected_sycl_example>(engine, config, get_weights_reorder(impl_params));
+    }
+};
+
+}  // namespace sycl
+}  // namespace cldnn
diff --git a/src/plugins/intel_gpu/src/graph/impls/sycl/primitive_sycl_base.h b/src/plugins/intel_gpu/src/graph/impls/sycl/primitive_sycl_base.h
@@ -7,7 +7,6 @@
 #include "primitive_inst.h"
 #include "intel_gpu/runtime/memory.hpp"
 #include "register.hpp"
-#include "utils.hpp"
 #include "runtime/ocl/ocl_event.hpp"
 
 #include <vector>
@@ -27,50 +26,21 @@ struct typed_primitive_sycl_impl : public typed_primitive_impl<PType> {
         : typed_primitive_impl<PType>(weights_reorder, "sycl_kernel"),
         _engine(&engine) { }
 
-    typed_primitive_sycl_impl()
-        : typed_primitive_impl<PType>({}, "undef"),
-          _engine(nullptr) {
+    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;
-    }
+    void set_arguments_impl(typed_primitive_inst<PType>& instance) override { }
+    void set_arguments_impl(typed_primitive_inst<PType>& instance, kernel_arguments_data& args) override { }
 
     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);
+            auto native_events = ::sycl::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)));
