[GPU] Enable blocked formats in Gather primitive (#11486)
* gather blocked format * enable double blocked * 5d test * support cross dimension * Add some disabled test for later use * Support non-default planar formats
This commit is contained in:
Felix Dohyun Kim authored and 
GitHub committed
1 parent
c1f4cc04de
commit
1c9cd04f96
5 files changed
+448
-45
No files matched your search
@@ -21,17 +21,41 @@ layout gather_inst::calc_output_layout(gather_node const& node) {
|
||||
auto input_layout = node.input(0).get_output_layout();
|
||||
std::vector<tensor::value_type> dims_converted(desc->output_shape.begin(), desc->output_shape.end());
|
||||
// extend shape to 4d
|
||||
for (size_t i = dims_converted.size(); i < 4; i++) {
|
||||
for (size_t i = dims_converted.size(); i < 4; i++)
|
||||
dims_converted.push_back(1);
|
||||
}
|
||||
auto output_format = format::get_default_format(dims_converted.size());
|
||||
|
||||
format output_format = input_layout.format;
|
||||
if (dims_converted.size() == 5) {
|
||||
switch (input_layout.format) {
|
||||
case format::bfyx:
|
||||
output_format = format::get_default_format(dims_converted.size());
|
||||
break;
|
||||
case format::b_fs_yx_fsv16:
|
||||
output_format = format::b_fs_zyx_fsv16;
|
||||
break;
|
||||
case format::b_fs_yx_fsv32:
|
||||
output_format = format::b_fs_zyx_fsv32;
|
||||
break;
|
||||
case format::bs_fs_yx_bsv16_fsv16:
|
||||
output_format = format::bs_fs_zyx_bsv16_fsv16;
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
} else if (dims_converted.size() == 6) {
|
||||
switch (input_layout.format) {
|
||||
case format::bfyx:
|
||||
case format::bfzyx:
|
||||
output_format = format::get_default_format(dims_converted.size());
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
}
|
||||
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, output_format, tensor(output_format, dims_converted)};
|
||||
return layout{output_type,
|
||||
output_format,
|
||||
tensor(format::get_default_format(dims_converted.size()), dims_converted)};
|
||||
}
|
||||
|
||||
std::string gather_inst::to_string(gather_node const& node) {
|
||||
|
||||
@@ -98,6 +98,24 @@ namespace detail {
|
||||
|
||||
attach_gather_impl::attach_gather_impl() {
|
||||
implementation_map<gather>::add(impl_types::ocl, gather_impl::create, {
|
||||
std::make_tuple(data_types::f32, format::fyxb),
|
||||
std::make_tuple(data_types::f16, format::fyxb),
|
||||
std::make_tuple(data_types::i32, format::fyxb),
|
||||
std::make_tuple(data_types::i8, format::fyxb),
|
||||
std::make_tuple(data_types::u8, format::fyxb),
|
||||
|
||||
std::make_tuple(data_types::f32, format::yxfb),
|
||||
std::make_tuple(data_types::f16, format::yxfb),
|
||||
std::make_tuple(data_types::i32, format::yxfb),
|
||||
std::make_tuple(data_types::i8, format::yxfb),
|
||||
std::make_tuple(data_types::u8, format::yxfb),
|
||||
|
||||
std::make_tuple(data_types::f32, format::byxf),
|
||||
std::make_tuple(data_types::f16, format::byxf),
|
||||
std::make_tuple(data_types::i32, format::byxf),
|
||||
std::make_tuple(data_types::i8, format::byxf),
|
||||
std::make_tuple(data_types::u8, format::byxf),
|
||||
|
||||
std::make_tuple(data_types::f32, format::bfyx),
|
||||
std::make_tuple(data_types::f16, format::bfyx),
|
||||
std::make_tuple(data_types::i32, format::bfyx),
|
||||
@@ -115,6 +133,84 @@ attach_gather_impl::attach_gather_impl() {
|
||||
std::make_tuple(data_types::i32, format::bfwzyx),
|
||||
std::make_tuple(data_types::i8, format::bfwzyx),
|
||||
std::make_tuple(data_types::u8, 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),
|
||||
|
||||
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),
|
||||
|
||||
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),
|
||||
|
||||
std::make_tuple(data_types::f32, format::fs_b_yx_fsv32),
|
||||
std::make_tuple(data_types::f16, format::fs_b_yx_fsv32),
|
||||
std::make_tuple(data_types::i32, format::fs_b_yx_fsv32),
|
||||
std::make_tuple(data_types::i8, format::fs_b_yx_fsv32),
|
||||
std::make_tuple(data_types::u8, format::fs_b_yx_fsv32),
|
||||
});
|
||||
}
|
||||
|
||||
|
||||
+50
-25
@@ -40,17 +40,16 @@ ParamsKey GatherKernelRef::GetSupportedKey() const {
|
||||
k.EnableInputDataType(Datatype::INT32);
|
||||
k.EnableInputDataType(Datatype::UINT8);
|
||||
k.EnableInputDataType(Datatype::INT8);
|
||||
|
||||
k.EnableOutputDataType(Datatype::F16);
|
||||
k.EnableOutputDataType(Datatype::F32);
|
||||
k.EnableOutputDataType(Datatype::INT32);
|
||||
k.EnableOutputDataType(Datatype::INT8);
|
||||
k.EnableOutputDataType(Datatype::UINT8);
|
||||
k.EnableInputLayout(DataLayout::bfyx);
|
||||
k.EnableOutputLayout(DataLayout::bfyx);
|
||||
k.EnableInputLayout(DataLayout::bfzyx);
|
||||
k.EnableOutputLayout(DataLayout::bfzyx);
|
||||
k.EnableInputLayout(DataLayout::bfwzyx);
|
||||
k.EnableOutputLayout(DataLayout::bfwzyx);
|
||||
|
||||
k.EnableAllInputLayout();
|
||||
k.EnableAllOutputLayout();
|
||||
|
||||
k.EnableTensorOffset();
|
||||
k.EnableTensorPitches();
|
||||
k.EnableBatching();
|
||||
@@ -62,8 +61,17 @@ 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)
|
||||
auto shape_raw = data_tensor.LogicalDims();
|
||||
auto shape = shape_raw;
|
||||
int shape_idx = 0;
|
||||
for (int i = 0; i < static_cast<int>(Tensor::DataChannelName::COUNT); i++) {
|
||||
int shape_raw_idx =
|
||||
data_tensor.Channelndex(data_tensor.GetLayout(), static_cast<Tensor::DataChannelName>(i));
|
||||
if (shape_raw_idx >= 0)
|
||||
shape[shape_idx++] = shape_raw[shape_raw_idx];
|
||||
}
|
||||
for (auto& i : shape) {
|
||||
if (i == 1)
|
||||
one_size_dims++;
|
||||
else
|
||||
break;
|
||||
@@ -82,10 +90,24 @@ static int64_t GetGatherBatchDim(const gather_params& params) {
|
||||
}
|
||||
|
||||
static inline std::string GetGatherMaxIndexDim(const gather_params& params) {
|
||||
return std::to_string(params.inputs[0].GetDims().at(params.inputs[0].GetDims().size() - GetGatherChannelIndex(params) - 1).v);
|
||||
switch (params.axis) {
|
||||
case GatherAxis::BATCH:
|
||||
return std::to_string(params.inputs[0].Batch().v);
|
||||
case GatherAxis::FEATURE:
|
||||
return std::to_string(params.inputs[0].Feature().v);
|
||||
case GatherAxis::W:
|
||||
return std::to_string(params.inputs[0].W().v);
|
||||
case GatherAxis::Z:
|
||||
return std::to_string(params.inputs[0].Z().v);
|
||||
case GatherAxis::Y:
|
||||
return std::to_string(params.inputs[0].Y().v);
|
||||
case GatherAxis::X:
|
||||
return std::to_string(params.inputs[0].X().v);
|
||||
}
|
||||
throw "Error";
|
||||
}
|
||||
|
||||
static inline std::string GetOrderString(std::vector<std::string>& order) {
|
||||
static inline std::string GetOrderString(const std::vector<std::string>& order) {
|
||||
std::string order_str = order[0];
|
||||
for (size_t i = 1; i < order.size(); i++)
|
||||
order_str += ", " + order[i];
|
||||
@@ -102,15 +124,13 @@ static inline std::vector<std::string> GetOrder(size_t size) {
|
||||
} else if (size == 6) {
|
||||
idx_order = {"b", "f", "w", "z", "y", "x"};
|
||||
}
|
||||
|
||||
return idx_order;
|
||||
}
|
||||
|
||||
static std::string GetDictionaryIndexOrder(const gather_params& params, size_t axis) {
|
||||
std::vector<std::string> idx_order = GetOrder(params.outputs[0].GetDims().size());
|
||||
|
||||
const std::string input_axis_index_macro = "INPUT_AXIS_INDEX";
|
||||
const std::string zeroVal = "0";
|
||||
auto idx_order = GetOrder(params.outputs[0].GetDims().size());
|
||||
auto input_axis_index_macro = "INPUT_AXIS_INDEX";
|
||||
auto zero_val = "0";
|
||||
|
||||
size_t dictionary_dims_num = GetNonEmptyDimsNumber(params.inputs[0]);
|
||||
size_t indices_dims_num = GetNonEmptyDimsNumber(params.outputs[0]) - dictionary_dims_num + 1;
|
||||
@@ -120,7 +140,7 @@ static std::string GetDictionaryIndexOrder(const gather_params& params, size_t a
|
||||
idx_order[i] = idx_order[i + indices_dims_num - 1];
|
||||
|
||||
for (size_t i = dictionary_dims_num; i < idx_order.size(); i++)
|
||||
idx_order[i] = zeroVal;
|
||||
idx_order[i] = zero_val;
|
||||
|
||||
// Fix size to inputs[0] dims size
|
||||
for (size_t i = 0; i < params.outputs[0].GetDims().size() - params.inputs[0].GetDims().size(); i++)
|
||||
@@ -131,10 +151,9 @@ static std::string GetDictionaryIndexOrder(const gather_params& params, size_t a
|
||||
return GetOrderString(idx_order);
|
||||
}
|
||||
|
||||
static std::string GetIndecesIdxOrder(const gather_params& params, size_t axis, int64_t batch_dim) {
|
||||
static std::string GetIndicesIdxOrder(const gather_params& params, size_t axis, int64_t batch_dim) {
|
||||
std::vector<std::string> idx_order = GetOrder(params.outputs[0].GetDims().size());
|
||||
|
||||
const std::string zero_val = "0";
|
||||
auto zero_val = "0";
|
||||
|
||||
size_t indices_dims_num = GetNonEmptyDimsNumber(params.inputs[1]);
|
||||
|
||||
@@ -159,24 +178,30 @@ CommonDispatchData GatherKernelRef::SetDefault(const gather_params& params, cons
|
||||
auto out_layout = params.outputs[0].GetLayout();
|
||||
std::vector<std::vector<Tensor::DataChannelName>> dims_by_gws;
|
||||
|
||||
if (out_layout == DataLayout::bfyx) {
|
||||
int rank = params.outputs[0].Dimentions();
|
||||
if (rank == 4) {
|
||||
dispatchData.gws = {output.X().v, output.Y().v, output.Feature().v * output.Batch().v};
|
||||
dims_by_gws = {{Tensor::DataChannelName::X},
|
||||
{Tensor::DataChannelName::Y},
|
||||
{Tensor::DataChannelName::FEATURE, Tensor::DataChannelName::BATCH}};
|
||||
} else if (out_layout == DataLayout::bfzyx) {
|
||||
} else if (rank == 5) {
|
||||
dispatchData.gws = {output.X().v, output.Y().v * output.Z().v, output.Feature().v * output.Batch().v};
|
||||
dims_by_gws = {{Tensor::DataChannelName::X},
|
||||
{Tensor::DataChannelName::Y, Tensor::DataChannelName::Z},
|
||||
{Tensor::DataChannelName::FEATURE, Tensor::DataChannelName::BATCH}};
|
||||
} else {
|
||||
dispatchData.gws = {output.X().v * output.Y().v, output.Z().v * output.W().v, output.Feature().v * output.Batch().v};
|
||||
} else if (rank == 6) {
|
||||
dispatchData.gws = {output.X().v * output.Y().v,
|
||||
output.Z().v * output.W().v,
|
||||
output.Feature().v * output.Batch().v};
|
||||
dims_by_gws = {{Tensor::DataChannelName::X, Tensor::DataChannelName::Y},
|
||||
{Tensor::DataChannelName::Z, Tensor::DataChannelName::W},
|
||||
{Tensor::DataChannelName::FEATURE, Tensor::DataChannelName::BATCH}};
|
||||
} else {
|
||||
IE_THROW() << "Unknown rank: rank=" << rank;
|
||||
}
|
||||
|
||||
dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo, in_layout, out_layout, dims_by_gws);
|
||||
dispatchData.lws =
|
||||
GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo, in_layout, out_layout, dims_by_gws);
|
||||
|
||||
return dispatchData;
|
||||
}
|
||||
@@ -185,7 +210,7 @@ JitConstants GatherKernelRef::GetJitConstants(const gather_params& params) const
|
||||
JitConstants jit = MakeBaseParamsJitConstants(params);
|
||||
|
||||
jit.AddConstant(MakeJitConstant("DICTIONARY_INDEX_ORDER", GetDictionaryIndexOrder(params, GetGatherChannelIndex(params))));
|
||||
jit.AddConstant(MakeJitConstant("INDICES_INDEX_ORDER", GetIndecesIdxOrder(params, GetGatherChannelIndex(params), GetGatherBatchDim(params))));
|
||||
jit.AddConstant(MakeJitConstant("INDICES_INDEX_ORDER", GetIndicesIdxOrder(params, GetGatherChannelIndex(params), GetGatherBatchDim(params))));
|
||||
if (params.support_neg_ind)
|
||||
jit.AddConstant(MakeJitConstant("INDEX_DIM", GetGatherMaxIndexDim(params)));
|
||||
|
||||
|
||||
@@ -30,27 +30,23 @@ KERNEL(gather_ref)(const __global INPUT0_TYPE* dictionary,
|
||||
#endif
|
||||
)
|
||||
{
|
||||
const uint b = (uint)get_global_id(2) / OUTPUT_FEATURE_NUM;
|
||||
const uint f = (uint)get_global_id(2) % OUTPUT_FEATURE_NUM;
|
||||
#if OUTPUT_DIMS == 6
|
||||
#define ORDER b,f,w,z,y,x
|
||||
const uint x = (uint)get_global_id(0) % OUTPUT_SIZE_X;
|
||||
const uint y = (uint)get_global_id(0) / OUTPUT_SIZE_X;
|
||||
const uint z = (uint)get_global_id(1) % OUTPUT_SIZE_Z;
|
||||
const uint w = (uint)get_global_id(1) / OUTPUT_SIZE_Z;
|
||||
const uint f = (uint)get_global_id(2) % OUTPUT_FEATURE_NUM;
|
||||
const uint b = (uint)get_global_id(2) / OUTPUT_FEATURE_NUM;
|
||||
const uint z = (uint)get_global_id(1) % OUTPUT_SIZE_Z;
|
||||
const uint y = (uint)get_global_id(0) / OUTPUT_SIZE_X;
|
||||
const uint x = (uint)get_global_id(0) % OUTPUT_SIZE_X;
|
||||
#elif OUTPUT_DIMS == 5
|
||||
#define ORDER b,f,z,y,x
|
||||
const uint x = (uint)get_global_id(0);
|
||||
const uint y = (uint)get_global_id(1) % OUTPUT_SIZE_Y;
|
||||
const uint z = (uint)get_global_id(1) / OUTPUT_SIZE_Y;
|
||||
const uint f = (uint)get_global_id(2) % OUTPUT_FEATURE_NUM;
|
||||
const uint b = (uint)get_global_id(2) / OUTPUT_FEATURE_NUM;
|
||||
const uint y = (uint)get_global_id(1) % OUTPUT_SIZE_Y;
|
||||
const uint x = (uint)get_global_id(0);
|
||||
#elif OUTPUT_DIMS == 4
|
||||
#define ORDER b,f,y,x
|
||||
const uint x = (uint)get_global_id(0);
|
||||
const uint y = (uint)get_global_id(1);
|
||||
const uint f = (uint)get_global_id(2) % OUTPUT_FEATURE_NUM;
|
||||
const uint b = (uint)get_global_id(2) / OUTPUT_FEATURE_NUM;
|
||||
const uint x = (uint)get_global_id(0);
|
||||
#endif
|
||||
|
||||
const uint indices_idx = GET_INDICES_INDEX(INDICES_INDEX_ORDER);
|
||||
|
||||
@@ -8,10 +8,272 @@
|
||||
#include <intel_gpu/primitives/gather.hpp>
|
||||
|
||||
#include <cstddef>
|
||||
#include <array>
|
||||
|
||||
using namespace cldnn;
|
||||
using namespace ::tests;
|
||||
|
||||
std::ostream& operator<<(std::ostream& os, FLOAT16 x) {
|
||||
return os << float(x);
|
||||
}
|
||||
|
||||
template <class T>
|
||||
int get_not_one_dim(const T& a) {
|
||||
int ret = a.size();
|
||||
while (ret - 1 >= 0 && a[ret - 1] == 1)
|
||||
ret--;
|
||||
return ret;
|
||||
};
|
||||
|
||||
template <class T>
|
||||
int get_linear_size(const T& a) {
|
||||
return std::accumulate(a.begin(), a.end(), 1, [](int x, int y) {
|
||||
return x * y;
|
||||
});
|
||||
};
|
||||
|
||||
using gather8_test_param = std::tuple<int, // batch_dim, value in [0,get_not_one_dim(dict))
|
||||
int, // axis, value in [batch_dim,get_not_one_dim(dict))
|
||||
format::type, // format of input0
|
||||
format::type, // format of input1
|
||||
std::vector<int>, // shape of input0, order = default_format
|
||||
std::vector<int>>; // shape of input1, order = default_format
|
||||
template <class T_dat, class T_ind, data_types T_dat_dt, data_types T_ind_dt>
|
||||
class gather8_test : public ::testing::TestWithParam<gather8_test_param> {
|
||||
public:
|
||||
int axis, batch_dim;
|
||||
std::array<format::type, 3> fmt;
|
||||
std::vector<int> shape_in[2];
|
||||
std::vector<int> shape_out;
|
||||
|
||||
void SetUp() override {
|
||||
auto& engine = get_test_engine();
|
||||
|
||||
std::tie(batch_dim, axis, fmt[0], fmt[1], shape_in[0], shape_in[1]) = GetParam();
|
||||
fmt[2] = fmt[0];
|
||||
|
||||
// refer: src/core/shape_inference/include/gather_shape_inference.hpp
|
||||
size_t out_dim = get_not_one_dim(shape_in[0]) - 1 + get_not_one_dim(shape_in[1]) - batch_dim;
|
||||
shape_out = std::vector<int>(std::max(shape_in[0].size(), out_dim), 1);
|
||||
for (int i = 0; i < batch_dim; i++) // batch_dim
|
||||
shape_out[i] = shape_in[0][i];
|
||||
for (int i = batch_dim; i < axis; i++) // before axis = shape_in[0][..]
|
||||
shape_out[i] = shape_in[0][i];
|
||||
for (int i = batch_dim; i < get_not_one_dim(shape_in[1]); i++) // axis = shape_in[1]
|
||||
shape_out[axis + (i - batch_dim)] = shape_in[1][i];
|
||||
for (int i = axis + 1; i < get_not_one_dim(shape_in[0]); i++) // after axis = shape_in[0][..]
|
||||
shape_out[axis + get_not_one_dim(shape_in[1]) - batch_dim + (i - axis - 1)] = shape_in[0][i];
|
||||
|
||||
auto dat = generate_random_1d<T_dat>(get_linear_size(shape_in[0]), -99, 99);
|
||||
auto input0 =
|
||||
engine.allocate_memory(layout(T_dat_dt,
|
||||
format::get_default_format(shape_in[0].size()),
|
||||
tensor(format::get_default_format(shape_in[0].size()), shape_in[0])));
|
||||
set_values(input0, dat);
|
||||
|
||||
auto ind =
|
||||
generate_random_1d<T_ind>(get_linear_size(shape_in[1]), -shape_in[0][axis], shape_in[0][axis] - 1, 1);
|
||||
auto input1 =
|
||||
engine.allocate_memory(layout(T_ind_dt,
|
||||
format::get_default_format(shape_in[1].size()),
|
||||
tensor(format::get_default_format(shape_in[1].size()), shape_in[1])));
|
||||
set_values(input1, ind);
|
||||
|
||||
topology reorder_topo;
|
||||
reorder_topo.add(input_layout("input0", input0->get_layout()));
|
||||
reorder_topo.add(input_layout("input1", input1->get_layout()));
|
||||
reorder_topo.add(reorder("reorder0", "input0", fmt[0], T_dat_dt));
|
||||
reorder_topo.add(reorder("reorder1", "input1", fmt[1], T_ind_dt));
|
||||
reorder_topo.add(gather("gather",
|
||||
"reorder0",
|
||||
"reorder1",
|
||||
axis,
|
||||
ov::Shape(shape_out.begin(), shape_out.end()),
|
||||
batch_dim,
|
||||
true));
|
||||
reorder_topo.add(reorder("reorder2", "gather", format::type::bfwzyx, T_dat_dt));
|
||||
network reorder_network(engine, reorder_topo);
|
||||
reorder_network.set_input_data("input0", input0);
|
||||
reorder_network.set_input_data("input1", input1);
|
||||
auto reorder_output = reorder_network.execute().at("reorder2").get_memory();
|
||||
cldnn::mem_lock<T_dat> reorder_output_ptr(reorder_output, get_test_stream());
|
||||
|
||||
topology planar_topo;
|
||||
planar_topo.add(input_layout("input0", input0->get_layout()));
|
||||
planar_topo.add(input_layout("input1", input1->get_layout()));
|
||||
planar_topo.add(
|
||||
gather("gather", "input0", "input1", axis, ov::Shape(shape_out.begin(), shape_out.end()), batch_dim, true));
|
||||
network planar_network(engine, planar_topo);
|
||||
planar_network.set_input_data("input0", input0);
|
||||
planar_network.set_input_data("input1", input1);
|
||||
auto planar_output = planar_network.execute().at("gather").get_memory();
|
||||
cldnn::mem_lock<T_dat> planar_output_ptr(planar_output, get_test_stream());
|
||||
|
||||
EXPECT_TRUE(
|
||||
!memcmp(reorder_output_ptr.data(), planar_output_ptr.data(), get_linear_size(shape_out) * sizeof(T_dat)));
|
||||
}
|
||||
};
|
||||
using gather8_test_f16i32 = gather8_test<FLOAT16, int, data_types::f16, data_types::i32>;
|
||||
using gather8_test_f32i8 = gather8_test<float, char, data_types::f32, data_types::i8>;
|
||||
using gather8_test_i32i32 = gather8_test<int, int, data_types::i32, data_types::i32>;
|
||||
TEST_P(gather8_test_f16i32, gather8_test_f16i32) {}
|
||||
TEST_P(gather8_test_f32i8, gather8_test_f32i8) {}
|
||||
TEST_P(gather8_test_i32i32, gather8_test_i32i32) {}
|
||||
|
||||
// Important testcases
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_bd0_d4_i1,
|
||||
gather8_test_f16i32,
|
||||
testing::Combine(testing::Values(0), // bdim in [0,get_not_one_dim(dict))
|
||||
testing::Values(0), // axis in [batch_dim,get_not_one_dim(dict))
|
||||
testing::Values(format::type::bfyx),
|
||||
testing::Values(format::type::bfyx),
|
||||
testing::Values(std::vector<int>{5, 44, 7, 8}),
|
||||
testing::Values(std::vector<int>{4, 1, 1, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_bd0_d2_i2,
|
||||
gather8_test_f16i32,
|
||||
testing::Combine(testing::Values(0),
|
||||
testing::Values(1),
|
||||
testing::Values(format::type::b_fs_yx_fsv4),
|
||||
testing::Values(format::type::b_fs_yx_fsv16),
|
||||
testing::Values(std::vector<int>{8, 67, 1, 1}),
|
||||
testing::Values(std::vector<int>{4, 56, 1, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_fs_b_yx_fsv32,
|
||||
gather8_test_f16i32,
|
||||
testing::Combine(testing::Values(0),
|
||||
testing::Values(0, 1, 2),
|
||||
testing::Values(format::type::fs_b_yx_fsv32),
|
||||
testing::Values(format::type::fs_b_yx_fsv32),
|
||||
testing::Values(std::vector<int>{3, 77, 4, 1}),
|
||||
testing::Values(std::vector<int>{2, 66, 1, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_fs_b_yx_fsv32_bd1,
|
||||
gather8_test_f16i32,
|
||||
testing::Combine(testing::Values(1),
|
||||
testing::Values(1, 2),
|
||||
testing::Values(format::type::fs_b_yx_fsv32),
|
||||
testing::Values(format::type::fs_b_yx_fsv32),
|
||||
testing::Values(std::vector<int>{3, 77, 44, 1}),
|
||||
testing::Values(std::vector<int>{3, 66, 55, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_fs_b_yx_fsv32_bd2,
|
||||
gather8_test_f16i32,
|
||||
testing::Combine(testing::Values(2),
|
||||
testing::Values(2, 3),
|
||||
testing::Values(format::type::fs_b_yx_fsv32),
|
||||
testing::Values(format::type::fs_b_yx_fsv32),
|
||||
testing::Values(std::vector<int>{3, 4, 44, 6}),
|
||||
testing::Values(std::vector<int>{3, 4, 5, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_bs_fs_yx_bsv16_fsv16_bd0_dim4_to_dim5,
|
||||
gather8_test_f16i32,
|
||||
testing::Combine(testing::Values(0),
|
||||
testing::Values(0, 2),
|
||||
testing::Values(format::type::bs_fs_yx_bsv16_fsv16),
|
||||
testing::Values(format::type::b_fs_yx_fsv32),
|
||||
testing::Values(std::vector<int>{3, 77, 44, 1}),
|
||||
testing::Values(std::vector<int>{3, 66, 55, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_b_fs_yx_fsv16_bd0_dim4_to_dim5,
|
||||
gather8_test_f16i32,
|
||||
testing::Combine(testing::Values(0),
|
||||
testing::Values(0, 2),
|
||||
testing::Values(format::type::b_fs_yx_fsv16),
|
||||
testing::Values(format::type::b_fs_yx_fsv4),
|
||||
testing::Values(std::vector<int>{3, 77, 44, 1}),
|
||||
testing::Values(std::vector<int>{3, 66, 55, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_bfyx_bd0_dim4_to_dim6,
|
||||
gather8_test_f16i32,
|
||||
testing::Combine(testing::Values(0),
|
||||
testing::Values(0, 2),
|
||||
testing::Values(format::type::bfyx),
|
||||
testing::Values(format::type::b_fs_yx_fsv4),
|
||||
testing::Values(std::vector<int>{3, 7, 4, 6}),
|
||||
testing::Values(std::vector<int>{3, 6, 5, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_bd0_d4_i1,
|
||||
gather8_test_f32i8,
|
||||
testing::Combine(testing::Values(0),
|
||||
testing::Values(2),
|
||||
testing::Values(format::type::bs_fs_yx_bsv16_fsv16),
|
||||
testing::Values(format::type::bs_fs_yx_bsv32_fsv16),
|
||||
testing::Values(std::vector<int>{5, 44, 7, 8}),
|
||||
testing::Values(std::vector<int>{4, 1, 1, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_bd0_d3_i3,
|
||||
gather8_test_f32i8,
|
||||
testing::Combine(testing::Values(0),
|
||||
testing::Values(1),
|
||||
testing::Values(format::type::b_fs_zyx_fsv16),
|
||||
testing::Values(format::type::bfzyx),
|
||||
testing::Values(std::vector<int>{8, 67, 3, 1, 1}),
|
||||
testing::Values(std::vector<int>{3, 56, 9, 1, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_b_fs_zyx_fsv32,
|
||||
gather8_test_f32i8,
|
||||
testing::Combine(testing::Values(1),
|
||||
testing::Values(2),
|
||||
testing::Values(format::type::b_fs_zyx_fsv32),
|
||||
testing::Values(format::type::b_fs_zyx_fsv16),
|
||||
testing::Values(std::vector<int>{8, 66, 3, 1, 1}),
|
||||
testing::Values(std::vector<int>{8, 56, 9, 1, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_b_fs_yx_fsv4,
|
||||
gather8_test_i32i32,
|
||||
testing::Combine(testing::Values(0),
|
||||
testing::Values(2),
|
||||
testing::Values(format::type::b_fs_yx_fsv4),
|
||||
testing::Values(format::type::b_fs_yx_fsv4),
|
||||
testing::Values(std::vector<int>{4, 6, 2, 1}),
|
||||
testing::Values(std::vector<int>{3, 5, 1, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_bfyx,
|
||||
gather8_test_i32i32,
|
||||
testing::Combine(testing::Values(0),
|
||||
testing::Values(1),
|
||||
testing::Values(format::type::bfyx),
|
||||
testing::Values(format::type::bfyx),
|
||||
testing::Values(std::vector<int>{4, 3, 2, 1}),
|
||||
testing::Values(std::vector<int>{5, 6, 1, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_byxf,
|
||||
gather8_test_i32i32,
|
||||
testing::Combine(testing::Values(0),
|
||||
testing::Values(2),
|
||||
testing::Values(format::type::byxf),
|
||||
testing::Values(format::type::byxf),
|
||||
testing::Values(std::vector<int>{4, 6, 2, 1}),
|
||||
testing::Values(std::vector<int>{3, 5, 1, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(gather8_fs_b_yx_fsv32_2,
|
||||
gather8_test_i32i32,
|
||||
testing::Combine(testing::Values(0),
|
||||
testing::Values(2),
|
||||
testing::Values(format::type::fs_b_yx_fsv32),
|
||||
testing::Values(format::type::fs_b_yx_fsv32),
|
||||
testing::Values(std::vector<int>{4, 6, 2, 3}),
|
||||
testing::Values(std::vector<int>{3, 1, 1, 1})));
|
||||
|
||||
// Disabled tests for faster CI
|
||||
// Remove DISABLED_ prefix to test these cases
|
||||
#define FORMAT4D \
|
||||
format::type::bfyx, format::type::b_fs_yx_fsv16, format::type::bs_fs_yx_bsv16_fsv16, format::type::byxf, \
|
||||
format::type::yxfb, format::type::fyxb, format::type::fs_b_yx_fsv32
|
||||
#define FORMAT5D format::type::bfzyx, format::type::b_fs_zyx_fsv16, format::type::bs_fs_zyx_bsv16_fsv16
|
||||
#define FORMAT6D format::type::bfwzyx
|
||||
INSTANTIATE_TEST_SUITE_P(DISABLED_gather8_4d_f16i32,
|
||||
gather8_test_f16i32,
|
||||
testing::Combine(testing::Values(0, 1),
|
||||
testing::Values(1, 2),
|
||||
testing::Values(FORMAT4D),
|
||||
testing::Values(FORMAT4D),
|
||||
testing::Values(std::vector<int>{3, 6, 2, 1}),
|
||||
testing::Values(std::vector<int>{3, 5, 1, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(DISABLED_gather8_5d_f16i32,
|
||||
gather8_test_f16i32,
|
||||
testing::Combine(testing::Values(0, 1),
|
||||
testing::Values(1, 2, 3),
|
||||
testing::Values(FORMAT5D),
|
||||
testing::Values(FORMAT5D),
|
||||
testing::Values(std::vector<int>{3, 6, 2, 7, 1}),
|
||||
testing::Values(std::vector<int>{3, 5, 1, 1, 1})));
|
||||
INSTANTIATE_TEST_SUITE_P(DISABLED_gather8_6d_f16i32,
|
||||
gather8_test_f16i32,
|
||||
testing::Combine(testing::Values(0, 1),
|
||||
testing::Values(1, 2, 3),
|
||||
testing::Values(FORMAT6D),
|
||||
testing::Values(FORMAT6D),
|
||||
testing::Values(std::vector<int>{3, 6, 2, 7, 1, 1}),
|
||||
testing::Values(std::vector<int>{3, 5, 4, 1, 1, 1})));
|
||||
|
||||
TEST(gather8_gpu_fp16, d323_axisY_bdim_m1) {
|
||||
// Dictionary : 3x2x3x4x2
|
||||
// Indexes : 3x2x3x1
|
||||
|
||||
Reference in new issue
Block a user