IntenseWebs/openvino
3
0 Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

[GPU] Added ScatterNDUpdate shape agnostic kernel (#15567)

Browse Source
This commit is contained in:
Roman Lyamin authored and GitHub committed 2023-02-08 16:27:34 +01:00
1 parent 8d2eae95a7
commit 23b863ffe8
5 files changed
+172 -110
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

src/plugins/intel_gpu/src/graph/impls/ocl/scatter_nd_update.cpp
+38 -86
View File
@@ -39,102 +39,54 @@ struct scatter_nd_update_impl : typed_primitive_impl_ocl<scatter_nd_update> {
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<scatter_nd_update>::add(impl_types::ocl, typed_primitive_impl_ocl<scatter_nd_update>::create<scatter_nd_update_impl>, {
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<scatter_nd_update>::add(impl_types::ocl,
shape_types::static_shape,
typed_primitive_impl_ocl<scatter_nd_update>::create<scatter_nd_update_impl>,
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<scatter_nd_update>::add(impl_types::ocl,
shape_types::dynamic_shape,
typed_primitive_impl_ocl<scatter_nd_update>::create<scatter_nd_update_impl>,
types,
dyn_formats);
}
} // namespace detail
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/scatter_nd_update_ref.cl
+2 -2
View File
@@ -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);
src/plugins/intel_gpu/src/kernel_selector/kernels/scatter_update/scatter_nd_update_kernel_ref.cpp
+59 -21
View File
@@ -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<std::string> 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<const scatter_nd_update_params&>(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<size_t> block_nd(rank + 1);
block_nd[rank] = 1;
for (int idx = (rank - 1); idx >= 0; idx--) {
std::vector<std::string> 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<scatter_nd_update_params*>(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<const scatter_nd_update_params&>(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<int>(newParams.inputs[0].LogicalDims().size());
int input2_rank = static_cast<int>(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<std::string, std::string> 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};
src/plugins/intel_gpu/src/kernel_selector/kernels/scatter_update/scatter_nd_update_kernel_ref.h
+1 -1
View File
@@ -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<FusedOpType> GetSupportedFusedOps() const override {
src/plugins/intel_gpu/tests/test_cases/scatter_nd_update_gpu_test.cpp
+72
View File
@@ -4,6 +4,7 @@
#include "test_utils.h"
#include "ngraph/runtime/reference/scatter_nd_update.hpp"
#include "scatter_nd_update_inst.h"
#include <intel_gpu/primitives/input_layout.hpp>
#include <intel_gpu/primitives/scatter_update.hpp>
@@ -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<float> output_ptr(output, get_test_stream());
std::vector<float> 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]);
}
}