diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_nd_update.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_nd_update.cpp index a83977ba37e..9b14746dde7 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_nd_update.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/scatter_nd_update.cpp @@ -39,102 +39,54 @@ struct scatter_nd_update_impl : typed_primitive_impl_ocl { return {params, optional_params}; } + + void update_dispatch_data(const kernel_impl_params& impl_param) override { + auto kernel_params = get_kernel_params(impl_param); + (_kernel_data.update_dispatch_data_func)(kernel_params.first, _kernel_data); + } }; namespace detail { attach_scatter_nd_update_impl::attach_scatter_nd_update_impl() { - implementation_map::add(impl_types::ocl, typed_primitive_impl_ocl::create, { - std::make_tuple(data_types::f32, format::bfyx), - std::make_tuple(data_types::f16, format::bfyx), - std::make_tuple(data_types::i32, format::bfyx), - std::make_tuple(data_types::i8, format::bfyx), - std::make_tuple(data_types::u8, format::bfyx), + auto types = { data_types::f32, data_types::f16, data_types::i32, data_types::i8, data_types::u8 }; + auto formats = { + format::bfyx, + format::b_fs_yx_fsv4, + format::b_fs_yx_fsv16, + format::b_fs_yx_fsv32, + format::bs_fs_yx_bsv4_fsv2, + format::bs_fs_yx_bsv4_fsv4, + format::bs_fs_yx_bsv8_fsv2, + format::bs_fs_yx_bsv8_fsv4, + format::bs_fs_yx_bsv16_fsv16, + format::bs_fs_yx_bsv32_fsv16, + format::bs_fs_yx_bsv32_fsv32, - std::make_tuple(data_types::f32, format::bfzyx), - std::make_tuple(data_types::f16, format::bfzyx), - std::make_tuple(data_types::i32, format::bfzyx), - std::make_tuple(data_types::i8, format::bfzyx), - std::make_tuple(data_types::u8, format::bfzyx), + format::bfzyx, + format::b_fs_zyx_fsv16, + format::b_fs_zyx_fsv32, - std::make_tuple(data_types::f32, format::bfwzyx), - std::make_tuple(data_types::f16, format::bfwzyx), - std::make_tuple(data_types::i32, format::bfwzyx), - std::make_tuple(data_types::i8, format::bfwzyx), - std::make_tuple(data_types::u8, format::bfwzyx), + format::bfwzyx + }; - std::make_tuple(data_types::f32, format::b_fs_yx_fsv4), - std::make_tuple(data_types::f16, format::b_fs_yx_fsv4), - std::make_tuple(data_types::i32, format::b_fs_yx_fsv4), - std::make_tuple(data_types::i8, format::b_fs_yx_fsv4), - std::make_tuple(data_types::u8, format::b_fs_yx_fsv4), + implementation_map::add(impl_types::ocl, + shape_types::static_shape, + typed_primitive_impl_ocl::create, + types, + formats); - std::make_tuple(data_types::f32, format::b_fs_yx_fsv16), - std::make_tuple(data_types::f16, format::b_fs_yx_fsv16), - std::make_tuple(data_types::i32, format::b_fs_yx_fsv16), - std::make_tuple(data_types::i8, format::b_fs_yx_fsv16), - std::make_tuple(data_types::u8, format::b_fs_yx_fsv16), + auto dyn_formats = { + format::bfyx, + format::bfzyx, + format::bfwzyx + }; - std::make_tuple(data_types::f32, format::b_fs_yx_fsv32), - std::make_tuple(data_types::f16, format::b_fs_yx_fsv32), - std::make_tuple(data_types::i32, format::b_fs_yx_fsv32), - std::make_tuple(data_types::i8, format::b_fs_yx_fsv32), - std::make_tuple(data_types::u8, format::b_fs_yx_fsv32), - - std::make_tuple(data_types::f32, format::b_fs_zyx_fsv16), - std::make_tuple(data_types::f16, format::b_fs_zyx_fsv16), - std::make_tuple(data_types::i32, format::b_fs_zyx_fsv16), - std::make_tuple(data_types::i8, format::b_fs_zyx_fsv16), - std::make_tuple(data_types::u8, format::b_fs_zyx_fsv16), - - std::make_tuple(data_types::f32, format::b_fs_zyx_fsv32), - std::make_tuple(data_types::f16, format::b_fs_zyx_fsv32), - std::make_tuple(data_types::i32, format::b_fs_zyx_fsv32), - std::make_tuple(data_types::i8, format::b_fs_zyx_fsv32), - std::make_tuple(data_types::u8, format::b_fs_zyx_fsv32), - - std::make_tuple(data_types::f32, format::bs_fs_yx_bsv4_fsv2), - std::make_tuple(data_types::f16, format::bs_fs_yx_bsv4_fsv2), - std::make_tuple(data_types::i32, format::bs_fs_yx_bsv4_fsv2), - std::make_tuple(data_types::i8, format::bs_fs_yx_bsv4_fsv2), - std::make_tuple(data_types::u8, format::bs_fs_yx_bsv4_fsv2), - - std::make_tuple(data_types::f32, format::bs_fs_yx_bsv4_fsv4), - std::make_tuple(data_types::f16, format::bs_fs_yx_bsv4_fsv4), - std::make_tuple(data_types::i32, format::bs_fs_yx_bsv4_fsv4), - std::make_tuple(data_types::i8, format::bs_fs_yx_bsv4_fsv4), - std::make_tuple(data_types::u8, format::bs_fs_yx_bsv4_fsv4), - - std::make_tuple(data_types::f32, format::bs_fs_yx_bsv8_fsv2), - std::make_tuple(data_types::f16, format::bs_fs_yx_bsv8_fsv2), - std::make_tuple(data_types::i32, format::bs_fs_yx_bsv8_fsv2), - std::make_tuple(data_types::i8, format::bs_fs_yx_bsv8_fsv2), - std::make_tuple(data_types::u8, format::bs_fs_yx_bsv8_fsv2), - - std::make_tuple(data_types::f32, format::bs_fs_yx_bsv8_fsv4), - std::make_tuple(data_types::f16, format::bs_fs_yx_bsv8_fsv4), - std::make_tuple(data_types::i32, format::bs_fs_yx_bsv8_fsv4), - std::make_tuple(data_types::i8, format::bs_fs_yx_bsv8_fsv4), - std::make_tuple(data_types::u8, format::bs_fs_yx_bsv8_fsv4), - - std::make_tuple(data_types::f32, format::bs_fs_yx_bsv16_fsv16), - std::make_tuple(data_types::f16, format::bs_fs_yx_bsv16_fsv16), - std::make_tuple(data_types::i32, format::bs_fs_yx_bsv16_fsv16), - std::make_tuple(data_types::i8, format::bs_fs_yx_bsv16_fsv16), - std::make_tuple(data_types::u8, format::bs_fs_yx_bsv16_fsv16), - - std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv16), - std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv16), - std::make_tuple(data_types::i32, format::bs_fs_yx_bsv32_fsv16), - std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv16), - std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv16), - - std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv32), - std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv32), - std::make_tuple(data_types::i32, format::bs_fs_yx_bsv32_fsv32), - std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv32), - std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv32), - }); + implementation_map::add(impl_types::ocl, + shape_types::dynamic_shape, + typed_primitive_impl_ocl::create, + types, + dyn_formats); } } // namespace detail diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/scatter_nd_update_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/scatter_nd_update_ref.cl index 747a8412dff..4ec9b665760 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/scatter_nd_update_ref.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/scatter_nd_update_ref.cl @@ -34,7 +34,8 @@ #define INDICES_MAX_DIM 6 -KERNEL(scatter_nd_update_ref)(const __global INPUT0_TYPE* data, +KERNEL(scatter_nd_update_ref)(OPTIONAL_SHAPE_INFO_ARG + const __global INPUT0_TYPE* data, const __global INPUT1_TYPE* indices, const __global INPUT2_TYPE* updates, __global OUTPUT_TYPE* output @@ -43,7 +44,6 @@ KERNEL(scatter_nd_update_ref)(const __global INPUT0_TYPE* data, #endif ) { - const uint dim0 = get_global_id(0); const uint dim1 = get_global_id(1); const uint dim2 = get_global_id(2); diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/scatter_update/scatter_nd_update_kernel_ref.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/scatter_update/scatter_nd_update_kernel_ref.cpp index 89415db7a33..d11630c978e 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/scatter_update/scatter_nd_update_kernel_ref.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/scatter_update/scatter_nd_update_kernel_ref.cpp @@ -27,6 +27,7 @@ ParamsKey ScatterNDUpdateKernelRef::GetSupportedKey() const { k.EnableTensorPitches(); k.EnableBatching(); k.EnableDifferentTypes(); + k.EnableDynamicShapesSupport(); return k; } @@ -44,7 +45,7 @@ static inline std::vector GetDefaultOrder(size_t size) { } ScatterNDUpdateKernelRef::DispatchData -ScatterNDUpdateKernelRef::SetDefault(const scatter_nd_update_params& params, const optional_params&, bool is_second) const { +ScatterNDUpdateKernelRef::SetDefault(const scatter_nd_update_params& params, bool is_second) const { DispatchData dispatchData; if (!is_second) { @@ -92,17 +93,20 @@ bool ScatterNDUpdateKernelRef::Validate(const Params& p, const optional_params& } const scatter_nd_update_params& params = static_cast(p); - auto input_dims = params.inputs[0].LogicalDims(); - auto indices_dims = params.inputs[1].LogicalDims(); - std::reverse(indices_dims.begin(), indices_dims.end()); auto indices_rank = params.indices_rank; if (indices_rank < 1) { return false; } - if (indices_dims[indices_rank - 1] > input_dims.size()) { - return false; + if (!params.has_dynamic_inputs()) { + auto input_dims = params.inputs[0].LogicalDims(); + auto indices_dims = params.inputs[1].LogicalDims(); + std::reverse(indices_dims.begin(), indices_dims.end()); + + if (indices_dims[indices_rank - 1] > input_dims.size()) { + return false; + } } for (auto& fused_op : params.fused_ops) { @@ -113,26 +117,37 @@ bool ScatterNDUpdateKernelRef::Validate(const Params& p, const optional_params& return true; } -static std::string GetInputBlockND(const scatter_nd_update_params& params, int num, const int rank) { +static std::string GetInputBlockND(const scatter_nd_update_params& params, size_t num, size_t rank) { const auto& input = params.inputs[num]; + auto input_dims = input.LogicalDims(); std::reverse(input_dims.begin(), input_dims.end()); + std::vector block_nd(rank + 1); block_nd[rank] = 1; - for (int idx = (rank - 1); idx >= 0; idx--) { + + std::vector block_nd_s(rank + 1); + block_nd_s[rank] = "1"; + size_t input_offset = num * 6; + + for (int32_t idx = (rank - 1); idx >= 0; --idx) { block_nd[idx] = input_dims[idx] * block_nd[idx + 1]; + + size_t dim_offset = idx < 2 ? idx : idx + 6 - rank; + block_nd_s[idx] = "(" + toCodeString(input.GetDims()[input.GetDims().size() - idx - 1], input_offset + dim_offset) + "*" + block_nd_s[idx + 1] + ")"; } - std::stringstream s; - for (int i = 0; i < (rank + 1); i++) { - if (i < rank) { - s << block_nd[i] << ","; - } else { - s << block_nd[i]; + std::string result; + if (input.is_dynamic()) { + for (auto& block : block_nd_s) { + result += block + ","; + } + } else { + for (size_t block : block_nd) { + result += toCodeString(block) + ","; } } - auto str_result = s.str(); - return str_result; + return result; } KernelsData ScatterNDUpdateKernelRef::GetKernelsData(const Params& params, const optional_params& options) const { @@ -144,27 +159,50 @@ KernelsData ScatterNDUpdateKernelRef::GetKernelsData(const Params& params, const scatter_nd_update_params& newParams = *static_cast(kd.params.get()); auto cldnn_jit = GetJitConstants(newParams); + kd.update_dispatch_data_func = [this](const Params& params, KernelData& kd) { + const auto& prim_params = static_cast(params); + OPENVINO_ASSERT(kd.kernels.size() == 2, "[GPU] Invalid kernels size for update dispatch data func"); + + for (size_t i = 0; i < 2; ++i) { + auto dispatchData = SetDefault(prim_params, (i == 1)); + kd.kernels[i].params.workGroups.global = dispatchData.gws; + kd.kernels[i].params.workGroups.local = dispatchData.lws; + } + }; + // First iter - copy input data to output data // Second iter - update values specified by updates at specific index position specified by indices for (int i = 0; i < 2; i++) { - auto dispatchData = SetDefault(newParams, options, (i == 1)); + auto dispatchData = SetDefault(newParams, (i == 1)); auto entry_point = GetEntryPoint(kernelName, newParams.layerID, params, options, i); if (i == 1) { - int input0_rank = static_cast(newParams.inputs[0].LogicalDims().size()); - int input2_rank = static_cast(newParams.inputs[2].LogicalDims().size()); + size_t input0_rank = newParams.inputs[0].LogicalDims().size(); + size_t input2_rank = newParams.inputs[2].LogicalDims().size(); cldnn_jit.AddConstant(MakeJitConstant("IS_SECOND_ITER", "true")); - cldnn_jit.AddConstant(MakeJitConstant("INDICES_LAST_DIM", dispatchData.indicesLastDim)); cldnn_jit.AddConstant(MakeJitConstant("INPUT0_BLOCK_ND", GetInputBlockND(newParams, 0, input0_rank))); cldnn_jit.AddConstant(MakeJitConstant("INPUT1_BLOCK_ND", GetInputBlockND(newParams, 1, newParams.indices_rank - 1))); cldnn_jit.AddConstant(MakeJitConstant("INPUT2_BLOCK_ND", GetInputBlockND(newParams, 2, input2_rank))); cldnn_jit.AddConstant(MakeJitConstant("INDICES_RANK", newParams.indices_rank)); + + const auto& ind_input = newParams.inputs[1]; + if (ind_input.is_dynamic()) { + size_t last_idx = newParams.indices_rank - 1; + size_t dim_offset = last_idx < 2 ? last_idx : 5; + size_t input_idx = last_idx < 2 ? ind_input.GetDims().size() - last_idx : 0; + + auto indices_last_dim = toCodeString(ind_input.GetDims()[input_idx], 6 + dim_offset); + cldnn_jit.AddConstant(MakeJitConstant("INDICES_LAST_DIM", indices_last_dim)); + } else { + cldnn_jit.AddConstant(MakeJitConstant("INDICES_LAST_DIM", dispatchData.indicesLastDim)); + } } std::pair jit = CreateJit(kernelName, cldnn_jit, entry_point); clKernelData& kernel = kd.kernels[i]; - FillCLKernelData(kernel, dispatchData, params.engineInfo, kernelName, jit, entry_point, "", false, false, 3, GetFusedPrimitiveInputsCount(params)); + FillCLKernelData(kernel, dispatchData, params.engineInfo, kernelName, jit, entry_point, + "", false, false, 3, GetFusedPrimitiveInputsCount(params), 1, newParams.has_dynamic_tensors()); } return {kd}; diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/scatter_update/scatter_nd_update_kernel_ref.h b/src/plugins/intel_gpu/src/kernel_selector/kernels/scatter_update/scatter_nd_update_kernel_ref.h index 4e4df91602c..a351320a34f 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/scatter_update/scatter_nd_update_kernel_ref.h +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/scatter_update/scatter_nd_update_kernel_ref.h @@ -32,7 +32,7 @@ public: ScatterNDUpdateKernelRef() : KernelBaseOpenCL("scatter_nd_update_ref") {} virtual ~ScatterNDUpdateKernelRef() {} virtual JitConstants GetJitConstants(const scatter_nd_update_params& params) const; - virtual DispatchData SetDefault(const scatter_nd_update_params& params, const optional_params&, bool is_second) const; + virtual DispatchData SetDefault(const scatter_nd_update_params& params, bool is_second) const; KernelsData GetKernelsData(const Params& params, const optional_params& options) const override; ParamsKey GetSupportedKey() const override; std::vector GetSupportedFusedOps() const override { diff --git a/src/plugins/intel_gpu/tests/test_cases/scatter_nd_update_gpu_test.cpp b/src/plugins/intel_gpu/tests/test_cases/scatter_nd_update_gpu_test.cpp index 356fcbb453d..f45da0ddd30 100644 --- a/src/plugins/intel_gpu/tests/test_cases/scatter_nd_update_gpu_test.cpp +++ b/src/plugins/intel_gpu/tests/test_cases/scatter_nd_update_gpu_test.cpp @@ -4,6 +4,7 @@ #include "test_utils.h" #include "ngraph/runtime/reference/scatter_nd_update.hpp" +#include "scatter_nd_update_inst.h" #include #include @@ -4383,3 +4384,74 @@ TEST(scatter_nd_update_gpu_fp16, d222222_i211111) { ASSERT_EQ(expected_results[i], half_to_float(output_ptr[i])); } } + +TEST(scatter_nd_update_gpu, dynamic) { + // Dictionary : 2x1x2x8 + // Indexes : 2x3 + // Updates : 2x8 + // Output : 2x1x2x8 + // Input values in fp32 + // + auto& engine = get_test_engine(); + + auto input1_layout = layout{ ov::PartialShape::dynamic(4), data_types::f32, format::bfyx }; + auto input2_layout = layout{ ov::PartialShape::dynamic(2), data_types::f32, format::bfyx }; + auto input3_layout = layout{ ov::PartialShape::dynamic(2), data_types::f32, format::bfyx }; + + auto input1 = engine.allocate_memory({ { 2, 1, 2, 8 }, data_types::f32, format::bfyx }); // Dictionary + auto input2 = engine.allocate_memory({ { 2, 3 }, data_types::f32, format::bfyx }); // Indexes + auto input3 = engine.allocate_memory({ { 2, 8 }, data_types::f32, format::bfyx }); // Updates + + set_values(input1, { + 0.f, 1.f, 2.f, 3.f, 4.f, 5.f, 6.f, 7.f, + 8.f, 9.f, 10.f, 11.f, 12.f, 13.f, 14.f, 15.f, + 16.f, 17.f, 18.f, 19.f, 20.f, 21.f, 22.f, 23.f, + 24.f, 25.f, 26.f, 27.f, 28.f, 29.f, 30.f, 31.f + }); + + set_values(input2, { + 0.f, 1.f, 1.f, 2.f, 2.f, 2.f + }); + + set_values(input3, { + 24.f, 24.f, 24.f, 24.f, 24.f, 24.f, 24.f, 24.f, + 42.f, 42.f, 42.f, 42.f, 42.f, 42.f, 42.f, 42.f + }); + + topology topology; + topology.add(input_layout("InputData", input1_layout)); + topology.add(input_layout("InputIndices", input2_layout)); + topology.add(input_layout("InputUpdates", input3_layout)); + topology.add( + scatter_nd_update("scatter_nd_update", input_info("InputData"), input_info("InputIndices"), input_info("InputUpdates"), 2) + ); + + ExecutionConfig config; + config.set_property(ov::intel_gpu::allow_new_shape_infer(true)); + network network(engine, topology, config); + + network.set_input_data("InputData", input1); + network.set_input_data("InputIndices", input2); + network.set_input_data("InputUpdates", input3); + + auto inst = network.get_primitive("scatter_nd_update"); + auto impl = inst->get_impl(); + ASSERT_TRUE(impl != nullptr); + ASSERT_TRUE(impl->is_dynamic()); + + auto outputs = network.execute(); + + auto output = outputs.at("scatter_nd_update").get_memory(); + cldnn::mem_lock output_ptr(output, get_test_stream()); + + std::vector expected_results = { + 0.f, 1.f, 2.f, 3.f, 4.f, 5.f, 6.f, 7.f, + 24.f, 24.f, 24.f, 24.f, 24.f, 24.f, 24.f, 24.f, + 16.f, 17.f, 18.f, 19.f, 20.f, 21.f, 22.f, 23.f, + 42.f, 42.f, 42.f, 42.f, 42.f, 42.f, 42.f, 42.f + }; + + for (size_t i = 0; i < expected_results.size(); ++i) { + ASSERT_EQ(expected_results[i], output_ptr[i]); + } +}