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

[IE CLDNN] Improve Gather performance and add fusing support (#736)

Browse Source
This commit is contained in:
Sergey Shlyapnikov authored and GitHub committed 2020-06-05 10:20:58 +03:00
1 parent 2100521a14
commit 6e491a89ad
6 files changed
+269 -59
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/gather/gather_kernel_ref.cpp
Vendored
+86 -44
View File
@@ -1,5 +1,5 @@
/*
// Copyright (c) 2019 Intel Corporation
// Copyright (c) 2019-2020 Intel Corporation
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
@@ -17,9 +17,10 @@
#include "gather_kernel_ref.h"
#include "kernel_selector_utils.h"
#include <string>
#include <vector>
namespace kernel_selector {
static int32_t GetGatherChannelIndex(const gather_params& params) {
static size_t GetGatherChannelIndex(const gather_params& params) {
Tensor::DataChannelName name = Tensor::DataChannelName::X;
switch (params.axis) {
@@ -46,8 +47,10 @@ ParamsKey GatherKernelRef::GetSupportedKey() const {
k.EnableOutputDataType(Datatype::F16);
k.EnableOutputDataType(Datatype::F32);
k.EnableOutputDataType(Datatype::INT32);
k.EnableAllInputLayout();
k.EnableAllOutputLayout();
k.EnableOutputDataType(Datatype::INT8);
k.EnableOutputDataType(Datatype::UINT8);
k.EnableInputLayout(DataLayout::bfyx);
k.EnableOutputLayout(DataLayout::bfyx);
k.EnableTensorOffset();
k.EnableTensorPitches();
k.EnableBatching();
@@ -55,41 +58,80 @@ ParamsKey GatherKernelRef::GetSupportedKey() const {
return k;
}
static size_t getPartSize(const gather_params& params, int32_t axis) {
size_t partSize = 1;
for (size_t i = params.inputs[0].Dimentions() - axis; i > 0; --i) partSize *= params.inputs[0].GetDims()[i - 1].v;
return partSize;
static size_t GetNonEmptyDimsNumber(const DataTensor& data_tensor) {
if (data_tensor.LogicalSize() != 1) {
// Count the number of "one size" dimensions starting with X to Batch
size_t one_size_dims = 0;
for (auto& i : data_tensor.GetDims()) {
if (i.v == 1)
one_size_dims++;
else
break;
}
return data_tensor.Dimentions() - one_size_dims;
} else {
return 1;
}
}
static size_t getNumberOfParts(const gather_params& params, size_t partSize) {
return params.inputs[0].LogicalSize() / partSize;
static inline std::string GetOrderString(std::vector<std::string>& order) {
std::string order_str = order[0];
for (size_t i = 1; i < order.size(); i++)
order_str += ", " + order[i];
return order_str;
}
static size_t getSliceSize(const gather_params& params, int32_t axis) {
size_t numberOfItemsInSlice = 1;
for (size_t i = params.inputs[0].Dimentions() - axis - 1; i > 0; --i)
numberOfItemsInSlice *= params.inputs[0].GetDims()[i - 1].v;
return numberOfItemsInSlice;
static std::string GetDictionaryIndexOrder(const gather_params& params, size_t axis) {
std::vector<std::string> default_order = { "b", "f", "y", "x" };
const std::string input_axis_index_macro = "INPUT_AXIS_INDEX";
const std::string zeroVal = "0";
size_t dictionary_dims_num = GetNonEmptyDimsNumber(params.inputs[0]);
size_t indices_dims_num = GetNonEmptyDimsNumber(params.inputs[1]);
// Shift indices of Gather dictionary input related to output dims
for (size_t i = axis + 1; i < dictionary_dims_num; i++)
default_order[i] = default_order[i + indices_dims_num - 1];
for (size_t i = dictionary_dims_num; i < default_order.size(); i++)
default_order[i] = zeroVal;
default_order[axis] = input_axis_index_macro;
return GetOrderString(default_order);
}
static std::string GetIndecesIdxOrder(const gather_params& params, size_t axis) {
std::vector<std::string> default_order = { "b", "f", "y", "x" };
const std::string zero_val = "0";
size_t indices_dims_num = GetNonEmptyDimsNumber(params.inputs[1]);
// Shift indices of Gather indices input related to output dims
for (size_t i = 0; i < indices_dims_num; i++)
default_order[i] = default_order[axis + i];
for (size_t i = indices_dims_num; i < default_order.size(); i++)
default_order[i] = zero_val;
return GetOrderString(default_order);
}
CommonDispatchData GatherKernelRef::SetDefault(const gather_params& params, const optional_params&) const {
CommonDispatchData runInfo;
const auto& output = params.output;
const int32_t axis = GetGatherChannelIndex(params);
std::vector<size_t> global = {output.Batch().v, output.Feature().v,output.X().v * output.Y().v};
std::vector<size_t> local = GetOptimalLocalWorkGroupSizes(global, params.engineInfo);
const size_t numberOfParts = params.inputs[0].LogicalSize() / getPartSize(params, axis);
size_t gws = numberOfParts * params.inputs[1].LogicalSize();
const size_t vectorSize = 16;
runInfo.gws0 = Align(gws, vectorSize);
runInfo.gws1 = 1;
runInfo.gws2 = 1;
runInfo.lws0 = vectorSize;
runInfo.lws1 = 1;
runInfo.lws2 = 1;
runInfo.gws0 = global[0];
runInfo.gws1 = global[1];
runInfo.gws2 = global[2];
runInfo.lws0 = local[0];
runInfo.lws1 = local[1];
runInfo.lws2 = local[2];
runInfo.fp16UnitUsed = params.inputs[0].GetDType() == Datatype::F16;
@@ -99,17 +141,13 @@ CommonDispatchData GatherKernelRef::SetDefault(const gather_params& params, cons
JitConstants GatherKernelRef::GetJitConstants(const gather_params& params) const {
JitConstants jit = MakeBaseParamsJitConstants(params);
int32_t axis = GetGatherChannelIndex(params);
size_t partSize = getPartSize(params, axis);
size_t sliceSize = getSliceSize(params, axis);
size_t numberOfParts = getNumberOfParts(params, partSize);
size_t numberOfIndexes = params.inputs[1].LogicalSize();
jit.AddConstant(MakeJitConstant("DICTIONARY_INDEX_ORDER", GetDictionaryIndexOrder(params, GetGatherChannelIndex(params))));
jit.AddConstant(MakeJitConstant("INDICES_INDEX_ORDER", GetIndecesIdxOrder(params, GetGatherChannelIndex(params))));
jit.AddConstant(MakeJitConstant("AXIS", axis));
jit.AddConstant(MakeJitConstant("PART_SIZE", partSize));
jit.AddConstant(MakeJitConstant("SLICE_SIZE", sliceSize));
jit.AddConstant(MakeJitConstant("PARTS_NUMBER", numberOfParts));
jit.AddConstant(MakeJitConstant("COMPUTATIONAL_OPERATIONS_NUMBER", numberOfParts * numberOfIndexes));
if (!params.fused_ops.empty()) {
FusedOpsConfiguration conf = { "", {"b", "f", "y", "x"}, "val", params.inputs[0].GetDType() };
jit.Merge(MakeFusedOpsJitConstants(params, {conf}));
}
return jit;
}
@@ -121,18 +159,22 @@ bool GatherKernelRef::Validate(const Params& p, const optional_params& o) const
const gather_params& params = static_cast<const gather_params&>(p);
if (params.inputs[0].GetDType() != params.output.GetDType())
return false;
for (auto& fused_op : params.fused_ops) {
if (!IsFusedPrimitiveSupported(fused_op))
return false;
}
return true;
}
KernelsData GatherKernelRef::GetKernelsData(const Params& params, const optional_params& options) const {
if (!Validate(params, options)) {
return {};
}
KernelData kd = KernelData::Default<gather_params>(params);
gather_params& newParams = *static_cast<gather_params*>(kd.params.get());
assert(params.GetType() == KernelType::GATHER);
auto runInfo = SetDefault(newParams, options);
auto entry_point = GetEntryPoint(kernelName, newParams.layerID, options);
auto cldnn_jit = GetJitConstants(newParams);
@@ -140,7 +182,7 @@ KernelsData GatherKernelRef::GetKernelsData(const Params& params, const optional
auto& kernel = kd.kernels[0];
FillCLKernelData(kernel, runInfo, params.engineInfo, kernelName, jit, entry_point, "", false, false, 2);
FillCLKernelData(kernel, runInfo, params.engineInfo, kernelName, jit, entry_point, "", false, false, 2, GetFusedPrimitiveInputsCount(params));
kd.estimatedTime = DONT_USE_IF_HAVE_SOMETHING_ELSE;
inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/gather/gather_kernel_ref.h
Vendored
+6 -1
View File
@@ -1,5 +1,5 @@
/*
// Copyright (c) 2019 Intel Corporation
// Copyright (c) 2019-2020 Intel Corporation
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
@@ -45,6 +45,11 @@ public:
virtual CommonDispatchData SetDefault(const gather_params& params, const optional_params&) const;
KernelsData GetKernelsData(const Params& params, const optional_params& options) const override;
ParamsKey GetSupportedKey() const override;
std::vector<FusedOpType> GetSupportedFusedOps() const override {
return { FusedOpType::QUANTIZE,
FusedOpType::SCALE,
FusedOpType::ACTIVATION };
}
protected:
bool Validate(const Params& p, const optional_params& o) const override;
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_ref.cl
Vendored
+31 -12
View File
@@ -1,4 +1,4 @@
// Copyright (c) 2019 Intel Corporation
// Copyright (c) 2019-2020 Intel Corporation
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
@@ -15,19 +15,38 @@
#include "include/include_all.cl"
KERNEL(gather_ref)(const __global INPUT0_TYPE* dictionary, const __global INPUT1_TYPE* indices, __global OUTPUT_TYPE* output)
#define INPUT_AXIS_INDEX (uint)indices[indices_idx]
#define GET_DICTIONARY_INDEX(idx_order) INPUT0_GET_INDEX(idx_order)
#define GET_INDICES_INDEX(idx_order) INPUT1_GET_INDEX(idx_order)
KERNEL(gather_ref)(const __global INPUT0_TYPE* dictionary,
const __global INPUT1_TYPE* indices,
__global OUTPUT_TYPE* output
#if HAS_FUSED_OPS_DECLS
, FUSED_OPS_DECLS
#endif
)
{
const uint workItemId = get_global_id(0);
const uint b = get_global_id(0);
const uint f = get_global_id(1);
const uint yx = get_global_id(2);
const uint y = yx / OUTPUT_SIZE_X;
const uint x = yx % OUTPUT_SIZE_X;
if (workItemId >= COMPUTATIONAL_OPERATIONS_NUMBER)
return;
const uint indices_idx = GET_INDICES_INDEX(INDICES_INDEX_ORDER);
const uint dictionary_idx = GET_DICTIONARY_INDEX(DICTIONARY_INDEX_ORDER);
const uint output_idx = OUTPUT_GET_INDEX(b, f, y, x);
uint partNumber = workItemId / INPUT1_LENGTH;
uint outputIndex = workItemId * SLICE_SIZE;
uint index = workItemId - (partNumber * INPUT1_LENGTH);
INPUT0_TYPE val = dictionary[dictionary_idx];
for (int k = 0; k < SLICE_SIZE; ++k)
{
output[outputIndex++] = ACTIVATION(dictionary[(partNumber * PART_SIZE) + ((uint) indices[index] * SLICE_SIZE) + k], ACTIVATION_PARAMS);
}
#if HAS_FUSED_OPS
FUSED_OPS;
output[output_idx] = TO_OUTPUT_TYPE(FUSED_OPS_RESULT);
#else
output[output_idx] = ACTIVATION(val, ACTIVATION_PARAMS);
#endif
}
#undef GET_INDICES_INDEX
#undef GET_DICTIONARY_INDEX
#undef INPUT_AXIS_INDEX
inference-engine/thirdparty/clDNN/src/gather.cpp
Vendored
+7 -2
View File
@@ -1,5 +1,5 @@
/*
// Copyright (c) 2019 Intel Corporation
// Copyright (c) 2019-2020 Intel Corporation
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
@@ -35,7 +35,12 @@ layout gather_inst::calc_output_layout(gather_node const& node) {
auto output_shape = desc->output_shape;
return layout{input_layout.data_type, input_format, output_shape};
auto output_type = input_layout.data_type;
if (node.has_fused_primitives()) {
output_type = node.get_fused_output_layout().data_type;
}
return layout{output_type, input_format, output_shape};
}
std::string gather_inst::to_string(gather_node const& node) {
inference-engine/thirdparty/clDNN/src/graph_optimizer/prepare_primitive_fusing.cpp
Vendored
+6
View File
@@ -371,6 +371,8 @@ void prepare_primitive_fusing::fuse_simple_primitives(program_impl &p) {
should_fuse |= input_data.is_type<lrn>();
should_fuse |= input_data.is_type<gather>();
if (!should_fuse)
return;
@@ -412,6 +414,8 @@ void prepare_primitive_fusing::fuse_simple_primitives(program_impl &p) {
should_fuse |= input_data.is_type<lrn>();
should_fuse |= input_data.is_type<gather>();
if (!should_fuse)
return;
@@ -481,6 +485,8 @@ void prepare_primitive_fusing::fuse_simple_primitives(program_impl &p) {
(input_data.get_dependency(0).get_output_layout().data_type == data_types::u8 ||
input_data.get_dependency(0).get_output_layout().data_type == data_types::i8 ||
input_data.get_output_layout().data_type == out_layout.data_type);
should_fuse |= input_data.is_type<gather>() && quantize_node.get_scale_shift_opt();
should_fuse |= input_data.is_type<permute>() &&
quantize_node.get_scale_shift_opt();
inference-engine/thirdparty/clDNN/tests/test_cases/fusings_gpu_test.cpp
Vendored
+133
View File
@@ -33,6 +33,7 @@
#include "api/mvn.hpp"
#include "api/deconvolution.hpp"
#include "api/permute.hpp"
#include "api/gather.hpp"
#include "test_utils/test_utils.h"
@@ -4059,6 +4060,138 @@ INSTANTIATE_TEST_CASE_P(DISABLED_fusings_gpu,
pooling_test_params{CASE_POOLING_I8_3, 2, 4, pooling_mode::average, "pooling_gpu_fs_bs_yx_bsv4_fsv32"},
}), );
/* ----------------------------------------------------------------------------------------------------- */
/* ------------------------------------------ Gather cases --------------------------------------------- */
/* ----------------------------------------------------------------------------------------------------- */
struct gather_test_params {
tensor dictionary_shape;
tensor indices_shape;
tensor out_shape;
cldnn::gather::gather_axis axis;
data_types data_type;
format input_format;
data_types default_type;
format default_format;
size_t expected_fused_primitives;
size_t expected_not_fused_primitives;
};
#define CASE_GATHER_FP32_1 {2, 3, 1, 4}, {4, 1, 1, 1}, {4, 3, 1, 4}, cldnn::gather::gather_axis::along_b, data_types::f32, format::bfyx, data_types::f32, format::bfyx
#define CASE_GATHER_FP32_2 {3, 2, 1, 2}, {2, 3, 1, 1}, {2, 3, 2, 2}, cldnn::gather::gather_axis::along_b, data_types::f32, format::bfyx, data_types::f32, format::bfyx
#define CASE_GATHER_FP32_3 {3, 1, 1, 2}, {2, 1, 1, 1}, {3, 2, 1, 2}, cldnn::gather::gather_axis::along_f, data_types::f32, format::bfyx, data_types::f32, format::bfyx
#define CASE_GATHER_FP32_4 {5, 3, 2, 2}, {3, 1, 1, 1}, {5, 2, 2, 3}, cldnn::gather::gather_axis::along_y, data_types::f32, format::bfyx, data_types::f32, format::bfyx
#define CASE_GATHER_FP32_5 {2, 3, 1, 2}, {1, 3, 1, 1}, {2, 3, 3, 1}, cldnn::gather::gather_axis::along_y, data_types::f32, format::bfyx, data_types::f32, format::bfyx
#define CASE_GATHER_FP16_1 {2, 3, 1, 4}, {4, 1, 1, 1}, {4, 3, 1, 4}, cldnn::gather::gather_axis::along_b, data_types::f16, format::bfyx, data_types::f16, format::bfyx
#define CASE_GATHER_FP16_2 {3, 2, 1, 2}, {2, 3, 1, 1}, {2, 3, 2, 2}, cldnn::gather::gather_axis::along_b, data_types::f16, format::bfyx, data_types::f16, format::bfyx
#define CASE_GATHER_FP16_3 {3, 1, 1, 2}, {2, 1, 1, 1}, {3, 2, 1, 2}, cldnn::gather::gather_axis::along_f, data_types::f16, format::bfyx, data_types::f16, format::bfyx
#define CASE_GATHER_FP16_4 {5, 3, 2, 2}, {3, 1, 1, 1}, {5, 2, 2, 3}, cldnn::gather::gather_axis::along_y, data_types::f16, format::bfyx, data_types::f16, format::bfyx
#define CASE_GATHER_FP16_5 {2, 3, 1, 2}, {1, 3, 1, 1}, {2, 3, 3, 1}, cldnn::gather::gather_axis::along_y, data_types::f16, format::bfyx, data_types::f16, format::bfyx
class GatherPrimitiveFusingTest : public ::BaseFusingTest<gather_test_params> {
public:
void execute(gather_test_params& p) {
auto input_prim = get_mem(get_input_layout(p));
network network_not_fused(this->engine, this->topology_non_fused, bo_not_fused);
network network_fused(this->engine, this->topology_fused, bo_fused);
network_fused.set_input_data("input", input_prim);
network_not_fused.set_input_data("input", input_prim);
compare(network_not_fused, network_fused, p);
}
layout get_input_layout(gather_test_params& p) {
return layout{ p.data_type, p.input_format, p.dictionary_shape };
}
layout get_indices_layout(gather_test_params& p) {
return layout{ p.data_type, format::bfyx, p.indices_shape };
}
size_t get_axis_dim(gather_test_params& p) {
switch (p.axis) {
case cldnn::gather::gather_axis::along_x:
return p.dictionary_shape.spatial[0];
case cldnn::gather::gather_axis::along_y:
return p.dictionary_shape.spatial[1];
case cldnn::gather::gather_axis::along_f:
return p.dictionary_shape.feature[0];
case cldnn::gather::gather_axis::along_b:
return p.dictionary_shape.batch[0];
default:
return 1;
}
}
layout get_per_channel_layout(gather_test_params& p) {
return layout{ p.default_type, p.default_format, tensor{1, p.out_shape.feature[0], 1, 1} };
}
};
class gather_quantize : public GatherPrimitiveFusingTest {};
TEST_P(gather_quantize, basic) {
auto p = GetParam();
create_topologies(input_layout("input", get_input_layout(p)),
data("gather_indices", get_mem(get_indices_layout(p), 0, static_cast<int>(get_axis_dim(p)))),
data("in_lo", get_mem(get_per_channel_layout(p), min_random, 0)),
data("in_hi", get_mem(get_per_channel_layout(p), 1, max_random)),
data("out_lo", get_mem(get_single_element_layout(p), -127)),
data("out_hi", get_mem(get_single_element_layout(p), 127)),
gather("gather_prim", "input", "gather_indices", p.axis, p.out_shape),
quantize("quantize", "gather_prim", "in_lo", "in_hi", "out_lo", "out_hi", 255, data_types::i8),
reorder("reorder_bfyx", "quantize", p.default_format, data_types::f32)
);
tolerance = 1.f;
execute(p);
}
INSTANTIATE_TEST_CASE_P(fusings_gpu, gather_quantize,
::testing::ValuesIn(std::vector<gather_test_params>{
gather_test_params{ CASE_GATHER_FP32_1, 2, 3 },
gather_test_params{ CASE_GATHER_FP32_2, 2, 3 },
gather_test_params{ CASE_GATHER_FP32_3, 2, 3 },
gather_test_params{ CASE_GATHER_FP32_4, 2, 3 },
gather_test_params{ CASE_GATHER_FP32_5, 2, 3 },
gather_test_params{ CASE_GATHER_FP16_1, 2, 3 },
gather_test_params{ CASE_GATHER_FP16_2, 2, 3 },
gather_test_params{ CASE_GATHER_FP16_3, 2, 3 },
gather_test_params{ CASE_GATHER_FP16_4, 2, 3 },
gather_test_params{ CASE_GATHER_FP16_5, 2, 3 },
}), );
class gather_scale_activation : public GatherPrimitiveFusingTest {};
TEST_P(gather_scale_activation, basic) {
auto p = GetParam();
create_topologies(input_layout("input", get_input_layout(p)),
data("gather_indices", get_mem(get_indices_layout(p), 0, static_cast<int>(get_axis_dim(p)))),
data("scale_data", get_mem(get_per_channel_layout(p), -10, 10)),
gather("gather_prim", "input", "gather_indices", p.axis, p.out_shape),
activation("activation", "gather_prim", activation_func::abs),
scale("scale", "activation", "scale_data"),
reorder("reorder_bfyx", "scale", p.default_format, data_types::f32)
);
tolerance = 1e-5f;
execute(p);
}
INSTANTIATE_TEST_CASE_P(fusings_gpu, gather_scale_activation,
::testing::ValuesIn(std::vector<gather_test_params>{
gather_test_params{ CASE_GATHER_FP32_1, 2, 4 },
gather_test_params{ CASE_GATHER_FP32_2, 2, 4 },
gather_test_params{ CASE_GATHER_FP32_3, 2, 4 },
gather_test_params{ CASE_GATHER_FP32_4, 2, 4 },
gather_test_params{ CASE_GATHER_FP32_5, 2, 4 },
gather_test_params{ CASE_GATHER_FP16_1, 2, 4 },
gather_test_params{ CASE_GATHER_FP16_2, 2, 4 },
gather_test_params{ CASE_GATHER_FP16_3, 2, 4 },
gather_test_params{ CASE_GATHER_FP16_4, 2, 4 },
gather_test_params{ CASE_GATHER_FP16_5, 2, 4 },
}), );
/* ------------------------------------------------------------------------------------------------------------ */
/* ---------------------------------------- PERMUTE FUSE cases -------------------------------------------------- */
/* ------------------------------------------------------------------------------------------------------------ */