[GPU] Multinomial-13 (#20616)

* Update random_uniform in multinomial reference implementation to generate values of the same type as the first input

* Add kernel selector for multinomial

* Add multinomial primitive and its instantiation

* Add calc_outut_layout implementation for slice_inst

* Multinomial builder

* Add multinomial kernel

* Add multinomial single layer test class

* Add GPU multinomial tests instantiation

src/core/reference/include/openvino/reference/multinomial.hpp

@@ -94,9 +94,9 @@ void multinomial(const T* probs,
     divide(cdf.data(), max_value_per_batch_divisor.data(), cdf.data(), total_inputs_elements_count, false);
 
     // Generate random probability samples
-    std::vector<double> uniform_samples(total_output_elements_count);
-    const double zero = 0;
-    const double one = 1;
+    std::vector<T> uniform_samples(total_output_elements_count);
+    const T zero = 0;
+    const T one = 1;
     const ov::Shape output_shape_shape{output_shape.size()};
     const std::vector<uint64_t> output_shape_u64(output_shape.begin(), output_shape.end());
     const std::pair<uint64_t, uint64_t> initial_state(0, 0);
@@ -105,7 +105,7 @@ void multinomial(const T* probs,
                    reinterpret_cast<const char*>(&one),
                    reinterpret_cast<char*>(uniform_samples.data()),
                    output_shape_shape,
-                   ov::element::f64,
+                   ov::element::from<T>(),
                    global_seed,
                    op_seed,
                    initial_state);

src/plugins/intel_gpu/include/intel_gpu/plugin/primitives_list.hpp

@@ -261,6 +261,9 @@ REGISTER_FACTORY(v12, GroupNormalization);
 REGISTER_FACTORY(v12, Pad);
 REGISTER_FACTORY(v12, ScatterElementsUpdate);
 
+// ------------------------------ Supported v13 ops ----------------------------- //
+REGISTER_FACTORY(v13, Multinomial);
+
 // --------------------------- Supported internal ops --------------------------- //
 REGISTER_FACTORY(internal, NonMaxSuppressionIEInternal);
 REGISTER_FACTORY(internal, GenerateProposalsIEInternal);

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

@@ -0,0 +1,90 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#pragma once
+#include "primitive.hpp"
+
+namespace cldnn {
+
+/// @brief
+/// @details
+struct multinomial : public primitive_base<multinomial> {
+    CLDNN_DECLARE_PRIMITIVE(multinomial)
+
+    multinomial() : primitive_base("", {}) {}
+
+    /// @brief Constructs multinomial primitive.
+    /// @param id This primitive id.
+    /// @param cdf Cumulative distribution of probabilties.
+    //  @param random_probabilities Random probability samples.
+    multinomial(const primitive_id& id,
+          const input_info& cdf,
+          const input_info& random_probabilities,
+          data_types output_data_type,
+          bool with_replacement,
+          bool log_probs,
+          std::uint64_t global_seed,
+          std::uint64_t op_seed,
+          std::int64_t num_samples,
+          const padding& output_padding = padding())
+        : primitive_base{id, {cdf, random_probabilities}, {output_padding}},
+          output_data_type {output_data_type},
+          with_replacement {with_replacement},
+          log_probs {log_probs},
+          global_seed {global_seed},
+          op_seed {global_seed},
+          num_samples {num_samples}
+    {}
+
+    data_types output_data_type;
+    bool with_replacement;
+    bool log_probs;
+    std::uint64_t global_seed;
+    std::uint64_t op_seed;
+    std::int64_t num_samples;
+
+    std::size_t hash() const override {
+        size_t seed = primitive::hash();
+        seed = hash_combine(seed, output_data_type);
+        seed = hash_combine(seed, with_replacement);
+        seed = hash_combine(seed, log_probs);
+        seed = hash_combine(seed, global_seed);
+        seed = hash_combine(seed, op_seed);
+        return hash_combine(seed, num_samples);
+    }
+
+    bool operator==(const primitive& rhs) const override {
+        if (!compare_common_params(rhs))
+            return false;
+        const multinomial& rhs_casted = downcast<const multinomial>(rhs);
+        return output_data_type == rhs_casted.output_data_type &&
+            with_replacement == rhs_casted.with_replacement &&
+            log_probs == rhs_casted.log_probs &&
+            global_seed == rhs_casted.global_seed &&
+            op_seed == rhs_casted.op_seed &&
+            num_samples == rhs_casted.num_samples;
+    }
+
+    void save(BinaryOutputBuffer& ob) const override {
+        primitive_base<multinomial>::save(ob);
+        ob << ov::element::Type(output_data_type).to_string();
+        ob << with_replacement;
+        ob << log_probs;
+        ob << global_seed;
+        ob << op_seed;
+        ob << num_samples;
+    }
+
+    void load(BinaryInputBuffer& ib) override {
+        primitive_base<multinomial>::load(ib);
+        std::string data_type;
+        ib >> data_type;
+        output_data_type = ov::element::Type(data_type);
+        ib >> with_replacement;
+        ib >> log_probs;
+        ib >> global_seed;
+        ib >> op_seed;
+        ib >> num_samples;
+    }
+};
+}  // namespace cldnn

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

@@ -19,26 +19,20 @@ struct slice : public primitive_base<slice> {
     /// @param inputs List of primitive ids.
     slice(const primitive_id& id,
           const std::vector<input_info>& inputs,
-          const tensor output_shape,
           const padding& output_padding = padding())
-        : primitive_base{id, inputs, {output_padding}},
-          output_shape {output_shape}
+        : primitive_base{id, inputs, {output_padding}}
     {}
 
-    tensor output_shape;
-
     bool operator==(const primitive& rhs) const override {
         return compare_common_params(rhs);
     }
 
     void save(BinaryOutputBuffer& ob) const override {
         primitive_base<slice>::save(ob);
-        ob << output_shape;
     }
 
     void load(BinaryInputBuffer& ib) override {
         primitive_base<slice>::load(ib);
-        ib >> output_shape;
     }
 };
 }  // namespace cldnn

src/plugins/intel_gpu/src/graph/impls/ocl/multinomial.cpp

@@ -0,0 +1,56 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#include "primitive_base.hpp"
+#include "multinomial_inst.h"
+#include "multinomial/multinomial_kernel_ref.h"
+#include "multinomial/multinomial_kernel_selector.h"
+
+namespace cldnn {
+namespace ocl {
+
+struct multinomial_impl : typed_primitive_impl_ocl<multinomial> {
+    using parent = typed_primitive_impl_ocl<multinomial>;
+    using parent::parent;
+    using kernel_selector_t = kernel_selector::multinomial_kernel_selector;
+    using kernel_params_t = std::pair<kernel_selector::multinomial_params, kernel_selector::multinomial_optional_params>;
+
+    DECLARE_OBJECT_TYPE_SERIALIZATION(cldnn::ocl::multinomial_impl)
+
+    std::unique_ptr<primitive_impl> clone() const override {
+        return make_unique<multinomial_impl>(*this);
+    }
+
+    static kernel_params_t get_kernel_params(const kernel_impl_params& impl_param, bool is_shape_agnostic = false) {
+        const auto& primitive = impl_param.typed_desc<multinomial>();
+        auto params = get_default_params<kernel_selector::multinomial_params>(impl_param, is_shape_agnostic);
+        params.inputs.push_back(convert_data_tensor(impl_param.get_input_layout(1)));
+        auto optional_params = get_default_optional_params<kernel_selector::multinomial_optional_params>(impl_param.get_program());
+        params.output_data_type = primitive->output_data_type;
+        params.with_replacement = primitive->with_replacement;
+        params.log_probs = primitive->log_probs;
+        return {params, optional_params};
+    }
+
+    void update_dispatch_data(const kernel_impl_params& impl_param) override {
+        auto kernel_params = get_kernel_params(impl_param, true);
+        (_kernel_data.update_dispatch_data_func)(kernel_params.first, _kernel_data);
+    }
+};
+
+namespace detail {
+
+attach_multinomial_impl::attach_multinomial_impl() {
+    auto types = {data_types::f16, data_types::f32};
+    implementation_map<multinomial>::add(impl_types::ocl, shape_types::static_shape,
+                                     typed_primitive_impl_ocl<multinomial>::create<multinomial_impl>,
+                                     types,
+                                     {format::bfyx});
+}
+
+}  // namespace detail
+}  // namespace ocl
+}  // namespace cldnn
+
+BIND_BINARY_BUFFER_WITH_TYPE(cldnn::ocl::multinomial_impl)
+BIND_BINARY_BUFFER_WITH_TYPE(cldnn::multinomial)

src/plugins/intel_gpu/src/graph/impls/ocl/register.cpp

@@ -44,6 +44,7 @@ void register_implementations() {
     REGISTER_OCL(lstm_gemm);
     REGISTER_OCL(lstm_elt);
     REGISTER_OCL(multiclass_nms);
+    REGISTER_OCL(multinomial);
     REGISTER_OCL(mutable_data);
     REGISTER_OCL(mvn);
     REGISTER_OCL(non_max_suppression);

src/plugins/intel_gpu/src/graph/impls/ocl/register.hpp

@@ -41,6 +41,7 @@
 #include "intel_gpu/primitives/lstm_dynamic_input.hpp"
 #include "intel_gpu/primitives/lstm_dynamic_timeloop.hpp"
 #include "intel_gpu/primitives/mutable_data.hpp"
+#include "intel_gpu/primitives/multinomial.hpp"
 #include "intel_gpu/primitives/mvn.hpp"
 #include "intel_gpu/primitives/non_max_suppression.hpp"
 #include "intel_gpu/primitives/normalize.hpp"
@@ -125,6 +126,7 @@ REGISTER_OCL(lrn);
 REGISTER_OCL(lstm_gemm);
 REGISTER_OCL(lstm_elt);
 REGISTER_OCL(multiclass_nms);
+REGISTER_OCL(multinomial);
 REGISTER_OCL(mutable_data);
 REGISTER_OCL(mvn);
 REGISTER_OCL(non_max_suppression);

src/plugins/intel_gpu/src/graph/include/multinomial_inst.h

@@ -0,0 +1,26 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#pragma once
+#include "intel_gpu/primitives/multinomial.hpp"
+#include "primitive_inst.h"
+
+namespace cldnn {
+
+using multinomial_node = typed_program_node<multinomial>;
+
+template <>
+class typed_primitive_inst<multinomial> : public typed_primitive_inst_base<multinomial> {
+    using parent = typed_primitive_inst_base<multinomial>;
+    using parent::parent;
+
+public:
+    static layout calc_output_layout(multinomial_node const& node, kernel_impl_params const& impl_param);
+    static std::string to_string(multinomial_node const& node);
+
+    typed_primitive_inst(network& network, multinomial_node const& desc);
+};
+
+using multinomial_inst = typed_primitive_inst<multinomial>;
+
+}  // namespace cldnn

src/plugins/intel_gpu/src/graph/include/slice_inst.h

@@ -16,6 +16,8 @@ class typed_primitive_inst<slice> : public typed_primitive_inst_base<slice> {
     using parent::parent;
 
 public:
+    template<typename ShapeType>
+        static std::vector<layout> calc_output_layouts(const slice_node& /*node*/, const kernel_impl_params& impl_param);
     static layout calc_output_layout(slice_node const& node, kernel_impl_params const& impl_param);
     static std::string to_string(slice_node const& node);
 

src/plugins/intel_gpu/src/graph/multinomial.cpp

@@ -0,0 +1,59 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#include "intel_gpu/runtime/error_handler.hpp"
+#include "intel_gpu/runtime/memory.hpp"
+#include "json_object.h"
+#include "primitive_type_base.h"
+#include "multinomial_inst.h"
+#include "multinomial_shape_inference.hpp"
+#include <string>
+
+namespace cldnn {
+GPU_DEFINE_PRIMITIVE_TYPE_ID(multinomial)
+
+layout multinomial_inst::calc_output_layout(multinomial_node const& node, kernel_impl_params const& impl_param) {
+    auto primitive = impl_param.typed_desc<multinomial>();
+    auto input_layout = impl_param.get_input_layout(0);
+    if (input_layout.get_shape().size() == 1) {
+        return {primitive->output_data_type, input_layout.format,
+            tensor{std::vector<tensor::value_type>{
+                static_cast<tensor::value_type>(primitive->num_samples)
+            }}};
+    } else {
+        return {primitive->output_data_type, input_layout.format,
+            tensor{std::vector<tensor::value_type>{
+                input_layout.batch(),
+                static_cast<tensor::value_type>(primitive->num_samples)
+            }}};
+    }
+}
+
+multinomial_inst::typed_primitive_inst(network& network, multinomial_node const& node)
+    : parent{network, node} {}
+
+std::string multinomial_inst::to_string(multinomial_node const& node) {
+    auto desc = node.get_primitive();
+    auto node_info = node.desc_to_json();
+    auto& cdf = node.input(0);
+    auto& random_probabilities = node.input(1);
+
+    std::stringstream primitive_description;
+
+    json_composite multinomial_info;
+    multinomial_info.add("cdf", cdf.id());
+    multinomial_info.add("random probabilities", random_probabilities.id());
+    multinomial_info.add("output data_type", desc->output_data_type);
+    multinomial_info.add("with replacement", desc->with_replacement);
+    multinomial_info.add("log probs", desc->log_probs);
+    multinomial_info.add("global seed", desc->global_seed);
+    multinomial_info.add("op seed", desc->op_seed);
+    multinomial_info.add("num samples", desc->num_samples);
+
+    node_info->add("multinomial info", multinomial_info);
+    node_info->dump(primitive_description);
+
+    return primitive_description.str();
+}
+
+}  // namespace cldnn

src/plugins/intel_gpu/src/graph/slice.cpp

@@ -4,6 +4,8 @@
 
 #include <slice_inst.h>
 #include "primitive_type_base.h"
+#include "openvino/op/slice.hpp"
+#include "slice_shape_inference.hpp"
 #include <sstream>
 #include <json_object.h>
 
@@ -14,11 +16,33 @@ slice_inst::typed_primitive_inst(network& network, slice_node const& node)
     : parent(network, node) {}
 
 layout slice_inst::calc_output_layout(slice_node const& node, kernel_impl_params const& impl_param) {
-    auto primitive = impl_param.typed_desc<slice>();
-    auto input_layout = impl_param.get_input_layout();
-    return {input_layout.data_type, input_layout.format, primitive->output_shape};
+    auto l = calc_output_layouts<ov::PartialShape>(node, impl_param)[0];
+    return layout{l.data_type, l.format, tensor{l.get_ordered_dims()}};
 }
 
+template<typename ShapeType>
+inline std::vector<layout> slice_inst::calc_output_layouts(const slice_node&, const kernel_impl_params& impl_param) {
+    std::vector<ShapeType> input_shapes{impl_param.input_layouts[0].get<ShapeType>()};
+    std::unordered_map<size_t, ov::Tensor> const_data;
+    for (std::size_t i = 1; i < impl_param.input_layouts.size(); i++) {
+        auto shape_len = shape_size(impl_param.input_layouts[i].get<ShapeType>().to_shape());
+        input_shapes.push_back({static_cast<ov::Dimension::value_type>(shape_len)});
+        auto gpu_mem = impl_param.memory_deps.at(i);
+        cldnn::mem_lock<uint8_t, mem_lock_type::read> gpu_mem_lock(gpu_mem, impl_param.get_stream());
+        const_data.emplace(i, make_tensor(layout { gpu_mem->get_layout().data_type, gpu_mem->get_layout().format,
+            tensor { std::vector<tensor::value_type>{ static_cast<tensor::value_type>(shape_len) } } }, gpu_mem_lock.data()));
+    }
+    ov::op::v8::Slice op;
+    auto output_shapes = shape_infer(&op, input_shapes, ov::make_tensor_accessor(const_data));
+    auto input_layout = impl_param.get_input_layout();
+    std::vector<layout> output_layouts;
+    for (size_t i = 0; i < output_shapes.size(); ++i) {
+        output_layouts.push_back(layout({output_shapes[i], input_layout.data_type, input_layout.format}));
+    }
+    return output_layouts;
+}
+
+
 std::string slice_inst::to_string(slice_node const& node) {
     auto node_info = node.desc_to_json();
     json_composite slice_info;

src/plugins/intel_gpu/src/kernel_selector/cl_kernels/multinomial_ref.cl

@@ -0,0 +1,68 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#if WITH_REPLACEMENT
+
+KERNEL (multinomial_ref)(  __global INPUT0_TYPE* cdf
+                           , __global INPUT1_TYPE* uniform_samples
+                           , __global OUTPUT_TYPE* output
+#if HAS_FUSED_OPS_DECLS
+                           , FUSED_OPS_DECLS
+#endif
+)
+{
+    unsigned i = get_global_id(0) * SAMPLES_SIZE;
+    unsigned j = get_global_id(1);
+    if (i >= BATCH_NUM * SAMPLES_SIZE || j >= SAMPLES_SIZE)
+        return;
+    unsigned i_translated = i / SAMPLES_SIZE * CLASS_SIZE;
+    unsigned selected_class_idx = CLASS_SIZE;
+    INPUT1_TYPE sample_value = uniform_samples[i + j];
+    for (unsigned k = 0; k < CLASS_SIZE; ++k) {
+        if (sample_value <= cdf[i_translated + k]) {
+            output[i + j] = (OUTPUT_TYPE) k;
+            selected_class_idx = k;
+            break;
+        }
+    }
+}
+
+#else
+
+KERNEL (multinomial_ref)(  __global INPUT0_TYPE* cdf
+                           , __global INPUT1_TYPE* uniform_samples
+                           , __global OUTPUT_TYPE* output
+#if HAS_FUSED_OPS_DECLS
+                           , FUSED_OPS_DECLS
+#endif
+)
+{
+    unsigned i = get_global_id(0) * SAMPLES_SIZE;
+    if (i >= BATCH_NUM * SAMPLES_SIZE)
+        return;
+    for (unsigned j = 0; j < SAMPLES_SIZE; ++j) {
+        unsigned i_translated = i / SAMPLES_SIZE * CLASS_SIZE;
+        unsigned selected_class_idx = CLASS_SIZE;
+        INPUT1_TYPE sample_value = uniform_samples[i + j];
+        for (unsigned k = 0; k < CLASS_SIZE; ++k) {
+            if (sample_value <= cdf[i_translated + k]) {
+                output[i + j] = (OUTPUT_TYPE) k;
+                selected_class_idx = k;
+                break;
+            }
+        }
+        // Additional step without replacement - change probability of a given class to 0, and update the cdf
+        INPUT0_TYPE class_probability = selected_class_idx ? cdf[i_translated + selected_class_idx] -
+                                                             cdf[i_translated + selected_class_idx - 1]
+                                                           : cdf[i_translated + selected_class_idx];
+        INPUT0_TYPE divisor = 1 - class_probability;
+        for (unsigned k = 0; k < CLASS_SIZE; ++k) {
+            if (k >= selected_class_idx)
+                cdf[i_translated + k] -= class_probability;
+            cdf[i_translated + k] /= divisor;
+        }
+    }
+}
+
+#endif

src/plugins/intel_gpu/src/kernel_selector/common_types.h

@@ -94,6 +94,7 @@ enum class KernelType {
     EYE,
     GENERATE_PROPOSALS,
     MULTICLASS_NMS,
+    MULTINOMIAL,
     UNIQUE_COUNT,
     UNIQUE_GATHER,
     RMS,

src/plugins/intel_gpu/src/kernel_selector/kernels/multinomial/multinomial_kernel_ref.cpp

@@ -0,0 +1,86 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#include "multinomial_kernel_ref.h"
+#include <kernel_selector_utils.h>
+
+namespace kernel_selector {
+
+namespace {
+
+std::size_t GetBatchNum(const multinomial_params &params) {
+    return params.inputs[0].Feature().v == 1 ? 1 : params.inputs[0].Batch().v;
+}
+
+std::size_t GetClassSize(const multinomial_params &params) {
+    return params.inputs[0].Feature().v == 1 ? params.inputs[0].Batch().v : params.inputs[0].Feature().v;
+}
+
+std::size_t GetSamplesSize(const multinomial_params &params) {
+    return params.inputs[0].Feature().v == 1 ? params.outputs[0].Batch().v : params.outputs[0].Feature().v;
+}
+
+} // anonymous namespace
+
+ParamsKey MultinomialKernelRef::GetSupportedKey() const {
+    ParamsKey k;
+    k.EnableInputDataType(Datatype::F16);
+    k.EnableInputDataType(Datatype::F32);
+    k.EnableOutputDataType(Datatype::INT32);
+    k.EnableOutputDataType(Datatype::INT64);
+    k.EnableInputLayout(DataLayout::bfyx);
+    k.EnableOutputLayout(DataLayout::bfyx);
+    k.EnableBatching();
+    k.EnableDifferentTypes();
+    return k;
+}
+
+CommonDispatchData MultinomialKernelRef::SetDefault(const multinomial_params &params,
+                                                    const optional_params&) const {
+    CommonDispatchData dispatch_data {};
+    dispatch_data.gws = {GetBatchNum(params), params.with_replacement ? GetSamplesSize(params) : 1, 1};
+    dispatch_data.lws = GetOptimalLocalWorkGroupSizes(dispatch_data.gws, params.engineInfo);
+    return dispatch_data;
+}
+
+JitConstants MultinomialKernelRef::GetJitConstants(const multinomial_params &params) const {
+    JitConstants jit = MakeBaseParamsJitConstants(params);
+    jit.AddConstants(
+        {
+        MakeJitConstant("BATCH_NUM", GetBatchNum(params)),
+        MakeJitConstant("CLASS_SIZE", GetClassSize(params)),
+        MakeJitConstant("SAMPLES_SIZE", GetSamplesSize(params)),
+        MakeJitConstant("WITH_REPLACEMENT", params.with_replacement)
+        });
+    return jit;
+}
+
+KernelsData MultinomialKernelRef::GetKernelsData(const Params &params,
+                                                 const optional_params &options) const {
+    if (!Validate(params, options)) {
+        return {};
+    }
+    KernelData kernel_data = KernelData::Default<multinomial_params>(params);
+    multinomial_params &new_params = dynamic_cast<multinomial_params&>(*kernel_data.params.get());
+    auto dispatch_data = SetDefault(new_params, options);
+    auto entry_point = GetEntryPoint(kernelName, new_params.layerID, params, options);
+    auto multinomial_specific_jit = GetJitConstants(new_params);
+    auto jit = CreateJit(kernelName, multinomial_specific_jit, entry_point);
+
+    FillCLKernelData(kernel_data.kernels[0], dispatch_data, params.engineInfo,
+                     kernelName, jit, entry_point, EXE_MODE_DEFAULT, false, false, 2);
+
+    return {kernel_data};
+}
+
+bool MultinomialKernelRef::Validate(const Params &p, const optional_params &o) const {
+    if (p.GetType() != KernelType::MULTINOMIAL || o.GetType() != KernelType::MULTINOMIAL) {
+        return false;
+    }
+    const multinomial_params &params = dynamic_cast<const multinomial_params&>(p);
+    if (params.inputs.empty())
+        return false;
+    return true;
+}
+
+} // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/kernels/multinomial/multinomial_kernel_ref.h

@@ -0,0 +1,44 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#pragma once
+#include "kernel_base_opencl.h"
+#include "kernel_selector_params.h"
+
+namespace kernel_selector {
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+// GroupNormalizationParams
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+struct multinomial_params : public base_params {
+    multinomial_params() : base_params(KernelType::MULTINOMIAL) {}
+
+    ov::element::Type_t output_data_type {};
+    bool with_replacement {};
+    bool log_probs {};
+
+    ParamsKey GetParamsKey() const override {
+        return base_params::GetParamsKey();
+    }
+};
+
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+// multinomial_optional_params
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+struct multinomial_optional_params : optional_params {
+    multinomial_optional_params() : optional_params(KernelType::MULTINOMIAL) {}
+};
+
+class MultinomialKernelRef : public KernelBaseOpenCL {
+public:
+    MultinomialKernelRef() : KernelBaseOpenCL{"multinomial_ref"} {}
+    KernelsData GetKernelsData(const Params& params, const optional_params& options) const override;
+    ParamsKey GetSupportedKey() const override;
+
+private:
+    JitConstants GetJitConstants(const multinomial_params &params) const;
+    CommonDispatchData SetDefault(const multinomial_params &params,
+                                  const optional_params&) const;
+    bool Validate(const Params &p, const optional_params &o) const override;
+};
+
+} // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/kernels/multinomial/multinomial_kernel_selector.cpp

@@ -0,0 +1,18 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#include "multinomial_kernel_selector.h"
+#include "multinomial_kernel_ref.h"
+
+namespace kernel_selector {
+
+multinomial_kernel_selector::multinomial_kernel_selector() {
+    Attach<MultinomialKernelRef>();
+}
+
+KernelsData multinomial_kernel_selector::GetBestKernels(const Params &params,
+                                                                const optional_params &options) const {
+    return GetNaiveBestKernel(params, options, KernelType::MULTINOMIAL);
+}
+
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/kernel_selector/kernels/multinomial/multinomial_kernel_selector.h

@@ -0,0 +1,19 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#pragma once
+#include "kernel_selector.h"
+
+namespace kernel_selector {
+class multinomial_kernel_selector : public kernel_selector_base {
+public:
+    static multinomial_kernel_selector& Instance() {
+        static multinomial_kernel_selector instance_;
+        return instance_;
+    }
+
+    multinomial_kernel_selector();
+
+    KernelsData GetBestKernels(const Params& params, const optional_params& options) const override;
+};
+}  // namespace kernel_selector

src/plugins/intel_gpu/src/plugin/ops/multinomial.cpp

@@ -0,0 +1,152 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#include "intel_gpu/plugin/program_builder.hpp"
+#include "intel_gpu/plugin/common_utils.hpp"
+
+#include "openvino/op/constant.hpp"
+#include "openvino/op/multinomial.hpp"
+
+#include "intel_gpu/primitives/activation.hpp"
+#include "intel_gpu/primitives/broadcast.hpp"
+#include "intel_gpu/primitives/cum_sum.hpp"
+#include "intel_gpu/primitives/data.hpp"
+#include "intel_gpu/primitives/eltwise.hpp"
+#include "intel_gpu/primitives/multinomial.hpp"
+#include "intel_gpu/primitives/random_uniform.hpp"
+#include "intel_gpu/primitives/reshape.hpp"
+#include "intel_gpu/primitives/shape_of.hpp"
+#include "intel_gpu/primitives/slice.hpp"
+
+#include <memory>
+#include <cstring>
+
+namespace ov {
+namespace intel_gpu {
+namespace {
+
+template<typename T>
+cldnn::data CreateScalarDataPrimitive(ProgramBuilder& p, const cldnn::primitive_id& name, T value) {
+    auto mem = p.get_engine().allocate_memory(
+        cldnn::layout{element::from<T>(), cldnn::format::bfyx, {1, 1, 1, 1}}, false);
+    cldnn::mem_lock<int8_t> host_mem{mem, p.get_engine().get_service_stream()};
+    std::memcpy(host_mem.data(), &value, sizeof value);
+    return {name, mem};
+}
+
+cldnn::data CreateShapeDataPrimitive(ProgramBuilder& p, const cldnn::primitive_id& name, Shape& value) {
+    auto mem = p.get_engine().allocate_memory(
+        cldnn::layout{element::Type_t::i64, cldnn::format::bfyx, {1, 1, 1, static_cast<int>(value.size())}}, false);
+    cldnn::mem_lock<int8_t> host_mem{mem, p.get_engine().get_service_stream()};
+    std::vector<std::int64_t> shape {};
+    std::copy(value.begin(), value.end(), std::back_inserter(shape));
+    std::memcpy(host_mem.data(), shape.data(), sizeof(std::int64_t) * shape.size());
+    return {name, mem};
+}
+
+static void CreateMultinomialOp(ProgramBuilder& p, const std::shared_ptr<ov::op::v13::Multinomial>& op) {
+    validate_inputs_count(op, { 2 });
+    auto inputs = p.GetInputInfo(op);
+    auto input_max_shape = op->get_input_partial_shape(0).get_max_shape();
+
+    const std::int64_t last_axis = input_max_shape.size() - 1;
+    cldnn::primitive_id cumsum {layer_type_name_ID(op) + "_cumsum"};
+    if (op->get_log_probs()) {
+        cldnn::primitive_id exp {layer_type_name_ID(op) + "_exp"};
+        cldnn::activation exp_prim {exp, inputs[0], cldnn::activation_func::exp};
+        p.add_primitive(*op, exp_prim);
+        cldnn::cum_sum cumsum_primitive{cumsum, exp, last_axis};
+        p.add_primitive(*op, cumsum_primitive);
+    } else {
+        cldnn::cum_sum cumsum_primitive{cumsum, inputs[0], last_axis};
+        p.add_primitive(*op, cumsum_primitive);
+    }
+
+    const cldnn::primitive_id slice_start {layer_type_name_ID(op) + "_slice_start"};
+    auto slice_start_prim = CreateScalarDataPrimitive(p, slice_start, static_cast<std::int64_t>(-1));
+    p.add_primitive(*op, slice_start_prim);
+
+    const cldnn::primitive_id slice_stop {layer_type_name_ID(op) + "_slice_stop"};
+    auto slice_stop_prim = CreateScalarDataPrimitive(p, slice_stop, static_cast<std::int64_t>(-2));
+    p.add_primitive(*op, slice_stop_prim);
+
+    const cldnn::primitive_id slice_step {layer_type_name_ID(op) + "_slice_step"};
+    auto slice_step_prim = CreateScalarDataPrimitive(p, slice_step, static_cast<std::int64_t>(-1));
+    p.add_primitive(*op, slice_step_prim);
+
+    const cldnn::primitive_id slice_axis {layer_type_name_ID(op) + "_slice_axis"};
+    auto slice_axis_prim = CreateScalarDataPrimitive(p, slice_axis, last_axis);
+    p.add_primitive(*op, slice_axis_prim);
+
+    const cldnn::primitive_id slice {layer_type_name_ID(op) + "_slice"};
+    cldnn::slice slice_prim {
+        slice,
+        { cumsum,
+            slice_start,
+            slice_stop,
+            slice_step,
+            slice_axis
+        }
+    };
+    p.add_primitive(*op, slice_prim);
+
+    const cldnn::primitive_id cdf {layer_type_name_ID(op) + "_cdf"};
+    cldnn::eltwise divide_prim{cdf, cumsum, slice, cldnn::eltwise_mode::div};
+    p.add_primitive(*op, divide_prim);
+
+    const cldnn::primitive_id input0_shape_of {layer_type_name_ID(op) + "_input0_shape_of"};
+    cldnn::shape_of input0_shape_of_prim{input0_shape_of, inputs[0], cldnn::data_types::i32};
+    p.add_primitive(*op, input0_shape_of_prim);
+
+    const cldnn::primitive_id cdf_reshaped {layer_type_name_ID(op) + "_cdf_reshaped"};
+    cldnn::reshape cdf_reshaped_prim {cdf_reshaped, cdf, input0_shape_of, false, op->get_input_partial_shape(0)};
+    p.add_primitive(*op, cdf_reshaped_prim);
+
+    const cldnn::primitive_id random {layer_type_name_ID(op) + "_random"};
+    const cldnn::primitive_id random_shape {layer_type_name_ID(op) + "_random_shape"};
+    auto random_shape_prim = CreateShapeDataPrimitive(p, random_shape, input_max_shape);
+    p.add_primitive(*op, random_shape_prim);
+    const cldnn::primitive_id random_minval {layer_type_name_ID(op) + "_random_minval"};
+    auto random_minval_prim = CreateScalarDataPrimitive(p, random_minval, 0.f);
+    p.add_primitive(*op, random_minval_prim);
+    const cldnn::primitive_id random_maxval {layer_type_name_ID(op) + "_random_maxval"};
+    auto random_maxval_prim = CreateScalarDataPrimitive(p, random_maxval, 1.f);
+    p.add_primitive(*op, random_maxval_prim);
+    cldnn::random_uniform random_prim {
+        random,
+        {random_shape, random_minval, random_maxval},
+        op->get_input_element_type(0),
+        op->get_global_seed(),
+        op->get_op_seed(),
+        input_max_shape
+    };
+    p.add_primitive(*op, random_prim);
+
+    auto const_num_samples = std::dynamic_pointer_cast<ov::op::v0::Constant>(op->get_input_node_shared_ptr(1));
+    OPENVINO_ASSERT(const_num_samples != nullptr, "[GPU] Unsupported num_samples node type in ", op->get_friendly_name(), " (", op->get_type_name(), ")");
+
+    std::int64_t num_samples{};
+    if (const_num_samples->get_output_element_type(0) == ov::element::Type_t::i32)
+        num_samples = const_num_samples->cast_vector<std::int32_t>(1)[0];
+    else
+        num_samples = const_num_samples->cast_vector<std::int64_t>(1)[0];
+
+    cldnn::multinomial multinomial_prim {
+        layer_type_name_ID(op),
+        cdf_reshaped,
+        random,
+        op->get_convert_type(),
+        op->get_with_replacement(),
+        op->get_log_probs(),
+        op->get_global_seed(),
+        op->get_op_seed(),
+        num_samples
+    };
+    p.add_primitive(*op, multinomial_prim);
+}
+} // namespace
+
+REGISTER_FACTORY_IMPL(v13, Multinomial);
+
+}  // namespace intel_gpu
+}  // namespace ov

src/plugins/intel_gpu/src/plugin/ops/slice.cpp

@@ -19,10 +19,7 @@ namespace {
 static void CreateSliceOp(ProgramBuilder& p, const std::shared_ptr<ov::op::v8::Slice>& op) {
     validate_inputs_count(op, { 4, 5 });
     auto inputs = p.GetInputInfo(op);
-    auto output_shape = tensor_from_dims(op->get_output_shape(0));
-    auto slice_prim = cldnn::slice(layer_type_name_ID(op),
-                                   inputs,
-                                   output_shape);
+    cldnn::slice slice_prim {layer_type_name_ID(op), inputs};
     p.add_primitive(*op, slice_prim);
 }
 

src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/multinomial.cpp

@@ -0,0 +1,49 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#include "single_layer_tests/multinomial.hpp"
+
+using namespace ov::test::subgraph;
+
+namespace {
+
+const std::vector<ov::test::ElementType> netPrecisions = {
+    ov::element::f16,
+    ov::element::f32,
+};
+
+const std::vector<ov::Shape> inputShapes = {
+    {1, 32},
+    {2, 28},
+    {32},
+};
+
+const std::vector<int64_t> numSamples = {
+    2, 4,
+};
+
+const std::vector<bool> withReplacement = {
+    false,
+    true
+};
+
+const std::vector<bool> logProbes = {
+    false,
+    true
+};
+
+INSTANTIATE_TEST_SUITE_P(
+    smoke_Multinomial,
+    MultinomialTest,
+    testing::Combine(testing::ValuesIn(netPrecisions),
+                     ::testing::Values(ov::element::undefined),
+                     ::testing::Values(ov::element::undefined),
+                     testing::ValuesIn(ov::test::static_shapes_to_test_representation(inputShapes)),
+                     testing::ValuesIn(numSamples),
+                     testing::Values(ov::element::i64),
+                     testing::ValuesIn(withReplacement),
+                     testing::ValuesIn(logProbes),
+                     testing::Values(ov::test::utils::DEVICE_GPU),
+                     testing::Values(ov::AnyMap())),
+                     MultinomialTest::getTestCaseName);
+} // anonymous namespace

src/plugins/intel_gpu/tests/unit/test_cases/slice.cpp

@@ -43,7 +43,7 @@ public:
             topology.add(data("axes", axes_));
             inputs.push_back(input_info("axes"));
         }
-        topology.add(slice("slice", inputs, tensor{output_shape_}));
+        topology.add(slice("slice", inputs));
 
         cldnn::network::ptr network = get_network(engine_, topology, get_test_default_config(engine_), get_test_stream_ptr(), is_caching_test);
 
@@ -112,11 +112,11 @@ TYPED_TEST(SliceTest, bfyx_positive_step) {
 TYPED_TEST(SliceTest, bfyx_negative_step) {
     this->input_shape_ = { 1, 2, 100, 12 };
     this->start_ = this->engine_.allocate_memory({ data_types::i64, format::bfyx, { 4, 1, 1, 1 } });
-    set_values<int64_t>(this->start_, { 1, 2, 5, 100 });
+    set_values<int64_t>(this->start_, { 0, 1, 5, 100 });
     this->stop_ = this->engine_.allocate_memory({ data_types::i64, format::bfyx, { 4, 1, 1, 1 } });
-    set_values<int64_t>(this->stop_, {0, 1, 0, 1});
+    set_values<int64_t>(this->stop_, {1, 0, 0, 1});
     this->step_ = this->engine_.allocate_memory({ data_types::i64, format::bfyx, { 4, 1, 1, 1 } });
-    set_values<int64_t>(this->step_, { -1, -1, -1, -10 });
+    set_values<int64_t>(this->step_, { 1, -1, -1, -10 });
     this->output_shape_ = { 1, 1, 5, 10 };
     this->expected_output_ = {
             1799, 1789, 1779, 1769, 1759, 1699, 1689, 1679, 1669, 1659,
@@ -167,11 +167,11 @@ TYPED_TEST(SliceTest, bfyx_positive_step_cached) {
 TYPED_TEST(SliceTest, bfyx_negative_step_cached) {
     this->input_shape_ = { 1, 2, 100, 12 };
     this->start_ = this->engine_.allocate_memory({ data_types::i64, format::bfyx, { 4, 1, 1, 1 } });
-    set_values<int64_t>(this->start_, { 1, 2, 5, 100 });
+    set_values<int64_t>(this->start_, { 0, 1, 5, 100 });
     this->stop_ = this->engine_.allocate_memory({ data_types::i64, format::bfyx, { 4, 1, 1, 1 } });
-    set_values<int64_t>(this->stop_, {0, 1, 0, 1});
+    set_values<int64_t>(this->stop_, {1, 0, 0, 1});
     this->step_ = this->engine_.allocate_memory({ data_types::i64, format::bfyx, { 4, 1, 1, 1 } });
-    set_values<int64_t>(this->step_, { -1, -1, -1, -10 });
+    set_values<int64_t>(this->step_, { 1, -1, -1, -10 });
     this->output_shape_ = { 1, 1, 5, 10 };
     this->expected_output_ = {
             1799, 1789, 1779, 1769, 1759, 1699, 1689, 1679, 1669, 1659,

src/plugins/template/tests/functional/op_reference/multinomial.cpp

@@ -157,7 +157,7 @@ std::vector<MultinomialParams> generateMultinomialParams() {
 }
 
 std::vector<MultinomialParams> generateMultinomialParams() {
-    std::vector<std::vector<MultinomialParams>> combo_params{generateMultinomialParams<ov::element::f32>()};
+    std::vector<std::vector<MultinomialParams>> combo_params{generateMultinomialParams<ov::element::f64>()};
     std::vector<MultinomialParams> test_params;
     for (auto& params : combo_params)
         std::move(params.begin(), params.end(), std::back_inserter(test_params));

src/tests/functional/plugin/shared/include/single_layer_tests/multinomial.hpp

@@ -0,0 +1,21 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#pragma once
+#include "shared_test_classes/single_layer/multinomial.hpp"
+
+namespace ov {
+namespace test {
+namespace subgraph {
+
+TEST_P(MultinomialTest, CompareWithRefs) {
+    run();
+}
+
+TEST_P(MultinomialTest, CompareQueryModel) {
+    query_model();
+}
+
+}  // namespace subgraph
+}  // namespace test
+}  // namespace ov

src/tests/functional/shared_test_classes/include/shared_test_classes/single_layer/multinomial.hpp

@@ -0,0 +1,37 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#pragma once
+#include "ov_models/builders.hpp"
+#include "common_test_utils/common_utils.hpp"
+#include "shared_test_classes/base/ov_subgraph.hpp"
+
+namespace ov {
+namespace test {
+namespace subgraph {
+
+using MultinomialTestParams = std::tuple<
+    ElementType,        // netPrecision
+    ElementType,        // inPrecision
+    ElementType,        // outPrecision
+    InputShape,         // Dynamic shape + Target static shapes
+    std::int64_t,       // Number of samples
+    element::Type_t,    // Output type attribute
+    bool,               // With replacement,
+    bool,               // Log probs;
+    TargetDevice,       // targetDevice
+    Config              // config
+    >;
+
+class MultinomialTest : public testing::WithParamInterface<MultinomialTestParams>,
+                               virtual public ov::test::SubgraphBaseTest {
+public:
+    static std::string getTestCaseName(const testing::TestParamInfo<MultinomialTestParams> &obj);
+
+protected:
+    void SetUp() override;
+};
+
+} // namespace subgraph
+} // namespace test
+} // namespace ov

src/tests/functional/shared_test_classes/src/base/utils/generate_inputs.cpp

@@ -1115,6 +1115,7 @@ InputsMap getInputMap() {
 #include "openvino/opsets/opset10_tbl.hpp"
 #include "openvino/opsets/opset11_tbl.hpp"
 #include "openvino/opsets/opset12_tbl.hpp"
+#include "openvino/opsets/opset13_tbl.hpp"
 
 #include "ov_ops/opset_private_tbl.hpp"
 #undef _OPENVINO_OP_REG

src/tests/functional/shared_test_classes/src/single_layer/multinomial.cpp

@@ -0,0 +1,76 @@
+// Copyright (C) 2023 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+#include <ie_plugin_config.hpp>
+#include "shared_test_classes/single_layer/multinomial.hpp"
+#include "ov_models/builders.hpp"
+#include "ov_models/utils/ov_helpers.hpp"
+
+namespace ov {
+namespace test {
+namespace subgraph {
+
+std::string MultinomialTest::getTestCaseName(const testing::TestParamInfo<MultinomialTestParams> &obj) {
+    ElementType netType, inType, outType;
+    InputShape shape;
+    std::int64_t numSamples;
+    element::Type_t outputType;
+    bool withReplacement;
+    bool logProbs;
+    TargetDevice targetDevice;
+    Config config;
+    std::tie(netType, inType, outType, shape, numSamples, outputType, withReplacement, logProbs, targetDevice, config) = obj.param;
+
+    std::ostringstream result;
+    result << "NetType=" << netType << "_";
+    result << "InType=" << inType << "_";
+    result << "OutType=" << outType << "_";
+    result << "IS=" << ov::test::utils::partialShape2str({shape.first}) << "_";
+    result << "TS=";
+    for (const auto& item : shape.second) {
+        result << ov::test::utils::vec2str(item) << "_";
+    }
+    result << "NumSamples=" << numSamples << "_";
+    result << "OutputType=" << outputType << "_";
+    result << "WithReplacement=" << withReplacement << "_";
+    result << "LogProbs=" << logProbs << "_";
+    result << "Device=" << targetDevice;
+
+    return result.str();
+}
+
+void MultinomialTest::SetUp() {
+    InputShape shape;
+    ElementType ngPrc;
+    std::int64_t numSamples;
+    element::Type_t outputType;
+    bool withReplacement;
+    bool logProbs;
+
+    std::tie(ngPrc, inType, outType, shape, numSamples, outputType, withReplacement, logProbs, targetDevice, configuration) = this->GetParam();
+    init_input_shapes({shape});
+    ov::ParameterVector params;
+    for (auto&& shape : inputDynamicShapes) {
+        params.push_back(std::make_shared<ov::op::v0::Parameter>(ngPrc, shape));
+    }
+
+    auto numSamplesConstant = std::make_shared<ngraph::opset3::Constant>(
+        ngraph::element::Type_t::i64, ov::Shape{1}, numSamples);
+    const auto paramOuts =
+        ngraph::helpers::convert2OutputVector(ngraph::helpers::castOps2Nodes<ngraph::op::Parameter>(params));
+
+    const auto multinomial = std::make_shared<ov::op::v13::Multinomial>(
+        paramOuts.at(0),
+        numSamplesConstant,
+        outputType,
+        withReplacement,
+        logProbs,
+        0,
+        2);
+
+    function = std::make_shared<ngraph::Function>(multinomial, params, "Multinomial");
+}
+
+} // namespace subgraph
+} // namespace test
+} // namespace ov