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

[GPU] Support 7D and 8D tensors (#16810)

Browse Source
This commit is contained in:
Vladimir Paramuzov authored and GitHub committed 2023-04-13 09:04:14 +04:00
1 parent dcf6fb1e1a
commit 67c07ccebe
58 files changed
+1505 -827
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

src/plugins/intel_gpu/include/intel_gpu/primitives/fully_connected.hpp
+4 -4
View File
@@ -19,14 +19,14 @@ namespace cldnn {
/// <tr><th>Data type <th>activation format <th>weights format
/// <tr><td rowspan="7">F32 <td rowspan="4">bfyx <td>yxfb
/// <tr> <td>fyxb
/// <tr> <td>bs_xs_xsv8_bsv8
/// <tr> <td>bs_x_bsv16
/// <tr> <td>bs_fs_fsv8_bsv8
/// <tr> <td>bs_f_bsv16
/// <tr> <td rowspan="3">yxfb <td>bfyx
/// <tr> <td>yxfb
/// <tr> <td>bs_xs_xsv8_bsv8
/// <tr> <td>bs_fs_fsv8_bsv8
/// <tr><td rowspan="4">F16 <td rowspan="3">bfyx <td>yxfb
/// <tr> <td>fyxb
/// <tr> <td>bs_x_bsv16
/// <tr> <td>bs_f_bsv16
/// <tr> <td >yxfb <td>bfyx
/// </table>
src/plugins/intel_gpu/include/intel_gpu/runtime/format.hpp
+15 -5
View File
@@ -57,6 +57,9 @@ struct format_traits {
static bool is_spatial_char(char c) { return std::string(spatial_chars()).find_first_of(c) != std::string::npos; }
/// @brief Checks if @p c represents group dimensions.
static bool is_group_char(char c) { return std::string(group_chars()).find_first_of(c) != std::string::npos; }
/// @brief Checks if order has @p c dimension.
bool has_dimension(char c) const { return order.find_first_of(c) != std::string::npos; }
};
/// @brief Represents memory formats (orders).
@@ -73,6 +76,8 @@ struct format {
bfyx, ///< the most common format for activations in clDNN.
bfzyx, ///< format for 5d data tensors
bfwzyx, ///< batch, feature, 4D spatial
bfuwzyx, ///< 7d tensor
bfvuwzyx, ///< 8d tensor
yxfb, ///< batch first, feature and than spatials
byxf, ///< used in bitmaps, input from user i.e b images of RGB format
fyxb, ///< format not used inside clDNN, but supported in reorder as extension
@@ -107,10 +112,10 @@ struct format {
bs_fs_zyx_bsv32_fsv32, ///< format used for big batches (batch and features blocked by 32)
bs_fs_zyx_bsv32_fsv16, ///< format used for big batches (batch blocked by 32, features blocked by 16)
fs_b_yx_fsv32, ///< format for input for fp16 primitives
bs_xs_xsv8_bsv8, ///< format used only for fully connected
bs_xs_xsv8_bsv16, ///< format used only for fully connected
bs_x_bsv16, ///< format used only for fully connected weights fp16 batch=1 : bs - batch slice
///< (responses slice), bsv16 - 16 values of single batch slice, x - flattened plane of (fyx)
bs_fs_fsv8_bsv8, ///< format used only for fully connected
bs_fs_fsv8_bsv16, ///< format used only for fully connected
bs_f_bsv16, ///< format used only for fully connected weights fp16 batch=1 : bs - batch slice
///< (responses slice), bsv16 - 16 values of single batch slice, f - flattened plane of (fyx)
b_fs_yx_32fp, ///< format for data for binary convolutions
winograd_2x3_s1_data, ///< format used for input for winograd convolution, F(2,3) -- filter 3x3 with stride 1
nv12, ///< format for media nv12 input
@@ -326,7 +331,8 @@ struct format {
static bool is_simple_data_format(type fmt) {
return (fmt == yxfb || fmt == byxf ||
fmt == bfyx || fmt == fyxb ||
fmt == bfzyx || fmt == bfwzyx);
fmt == bfzyx || fmt == bfwzyx ||
fmt == bfuwzyx || fmt == bfvuwzyx);
}
static format get_default_format(size_t rank, bool is_weights = false, bool is_grouped = false);
@@ -393,6 +399,10 @@ struct format {
std::string to_string() const;
};
inline std::ostream& operator<<(std::ostream& os, const format& fmt) {
return os << fmt.to_string();
}
/// @}
/// @}
} // namespace cldnn
src/plugins/intel_gpu/include/intel_gpu/runtime/layout.hpp
+1 -1
View File
@@ -527,7 +527,7 @@ struct layout {
// for smaller buffer which, currently, should always be performed
bool identical(const layout& other) const;
static size_t max_rank() { return 6; }
static size_t max_rank() { return 8; }
static ov::PartialShape transform(const ov::PartialShape& pshape, cldnn::format old_fmt, cldnn::format new_fmt);
size_t hash() const {
src/plugins/intel_gpu/include/intel_gpu/runtime/tensor.hpp
+32 -70
View File
@@ -31,7 +31,7 @@ namespace cldnn {
constexpr int32_t tensor_batch_dim_max = 1;
constexpr int32_t tensor_feature_dim_max = 1;
constexpr int32_t tensor_spatial_dim_max = 4;
constexpr int32_t tensor_spatial_dim_max = 6;
constexpr int32_t tensor_group_dim_max = 1;
constexpr int32_t tensor_dim_max = tensor_batch_dim_max + tensor_feature_dim_max + tensor_spatial_dim_max + tensor_group_dim_max;
@@ -346,7 +346,7 @@ public:
delim = ",";
}
std::vector<std::string> spatial_dim_names = {", x", ", y", ", z", ", w"};
std::vector<std::string> spatial_dim_names = {", x", ", y", ", z", ", w", ", u", ", v"};
for (size_t i = 0; i < spatial.size(); ++i) {
out << spatial_dim_names[i] << ":" << spatial[i];
}
@@ -471,84 +471,46 @@ public:
* @endcode
*/
tensor transform(cldnn::format new_fmt, value_type default_size) const {
cldnn::format format = cldnn::format::bfwzyx;
cldnn::format format = cldnn::format::bfvuwzyx;
auto val_order = format.internal_order();
auto new_order = new_fmt.internal_order();
std::vector<value_type> old_sizes = sizes();
std::vector<value_type> new_sizes(old_sizes.size(), default_size);
auto tmp = 1;
auto tmp_z = 1;
auto tmp_w = 1;
for (size_t i = 0; i < format.order().size(); i++) {
auto c = val_order[i];
// skip f and y, z for the formats that do not have it
if (((new_fmt == format::bs_xs_xsv8_bsv8) ||
(new_fmt == format::bs_xs_xsv8_bsv16) ||
(new_fmt == format::os_i_osv8__ai8) ||
(new_fmt == format::os_i_osv16__ai8) ||
(new_fmt == format::bs_x_bsv16)) &&
((c == 'f') ||
(c == 'y') ||
(c == 'z') ||
(c == 'w'))) {
if (new_order[i] == '?')
new_sizes[i] = default_size;
const auto& new_traits = format::traits(new_fmt);
const cldnn::format default_fmt = cldnn::format::bfvuwzyx;
static const std::map<char, char> flatten_mapping = {
{ 'v', 'u'},
{ 'u', 'w'},
{ 'w', 'z'},
{ 'z', 'y'}
};
tmp *= old_sizes[i];
continue;
for (size_t i = 0; i < default_fmt.order().size(); i++) {
auto target_dim = val_order[i]; //bfxywzuv
while (!new_traits.has_dimension(target_dim)) {
if (flatten_mapping.find(target_dim) != flatten_mapping.end()) {
target_dim = flatten_mapping.at(target_dim);
} else {
target_dim = new_fmt.order().back();
}
}
// skip z for the formats that do not have it
if (((new_fmt != format::bfzyx && new_fmt != format::b_fs_zyx_fsv16 && new_fmt != format::b_fs_zyx_fsv32 && new_fmt != format::bzyxf &&
new_fmt != format::bfwzyx && new_fmt != format::bs_fs_zyx_bsv16_fsv16 && new_fmt != format::bs_fs_zyx_bsv16_fsv32 &&
new_fmt != format::bs_fs_zyx_bsv32_fsv16 && new_fmt != format::bs_fs_zyx_bsv32_fsv32 &&
new_fmt != format::b_fs_zyx_fsv2 && new_fmt != format::b_fs_zyx_fsv4 &&
new_fmt != format::bs_fs_zyx_bsv8_fsv2 && new_fmt != format::bs_fs_zyx_bsv8_fsv4 &&
new_fmt != format::bs_fs_zyx_bsv16_fsv2 && new_fmt != format::bs_fs_zyx_bsv16_fsv4)) && (c == 'z')) {
if (new_order[i] == '?')
new_sizes[i] = default_size;
tmp_z *= old_sizes[i];
continue;
auto new_pos = new_order.find(target_dim);
if (new_pos != std::string::npos) {
if (new_sizes[new_pos] == -1) {
new_sizes[new_pos] = old_sizes[i];
} else {
new_sizes[new_pos] *= old_sizes[i];
}
}
if (new_fmt != format::bfwzyx && c == 'w') {
if (new_order[i] == '?')
new_sizes[i] = default_size;
if (new_fmt == format::bfzyx || new_fmt == format::b_fs_zyx_fsv16 ||
new_fmt == format::bs_fs_zyx_bsv16_fsv16 || new_fmt == format::b_fs_zyx_fsv32 ||
new_fmt == format::bs_fs_zyx_bsv16_fsv32)
tmp_w *= old_sizes[i];
else
tmp_z *= old_sizes[i];
continue;
}
auto new_pos = new_order.find(c);
if (new_pos == std::string::npos)
throw std::invalid_argument("cannot convert to new format");
new_sizes[new_pos] = old_sizes[i];
}
// in case of formats with smaller number of dimensions than input, flatten is performed below
if (tmp != 1 || tmp_z != 1 || tmp_w != 1) {
for (size_t i = 0; i < format.order().size(); i++) {
auto c = val_order[i];
if (c == 'x') {
auto new_pos = new_order.find(c);
new_sizes[new_pos] *= tmp;
}
if (c == 'y') {
auto new_pos = new_order.find(c);
if (new_pos != std::string::npos)
new_sizes[new_pos] *= tmp_z;
}
if (c == 'z') {
auto new_pos = new_order.find(c);
if (new_pos != std::string::npos)
new_sizes[new_pos] *= tmp_w;
}
for (size_t i = 0; i < new_order.size(); i++) {
auto c = new_order[i]; //bfxywz
if (c == '?')
continue;
if (new_sizes[i] == -1) {
new_sizes[i] = 1;
}
}
src/plugins/intel_gpu/src/graph/fully_connected.cpp
+2 -2
View File
@@ -50,8 +50,8 @@ format::type get_preferred_format(fully_connected_node const& node, const kernel
if (data_type_traits::is_floating_point(input_layout.data_type) &&
(is_batch_after_spatial(input_layout.format.order()) ||
input_layout.format == format::bs_x_bsv16 ||
input_layout.format == format::bs_xs_xsv8_bsv8))
input_layout.format == format::bs_f_bsv16 ||
input_layout.format == format::bs_fs_fsv8_bsv8))
return format::yxfb;
bool no_spatial_padding = true;
src/plugins/intel_gpu/src/graph/graph_optimizer/remove_redundant_reorders.cpp
+2 -2
View File
@@ -20,8 +20,8 @@
using namespace cldnn;
#define LOG_NODE_REMOVAL(id) GPU_DEBUG_LOG_PASS << "Remove node: " << (id) << std::endl;
#define LOG_NODE_REPLACEMENT(id) GPU_DEBUG_LOG_PASS << "Replace node: " << (id) << std::endl;
#define LOG_NODE_REMOVAL(id) GPU_DEBUG_LOG_PASS << __func__ << ":" << __LINE__ << ": remove node: " << (id) << std::endl;
#define LOG_NODE_REPLACEMENT(id) GPU_DEBUG_LOG_PASS << __func__ << ":" << __LINE__ << ": replace node: " << (id) << std::endl;
remove_redundant_reorders::remove_redundant_reorders(layout_optimizer& lo_ref, bool enable_reorder_fusing, bool update_implementations,
bool remove_output_reorders)
src/plugins/intel_gpu/src/graph/graph_optimizer/reorder_inputs.cpp
+4 -1
View File
@@ -13,7 +13,6 @@
#include "mvn_inst.h"
#include "to_string_utils.h"
#include "pooling_inst.h"
#include "reshape_inst.h"
#ifdef ENABLE_ONEDNN_FOR_GPU
#include "gemm_inst.h"
@@ -285,6 +284,8 @@ void propagate_formats_rec(std::map<program_node*, format::type>& fmt_map,
for (auto next : travel_direction_wrapper<dir>::next_nodes(node)) {
if (!next->is_in_data_flow())
continue;
if (!can_propagate_formats_rec<dir>(fmt_map, lo, node, next, fmt))
continue;
propagate_formats_rec<dir>(fmt_map, lo, node, next, fmt);
}
}
@@ -511,6 +512,8 @@ void minimize_local_reorders(program& p, std::map<program_node*, format::type>&
continue;
for (auto new_fmt : local_formats) {
if (fmt_map.at(node) != format::any && format::dimension(fmt_map.at(node)) != format::dimension(new_fmt))
continue;
fmt_map.at(node) = new_fmt;
auto reorders_cnt = count_reorders(fmt_map, lo, node);
src/plugins/intel_gpu/src/graph/impls/ocl/eltwise.cpp
+18 -1
View File
@@ -133,6 +133,7 @@ public:
input_pshape = extend_shape_to_rank_from_begin(input_pshape, out_pshape.size());
}
input_layout.set_partial_shape(extend_shape_to_rank_from_end(input_pshape));
input_layout.format = format::adjust_to_rank(input_layout.format, input_pshape.size());
}
return updated_impl_params;
@@ -163,7 +164,9 @@ attach_eltwise_impl::attach_eltwise_impl() {
auto dyn_formats = {
format::bfyx,
format::bfzyx,
format::bfwzyx
format::bfwzyx,
format::bfuwzyx,
format::bfvuwzyx,
};
implementation_map<eltwise>::add(impl_types::ocl,
@@ -213,6 +216,20 @@ attach_eltwise_impl::attach_eltwise_impl() {
std::make_tuple(data_types::i32, format::bfwzyx),
std::make_tuple(data_types::i64, format::bfwzyx),
std::make_tuple(data_types::f32, format::bfuwzyx),
std::make_tuple(data_types::f16, format::bfuwzyx),
std::make_tuple(data_types::i8, format::bfuwzyx),
std::make_tuple(data_types::u8, format::bfuwzyx),
std::make_tuple(data_types::i32, format::bfuwzyx),
std::make_tuple(data_types::i64, format::bfuwzyx),
std::make_tuple(data_types::f32, format::bfvuwzyx),
std::make_tuple(data_types::f16, format::bfvuwzyx),
std::make_tuple(data_types::i8, format::bfvuwzyx),
std::make_tuple(data_types::u8, format::bfvuwzyx),
std::make_tuple(data_types::i32, format::bfvuwzyx),
std::make_tuple(data_types::i64, format::bfvuwzyx),
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::i8, format::b_fs_zyx_fsv16),
src/plugins/intel_gpu/src/graph/impls/ocl/fully_connected.cpp
+3
View File
@@ -151,6 +151,9 @@ attach_fully_connected_impl::attach_fully_connected_impl() {
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::f16, format::fs_b_yx_fsv32),
std::make_tuple(data_types::f32, format::bs_fs_fsv8_bsv8),
std::make_tuple(data_types::f16, format::bs_fs_fsv8_bsv8),
std::make_tuple(data_types::f16, format::bs_fs_fsv8_bsv16),
});
}
src/plugins/intel_gpu/src/graph/impls/ocl/kernel_selector_helper.cpp
+16 -8
View File
@@ -223,9 +223,9 @@ kernel_selector::data_layout to_data_layout(format f) {
return kernel_selector::data_layout::b_fs_zyx_fsv4;
case format::b_fs_zyx_fsv32:
return kernel_selector::data_layout::b_fs_zyx_fsv32;
case format::bs_x_bsv16:
case format::bs_f_bsv16:
return kernel_selector::data_layout::bs_f_bsv16__af8;
case format::bs_xs_xsv8_bsv8:
case format::bs_fs_fsv8_bsv8:
return kernel_selector::data_layout::bs_f_bsv8__af8;
case format::winograd_2x3_s1_data:
return kernel_selector::data_layout::winograd_2x3_s1_data;
@@ -239,6 +239,10 @@ kernel_selector::data_layout to_data_layout(format f) {
return kernel_selector::data_layout::fs_b_yx_fsv32;
case format::bfwzyx:
return kernel_selector::data_layout::bfwzyx;
case format::bfuwzyx:
return kernel_selector::data_layout::bfuwzyx;
case format::bfvuwzyx:
return kernel_selector::data_layout::bfvuwzyx;
case format::b_fs_zyx_fsv16:
return kernel_selector::data_layout::b_fs_zyx_fsv16;
case format::bs_fs_yx_bsv16_fsv32:
@@ -282,7 +286,7 @@ kernel_selector::data_layout to_data_layout(format f) {
case format::image_2d_rgba:
return kernel_selector::data_layout::image_2d_rgba;
default:
throw std::invalid_argument("Format f (" + std::to_string((int32_t)f.value) + ") is not a proper data layout");
OPENVINO_THROW("[GPU] Can't convert tensor format to kernel selector format as f=", f, " is not handled");
}
}
@@ -311,9 +315,9 @@ cldnn::format from_data_layout(kernel_selector::data_layout l) {
case kernel_selector::data_layout::b_fs_zyx_fsv32:
return cldnn::format::b_fs_zyx_fsv32;
case kernel_selector::data_layout::bs_f_bsv8__af8:
return cldnn::format::bs_xs_xsv8_bsv8;
return cldnn::format::bs_fs_fsv8_bsv8;
case kernel_selector::data_layout::bs_f_bsv16__af8:
return cldnn::format::bs_x_bsv16;
return cldnn::format::bs_f_bsv16;
case kernel_selector::data_layout::winograd_2x3_s1_data:
return cldnn::format::winograd_2x3_s1_data;
case kernel_selector::data_layout::b_fs_yx_32fp:
@@ -324,6 +328,10 @@ cldnn::format from_data_layout(kernel_selector::data_layout l) {
return cldnn::format::fs_b_yx_fsv32;
case kernel_selector::data_layout::bfwzyx:
return cldnn::format::bfwzyx;
case kernel_selector::data_layout::bfuwzyx:
return cldnn::format::bfuwzyx;
case kernel_selector::data_layout::bfvuwzyx:
return cldnn::format::bfvuwzyx;
case kernel_selector::data_layout::bs_fs_yx_bsv16_fsv16:
return cldnn::format::bs_fs_yx_bsv16_fsv16;
case kernel_selector::data_layout::bs_fs_zyx_bsv32_fsv16:
@@ -482,12 +490,12 @@ kernel_selector::weights_layout to_weights_layout(format f, bool is_grouped) {
return kernel_selector::weights_layout::oizyx;
case format::iozyx:
return kernel_selector::weights_layout::iozyx;
case format::bs_xs_xsv8_bsv8:
case format::bs_fs_fsv8_bsv8:
case format::os_i_osv8__ai8:
return kernel_selector::weights_layout::os_i_osv8__ai8;
case format::os_i_osv16__ai8:
return kernel_selector::weights_layout::os_i_osv16__ai8;
case format::bs_x_bsv16:
case format::bs_f_bsv16:
return kernel_selector::weights_layout::os_i_osv16;
case format::os_is_zyx_isv16_osv16:
return kernel_selector::weights_layout::os_is_zyx_isv16_osv16;
@@ -656,7 +664,7 @@ cldnn::format::type from_weights_layout(kernel_selector::weights_layout l) {
case kernel_selector::weights_layout::os_iyx_osv64:
return cldnn::format::os_iyx_osv64;
case kernel_selector::weights_layout::os_i_osv16:
return cldnn::format::bs_x_bsv16;
return cldnn::format::bs_f_bsv16;
case kernel_selector::weights_layout::os_i_osv8__ai8:
return cldnn::format::os_i_osv8__ai8;
case kernel_selector::weights_layout::os_i_osv16__ai8:
src/plugins/intel_gpu/src/graph/impls/ocl/permute.cpp
+3 -1
View File
@@ -84,7 +84,9 @@ attach_permute_impl::attach_permute_impl() {
auto dyn_formats = {
format::bfyx,
format::bfzyx,
format::bfwzyx
format::bfwzyx,
format::bfuwzyx,
format::bfvuwzyx,
};
implementation_map<permute>::add(impl_types::ocl,
src/plugins/intel_gpu/src/graph/impls/ocl/primitive_base.hpp
+9 -2
View File
@@ -248,14 +248,21 @@ protected:
is_output_event = instance.is_output_event();
}
auto& params = _kernel_data.kernels[kd_idx].params;
auto args = get_arguments(instance);
args.scalars = &_kernel_data.kernels[kd_idx].params.scalars;
args.scalars = &params.scalars;
for (const auto& m : instance.get_intermediates_memories()) {
args.intermediates.push_back(m);
}
auto ev = stream.enqueue_kernel(*_kernels[kd_idx], _kernel_data.kernels[kd_idx].params, args, tmp_events, is_output_event);
const auto& gws = params.workGroups.global;
const auto& lws = params.workGroups.local;
GPU_DEBUG_TRACE_DETAIL << "Enqueue kernel " << kd_idx << ": gws=[" << gws[0] << ", " << gws[1] << ", " << gws[2] << "] "
<< "lws=[" << lws[0] << ", " << lws[1] << ", " << lws[2] << "]" << std::endl;
auto ev = stream.enqueue_kernel(*_kernels[kd_idx], params, args, tmp_events, is_output_event);
new_events.push_back(ev);
all_events.push_back(ev);
src/plugins/intel_gpu/src/graph/impls/ocl/reduce.cpp
+5 -1
View File
@@ -105,6 +105,8 @@ attach_reduce_impl::attach_reduce_impl() {
format::bfyx,
format::bfzyx,
format::bfwzyx,
format::bfuwzyx,
format::bfvuwzyx,
format::b_fs_yx_fsv16,
format::b_fs_yx_fsv32,
format::b_fs_zyx_fsv16
@@ -119,7 +121,9 @@ attach_reduce_impl::attach_reduce_impl() {
auto dyn_formats = {
format::bfyx,
format::bfzyx,
format::bfwzyx
format::bfwzyx,
format::bfuwzyx,
format::bfvuwzyx
};
implementation_map<reduce>::add(impl_types::ocl,
src/plugins/intel_gpu/src/graph/impls/ocl/reshape.cpp
+3 -1
View File
@@ -49,7 +49,9 @@ attach_reshape_impl::attach_reshape_impl() {
auto dyn_formats = {
format::bfyx,
format::bfzyx,
format::bfwzyx
format::bfwzyx,
format::bfuwzyx,
format::bfvuwzyx
};
implementation_map<reshape>::add(impl_types::ocl,
src/plugins/intel_gpu/src/graph/layout_optimizer.cpp
+15 -2
View File
@@ -440,6 +440,10 @@ bool layout_optimizer::can_fuse_reorder_to_prev(program_node& prev, reorder_node
&& (!prev.as<permute>().is_rotating_except_batch())) {
return false;
}
// permute kernel doesn't support reorder fusion for ranks > 6
if (fmt_prev.dimension() > 6 || fmt_next.dimension() > 6)
return false;
return true;
}
@@ -1582,8 +1586,17 @@ format layout_optimizer::get_preferred_format(program_node& node) {
auto dep_size = node.get_dependencies().size();
for (size_t i = 0; i < dep_size; i++) {
auto in_lay_rank = node.get_dependency(i).get_output_layout(false).get_rank();
if (in_lay_rank != out_lay_rank)
node.set_preferred_input_fmt(i, get_preferred_format(node.get_dependency(i)));
if (in_lay_rank != out_lay_rank) {
auto fmt = get_preferred_format(node.get_dependency(i));
// Check if selected format can be adjusted to the required output rank
// If no, use default fotmat instead
try {
format::adjust_to_rank(fmt, out_lay_rank);
} catch (ov::Exception&) {
fmt = format::get_default_format(out_lay_rank);
}
node.set_preferred_input_fmt(i, fmt);
}
}
// shape_infer_dep should be plain format because the memory is being read by ngraph shape infer as is
src/plugins/intel_gpu/src/graph/network.cpp
+3
View File
@@ -131,6 +131,7 @@ void dump_perf_data_raw(std::string dump_path, const std::list<std::shared_ptr<p
}
}
float convert_element(int64_t i) { return static_cast<float>(i); }
float convert_element(int32_t i) { return static_cast<float>(i); }
float convert_element(float f) { return f; }
@@ -265,6 +266,8 @@ void log_memory_to_file(memory::ptr mem, stream& stream, std::string layerName,
dump<half_t>(mem, stream, file_stream, dump_raw);
else if (mem_dt == cldnn::data_types::bin)
dump<uint32_t>(mem, stream, file_stream, dump_raw);
else if (mem_dt == cldnn::data_types::i64)
dump<int64_t>(mem, stream, file_stream, dump_raw);
else if (mem_dt == cldnn::data_types::i32)
dump<int32_t>(mem, stream, file_stream, dump_raw);
else if (mem_dt == cldnn::data_types::i8)
src/plugins/intel_gpu/src/graph/reorder.cpp
+1 -1
View File
@@ -146,7 +146,7 @@ layout reorder_inst::calc_output_layout(reorder_node const& node, kernel_impl_pa
"Conversion of weights from winograd to standard domain is currently unsupported");
}
if ((ofmt == format::bs_xs_xsv8_bsv8 || ofmt == format::os_i_osv8__ai8 || ofmt == format::os_i_osv16__ai8 || ofmt == format::bs_x_bsv16 ||
if ((ofmt == format::bs_fs_fsv8_bsv8 || ofmt == format::os_i_osv8__ai8 || ofmt == format::os_i_osv16__ai8 || ofmt == format::bs_f_bsv16 ||
ofmt == format::bfzyx || ifmt == format::bfzyx || ofmt == format::b_fs_zyx_fsv16 || ifmt == format::b_fs_zyx_fsv16 ||
ofmt == format::bs_fs_zyx_bsv16_fsv16 || ifmt == format::bs_fs_zyx_bsv16_fsv16 ||
ofmt == format::bs_fs_zyx_bsv16_fsv32 || ifmt == format::bs_fs_zyx_bsv16_fsv32 ||
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bs_f_bsv16_b1.cl
+2 -2
View File
@@ -160,7 +160,7 @@
REQD_SUB_GROUP_SIZE(SUB_GROUP_SIZE)
__attribute__((reqd_work_group_size(SUB_GROUP_SIZE, 1, 1)))
KERNEL (fully_connected_gpu_bx_bs_x_bsv16_b1)(
KERNEL (fully_connected_gpu_bx_bs_f_bsv16_b1)(
const __global UNIT_TYPE* input,
__global UNIT_TYPE* output,
const __global UNIT_TYPE* weight
@@ -187,7 +187,7 @@ KERNEL (fully_connected_gpu_bx_bs_x_bsv16_b1)(
// Input base offset in bytes (bfyx/bx format of input).
const uint input_base = 0;
// Filter base offset in bytes (bs_x_bsv16 format of weights).
// Filter base offset in bytes (bs_f_bsv16 format of weights).
const uint filter_base = sg_id * input_byte_size * RESPONSES_PER_SG_EXEC;
// [SCATTERED] Output base identifier/element offset to use (bx format of output).
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bs_f_bsv8_af8_vload.cl
+1 -1
View File
@@ -57,7 +57,7 @@
__attribute__((reqd_work_group_size(SUB_GROUP_SIZE, 1, 1)))
REQD_SUB_GROUP_SIZE(SUB_GROUP_SIZE)
KERNEL (fully_connected_gpu_xb_bs_xs_xsv8_bsv8_vload)(
KERNEL (fully_connected_gpu_xb_bs_fs_fsv8_bsv8_vload)(
const __global UNIT_TYPE* input,
__global UNIT_TYPE* output,
const __global UNIT_TYPE* weight
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/generic_eltwise_ref.cl
+75 -1
View File
@@ -25,7 +25,81 @@ KERNEL(eltwise)(
)
{
#if OUTPUT_DIMS == 6 // 4D spatial
#if OUTPUT_DIMS == 8 // 6D spatial
#if ELTWISE_LAYOUT_BASED || QUANTIZATION_TERM || ELTWISE_BROADCAST
uint data_idx = get_global_id(GWS_YX);
const uint d1 = data_idx % OUTPUT_SIZE_X; // X
data_idx = data_idx / OUTPUT_SIZE_X;
const uint d2 = data_idx % OUTPUT_SIZE_Y; // Y
data_idx = data_idx / OUTPUT_SIZE_Y;
const uint d3 = data_idx % OUTPUT_SIZE_Z; // Z
data_idx = data_idx / OUTPUT_SIZE_Z;
const uint d4 = data_idx % OUTPUT_SIZE_W; // W
data_idx = data_idx / OUTPUT_SIZE_W;
const uint d5 = data_idx % OUTPUT_SIZE_U; // U
data_idx = data_idx / OUTPUT_SIZE_U;
const uint d6 = data_idx % OUTPUT_SIZE_V; // V
const uint d7 = get_global_id(GWS_FEATURE); // Feature
const uint d8 = get_global_id(GWS_BATCH); // Batch
uint output_offset = OUTPUT_GET_INDEX(d8, d7, d6, d5, d4, d3, d2, d1);
#elif ELTWISE_NO_PITCH_SAME_DIMS
const uint d1 = get_global_id(0);
uint output_offset = OUTPUT_OFFSET + d1;
#else
const uint d1 = get_global_id(0);
const uint d2 = (uint)get_global_id(1) % OUTPUT_SIZES[1];
const uint d3 = (uint)get_global_id(1) / OUTPUT_SIZES[1] % OUTPUT_SIZES[2];
const uint d4 = (uint)get_global_id(1) / OUTPUT_SIZES[1] / OUTPUT_SIZES[2] % OUTPUT_SIZES[3];
const uint d5 = (uint)get_global_id(1) / OUTPUT_SIZES[1] / OUTPUT_SIZES[2] / OUTPUT_SIZES[3] % OUTPUT_SIZES[4];
const uint d6 = (uint)get_global_id(1) / OUTPUT_SIZES[1] / OUTPUT_SIZES[2] / OUTPUT_SIZES[3] / OUTPUT_SIZES[4];
const uint d7 = (uint)get_global_id(2) % OUTPUT_SIZES[6];
const uint d8 = (uint)get_global_id(2) / OUTPUT_SIZES[6];
uint output_offset = OUTPUT_GET_INDEX(d8, d7, d6, d5, d4, d3, d2, d1);
#endif
#elif OUTPUT_DIMS == 7 // 5D spatial
#if ELTWISE_LAYOUT_BASED || QUANTIZATION_TERM || ELTWISE_BROADCAST
uint data_idx = get_global_id(GWS_YX);
const uint d1 = data_idx % OUTPUT_SIZE_X; // X
data_idx = data_idx / OUTPUT_SIZE_X;
const uint d2 = data_idx % OUTPUT_SIZE_Y; // Y
data_idx = data_idx / OUTPUT_SIZE_Y;
const uint d3 = data_idx % OUTPUT_SIZE_Z; // Z
data_idx = data_idx / OUTPUT_SIZE_Z;
const uint d4 = data_idx % OUTPUT_SIZE_W; // W
data_idx = data_idx / OUTPUT_SIZE_W;
const uint d5 = data_idx % OUTPUT_SIZE_U; // U
const uint d6 = get_global_id(GWS_FEATURE); // Feature
const uint d7 = get_global_id(GWS_BATCH); // Batch
uint output_offset = OUTPUT_GET_INDEX(d7, d6, d5, d4, d3, d2, d1);
#elif ELTWISE_NO_PITCH_SAME_DIMS
const uint d1 = get_global_id(0);
uint output_offset = OUTPUT_OFFSET + d1;
#else
const uint d1 = get_global_id(0);
const uint d2 = (uint)get_global_id(1) % OUTPUT_SIZES[1];
const uint d3 = (uint)get_global_id(1) / OUTPUT_SIZES[1] % OUTPUT_SIZES[2];
const uint d4 = (uint)get_global_id(1) / OUTPUT_SIZES[1] / OUTPUT_SIZES[2] % OUTPUT_SIZES[3];
const uint d5 = (uint)get_global_id(1) / OUTPUT_SIZES[1] / OUTPUT_SIZES[2] / OUTPUT_SIZES[3];
const uint d6 = (uint)get_global_id(2) % OUTPUT_SIZES[5];
const uint d7 = (uint)get_global_id(2) / OUTPUT_SIZES[5];
uint output_offset = OUTPUT_GET_INDEX(d7, d6, d5, d4, d3, d2, d1);
#endif
#elif OUTPUT_DIMS == 6 // 4D spatial
#if ELTWISE_LAYOUT_BASED || QUANTIZATION_TERM || ELTWISE_BROADCAST
uint data_idx = get_global_id(GWS_YX);
const uint d1 = data_idx % OUTPUT_SIZE_X; // X
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/batch_headers/fetch_data.cl
+62 -1
View File
@@ -76,6 +76,68 @@
(i4)*CAT(prefix, _PITCHES)[4] + \
(i5)*CAT(prefix, _PITCHES)[5]
#define GET_DATA_INDEX_7D(prefix, b, f, u, w, z, y, x) \
CAT(prefix, _OFFSET) + \
(x)*CAT(prefix, _X_PITCH) + \
(y)*CAT(prefix, _Y_PITCH) + \
(z)*CAT(prefix, _Z_PITCH) + \
(w)*CAT(prefix, _W_PITCH) + \
(u)*CAT(prefix, _U_PITCH) + \
(f)*CAT(prefix, _FEATURE_PITCH) + \
(b)*CAT(prefix, _BATCH_PITCH)
#define GET_DATA_INDEX_7D_SAFE(prefix, b, f, u, w, z, y, x) \
CAT(prefix, _OFFSET) + \
(x % CAT(prefix, _SIZE_X ))*CAT(prefix, _X_PITCH) + \
(y % CAT(prefix, _SIZE_Y ))*CAT(prefix, _Y_PITCH) + \
(z % CAT(prefix, _SIZE_Z ))*CAT(prefix, _Z_PITCH) + \
(w % CAT(prefix, _SIZE_W ))*CAT(prefix, _W_PITCH) + \
(u % CAT(prefix, _SIZE_U ))*CAT(prefix, _U_PITCH) + \
(f % CAT(prefix, _FEATURE_NUM))*CAT(prefix, _FEATURE_PITCH) + \
(b % CAT(prefix, _BATCH_NUM ))*CAT(prefix, _BATCH_PITCH)
#define GET_DATA_INDEX_7D_RAW(prefix, i0, i1, i2, i3, i4, i5, i6) \
CAT(prefix, _OFFSET) + \
(i0)*CAT(prefix, _PITCHES)[0] + \
(i1)*CAT(prefix, _PITCHES)[1] + \
(i2)*CAT(prefix, _PITCHES)[2] + \
(i3)*CAT(prefix, _PITCHES)[3] + \
(i4)*CAT(prefix, _PITCHES)[4] + \
(i5)*CAT(prefix, _PITCHES)[5] + \
(i6)*CAT(prefix, _PITCHES)[6]
#define GET_DATA_INDEX_8D(prefix, b, f, v, u, w, z, y, x) \
CAT(prefix, _OFFSET) + \
(x)*CAT(prefix, _X_PITCH) + \
(y)*CAT(prefix, _Y_PITCH) + \
(z)*CAT(prefix, _Z_PITCH) + \
(w)*CAT(prefix, _W_PITCH) + \
(u)*CAT(prefix, _U_PITCH) + \
(v)*CAT(prefix, _V_PITCH) + \
(f)*CAT(prefix, _FEATURE_PITCH) + \
(b)*CAT(prefix, _BATCH_PITCH)
#define GET_DATA_INDEX_8D_SAFE(prefix, b, f, v, u, w, z, y, x) \
CAT(prefix, _OFFSET) + \
(x % CAT(prefix, _SIZE_X ))*CAT(prefix, _X_PITCH) + \
(y % CAT(prefix, _SIZE_Y ))*CAT(prefix, _Y_PITCH) + \
(z % CAT(prefix, _SIZE_Z ))*CAT(prefix, _Z_PITCH) + \
(w % CAT(prefix, _SIZE_W ))*CAT(prefix, _W_PITCH) + \
(u % CAT(prefix, _SIZE_U ))*CAT(prefix, _U_PITCH) + \
(v % CAT(prefix, _SIZE_V ))*CAT(prefix, _V_PITCH) + \
(f % CAT(prefix, _FEATURE_NUM))*CAT(prefix, _FEATURE_PITCH) + \
(b % CAT(prefix, _BATCH_NUM ))*CAT(prefix, _BATCH_PITCH)
#define GET_DATA_INDEX_8D_RAW(prefix, i0, i1, i2, i3, i4, i5, i6, i7) \
CAT(prefix, _OFFSET) + \
(i0)*CAT(prefix, _PITCHES)[0] + \
(i1)*CAT(prefix, _PITCHES)[1] + \
(i2)*CAT(prefix, _PITCHES)[2] + \
(i3)*CAT(prefix, _PITCHES)[3] + \
(i4)*CAT(prefix, _PITCHES)[4] + \
(i5)*CAT(prefix, _PITCHES)[5] + \
(i6)*CAT(prefix, _PITCHES)[6] + \
(i7)*CAT(prefix, _PITCHES)[7]
#define GET_DATA_BS_FYX_BSV8_INDEX(prefix, b, f, y, x, sub_group_size) \
CAT(prefix, _OFFSET) + \
@@ -1194,4 +1256,3 @@ inline uint get_bs_fs_zyx_bsv_fsv_index(uint b, uint f, uint z, uint y, uint x,
CAT(prefix, _PAD_AFTER_SIZE_Y), \
CAT(prefix, _PAD_BEFORE_SIZE_X), \
CAT(prefix, _PAD_AFTER_SIZE_X), 16, 16)
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/fetch_utils.cl
+18 -2
View File
@@ -4,26 +4,42 @@
#include "batch_headers/fetch_data.cl"
inline uint FUNC(get_input_index)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x) {
inline uint FUNC(get_input_index)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint v, uint u, uint w, uint z, uint y, uint x) __attribute__((overloadable)) {
#if INPUT0_DIMS < 5
return INPUT0_GET_INDEX(b, f, y, x);
#elif INPUT0_DIMS == 5
return INPUT0_GET_INDEX(b, f, z, y, x);
#elif INPUT0_DIMS == 6
return INPUT0_GET_INDEX(b, f, w, z, y, x);
#elif INPUT0_DIMS == 7
return INPUT0_GET_INDEX(b, f, u, w, z, y, x);
#elif INPUT0_DIMS == 8
return INPUT0_GET_INDEX(b, f, v, u, w, z, y, x);
#else
#error [GPU] Unsupported input tensor rank in get_input_index function
#endif
}
inline uint FUNC(get_output_index)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x) {
inline uint FUNC(get_input_index)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x) __attribute__((overloadable)) {
return FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR b, f, 0, 0, w, z, y, x);
}
inline uint FUNC(get_output_index)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint v, uint u, uint w, uint z, uint y, uint x) __attribute__((overloadable)) {
#if OUTPUT_DIMS < 5
return OUTPUT_GET_INDEX(b, f, y, x);
#elif OUTPUT_DIMS == 5
return OUTPUT_GET_INDEX(b, f, z, y, x);
#elif OUTPUT_DIMS == 6
return OUTPUT_GET_INDEX(b, f, w, z, y, x);
#elif OUTPUT_DIMS == 7
return OUTPUT_GET_INDEX(b, f, u, w, z, y, x);
#elif OUTPUT_DIMS == 8
return OUTPUT_GET_INDEX(b, f, v, u, w, z, y, x);
#else
#error [GPU] Unsupported output tensor rank in get_output_index function
#endif
}
inline uint FUNC(get_output_index)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x) __attribute__((overloadable)) {
return FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR b, f, 0, 0, w, z, y, x);
}
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/reshape_dims.cl
+44 -257
View File
@@ -5,282 +5,69 @@
// ! Huge warning !
// Reshape functions below assume both input and output formats are bfyx, bfzyx, bfwzyx, etc.
inline uint8 FUNC(reshape_2_to_4)(uint o, uint i, uint y, uint x, uint dst_size_y, uint dst_size_x)
{
uint _i = i / (dst_size_y*dst_size_x);
uint _yx = i % (dst_size_y*dst_size_x);
uint _y = _yx / dst_size_x;
uint _x = _yx % dst_size_x;
return (uint8)(0, o, _i, 0, 0, _y,_x, 0);
}
inline uint8 FUNC(reshape_4_to_2)(uint o, uint i, uint y, uint x, uint src_size_y, uint src_size_x)
{
uint _i = i*src_size_y*src_size_x + y*src_size_x + x;
return (uint8)(0, o, _i, 0, 0, 0, 0, 0);
}
inline uint8 FUNC(reshape_4_to_5)(uint o, uint i, uint y, uint x,
uint src_size_f, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_z, uint dst_size_y, uint dst_size_x)
{
const uint src_pitch_x = 1;
const uint src_pitch_y = src_pitch_x * src_size_x;
const uint src_pitch_f = src_pitch_y * src_size_y;
const uint src_pitch_b = src_pitch_f * src_size_f;
uint flat_idx = x * src_pitch_x + y * src_pitch_y + i * src_pitch_f + o * src_pitch_b;
uint dst_x = flat_idx % dst_size_x;
flat_idx /= dst_size_x;
uint dst_y = flat_idx % dst_size_y;
flat_idx /= dst_size_y;
uint dst_z = flat_idx % dst_size_z;
flat_idx /= dst_size_z;
uint dst_f = flat_idx % dst_size_f;
flat_idx /= dst_size_f;
uint dst_b = flat_idx;
return (uint8)(0, dst_b, dst_f, 0, dst_z, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_5_to_4)(uint o, uint i, uint z, uint y, uint x,
uint src_size_f, uint src_size_z, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_y, uint dst_size_x)
{
const uint src_pitch_x = 1;
const uint src_pitch_y = src_pitch_x * src_size_x;
const uint src_pitch_z = src_pitch_y * src_size_y;
const uint src_pitch_f = src_pitch_z * src_size_z;
const uint src_pitch_b = src_pitch_f * src_size_f;
uint flat_idx = x * src_pitch_x + y * src_pitch_y + z * src_pitch_z + i * src_pitch_f + o * src_pitch_b;
uint dst_x = flat_idx % dst_size_x;
flat_idx /= dst_size_x;
uint dst_y = flat_idx % dst_size_y;
flat_idx /= dst_size_y;
uint dst_f = flat_idx % dst_size_f;
flat_idx /= dst_size_f;
uint dst_b = flat_idx;
return (uint8)(0, dst_b, dst_f, 0, 0, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_4_to_6)(uint o, uint i, uint y, uint x,
uint src_size_f, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_w, uint dst_size_z, uint dst_size_y, uint dst_size_x)
{
const uint src_pitch_x = 1;
const uint src_pitch_y = src_pitch_x * src_size_x;
const uint src_pitch_f = src_pitch_y * src_size_y;
const uint src_pitch_b = src_pitch_f * src_size_f;
uint flat_idx = x * src_pitch_x + y * src_pitch_y + i * src_pitch_f + o * src_pitch_b;
uint dst_x = flat_idx % dst_size_x;
flat_idx /= dst_size_x;
uint dst_y = flat_idx % dst_size_y;
flat_idx /= dst_size_y;
uint dst_z = flat_idx % dst_size_z;
flat_idx /= dst_size_z;
uint dst_w = flat_idx % dst_size_w;
flat_idx /= dst_size_w;
uint dst_f = flat_idx % dst_size_f;
flat_idx /= dst_size_f;
uint dst_b = flat_idx;
return (uint8)(0, dst_b, dst_f, dst_w, dst_z, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_6_to_4)(uint o, uint i, uint w, uint z, uint y, uint x,
uint src_size_f, uint src_size_w, uint src_size_z, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_y, uint dst_size_x)
{
const uint src_pitch_x = 1;
const uint src_pitch_y = src_pitch_x * src_size_x;
const uint src_pitch_z = src_pitch_y * src_size_y;
const uint src_pitch_w = src_pitch_z * src_size_z;
const uint src_pitch_f = src_pitch_w * src_size_w;
const uint src_pitch_b = src_pitch_f * src_size_f;
uint flat_idx = x * src_pitch_x + y * src_pitch_y + z * src_pitch_z + w * src_pitch_w + i * src_pitch_f + o * src_pitch_b;
uint dst_x = flat_idx % dst_size_x;
flat_idx /= dst_size_x;
uint dst_y = flat_idx % dst_size_y;
flat_idx /= dst_size_y;
uint dst_f = flat_idx % dst_size_f;
flat_idx /= dst_size_f;
uint dst_b = flat_idx;
return (uint8)(0, dst_b, dst_f, 0, 0, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_6_to_5)(uint o, uint i, uint w, uint z, uint y, uint x,
uint src_size_f, uint src_size_w, uint src_size_z, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_z, uint dst_size_y, uint dst_size_x)
{
const uint src_pitch_x = 1;
const uint src_pitch_y = src_pitch_x * src_size_x;
const uint src_pitch_z = src_pitch_y * src_size_y;
const uint src_pitch_w = src_pitch_z * src_size_z;
const uint src_pitch_f = src_pitch_w * src_size_w;
const uint src_pitch_b = src_pitch_f * src_size_f;
uint flat_idx = x * src_pitch_x + y * src_pitch_y + z * src_pitch_z + w * src_pitch_w + i * src_pitch_f + o * src_pitch_b;
uint dst_x = flat_idx % dst_size_x;
flat_idx /= dst_size_x;
uint dst_y = flat_idx % dst_size_y;
flat_idx /= dst_size_y;
uint dst_z = flat_idx % dst_size_z;
flat_idx /= dst_size_z;
uint dst_f = flat_idx % dst_size_f;
flat_idx /= dst_size_f;
uint dst_b = flat_idx;
return (uint8)(0, dst_b, dst_f, 0, dst_z, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_grouped)(uint g, uint o, uint i, uint z, uint y, uint x, uint src_size_ofm, uint dst_size_ofm)
{
const uint flat_ofm = g * src_size_ofm + o;
const uint dst_ofm = flat_ofm % dst_size_ofm;
const uint dst_g = flat_ofm / dst_size_ofm;
const uint dst_ifm = i;
const uint dst_z = z;
const uint dst_y = y;
const uint dst_x = x;
return (uint8)(dst_g, dst_ofm, dst_ifm, 0, dst_z, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_dims)(
uint o, uint i, uint w, uint z, uint y, uint x,
uint src_size_f, uint src_size_w, uint src_size_z, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_w, uint dst_size_z, uint dst_size_y, uint dst_size_x,
uint b, uint f, uint v, uint u, uint w, uint z, uint y, uint x,
uint src_size_f, uint src_size_v, uint src_size_u, uint src_size_w, uint src_size_z, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_v, uint dst_size_u, uint dst_size_w, uint dst_size_z, uint dst_size_y, uint dst_size_x,
uint src_dims, uint dst_dims)
{
if (src_dims == 4 && dst_dims == 2)
{
return FUNC_CALL(reshape_4_to_2)(o,i,y,x,src_size_y,src_size_x);
}
else if (src_dims == 2 && dst_dims == 4)
{
return FUNC_CALL(reshape_2_to_4)(o,i,y,x,dst_size_y,dst_size_x);
}
else if (src_dims == 4 && dst_dims == 6)
{
return FUNC_CALL(reshape_4_to_6)(o, i, y, x, src_size_f, src_size_y, src_size_x, dst_size_f, dst_size_w, dst_size_z, dst_size_y, dst_size_x);
}
else if (src_dims == 6 && dst_dims == 4)
{
return FUNC_CALL(reshape_6_to_4)(o, i, w, z, y, x, src_size_f, src_size_w, src_size_z, src_size_y, src_size_x, dst_size_f, dst_size_y, dst_size_x);
}
else if (src_dims == 4 && dst_dims == 5)
{
return FUNC_CALL(reshape_4_to_5)(o, i, y, x, src_size_f, src_size_y, src_size_x, dst_size_f, dst_size_z, dst_size_y, dst_size_x);
}
else if (src_dims == 5 && dst_dims == 4)
{
return FUNC_CALL(reshape_5_to_4)(o, i, z, y, x, src_size_f, src_size_z, src_size_y, src_size_x, dst_size_f, dst_size_y, dst_size_x);
}
else if (src_dims == 6 && dst_dims == 5)
{
return FUNC_CALL(reshape_6_to_5)(o, i, w, z, y, x, src_size_f, src_size_w, src_size_z, src_size_y, src_size_x, dst_size_f, dst_size_z, dst_size_y, dst_size_x);
}
return (uint8)(0, o, i, w, z, y, x, 0);
}
inline uint8 FUNC(reshape_dims_with_groups)(
uint g, uint o, uint i, uint w, uint z, uint y, uint x,
uint src_size_ofm, uint src_size_ifm, uint src_size_w, uint src_size_z, uint src_size_y, uint src_size_x,
uint dst_size_ofm, uint dst_size_ifm, uint dst_size_w, uint dst_size_z, uint dst_size_y, uint dst_size_x,
uint src_dims, uint dst_dims, uint src_size_groups, uint dst_size_groups)
{
if (src_dims == 5 && dst_dims == 4) // goiyx -> oiyx
{
return FUNC_CALL(reshape_grouped)(g, o, i, 0, y, x, src_size_ofm, dst_size_ofm);
}
else if (src_dims == 6 && dst_dims == 5) // goizyx -> oizyx or goizyx -> goiyx
{
return FUNC_CALL(reshape_grouped)(g, o, i, z, y, x, src_size_ofm, dst_size_ofm);
}
else if (src_dims == 6 && dst_dims == 4) // goizyx -> oiyx
{
return FUNC_CALL(reshape_grouped)(g, o, i, 0, y, x, src_size_ofm, dst_size_ofm);
}
return (uint8)(g, o, i, w, z, y, x, 0);
if (dst_dims == src_dims) {
return (uint8)(b, f, v, u, w, z, y, x);
}
const uint src_pitch_x = 1;
const uint src_pitch_y = src_pitch_x * src_size_x;
const uint src_pitch_z = src_pitch_y * src_size_y;
const uint src_pitch_w = src_pitch_z * src_size_z;
const uint src_pitch_u = src_pitch_w * src_size_w;
const uint src_pitch_v = src_pitch_u * src_size_u;
const uint src_pitch_f = src_pitch_v * src_size_v;
const uint src_pitch_b = src_pitch_f * src_size_f;
uint flat_idx = x * src_pitch_x
+ y * src_pitch_y
+ z * src_pitch_z
+ w * src_pitch_w
+ u * src_pitch_u
+ v * src_pitch_v
+ f * src_pitch_f
+ b * src_pitch_b;
uint dst_x = flat_idx % dst_size_x;
flat_idx /= dst_size_x;
uint dst_y = flat_idx % dst_size_y;
flat_idx /= dst_size_y;
uint dst_z = flat_idx % dst_size_z;
flat_idx /= dst_size_z;
uint dst_w = flat_idx % dst_size_w;
flat_idx /= dst_size_w;
uint dst_u = flat_idx % dst_size_u;
flat_idx /= dst_size_u;
uint dst_v = flat_idx % dst_size_v;
flat_idx /= dst_size_v;
uint dst_f = flat_idx % dst_size_f;
flat_idx /= dst_size_f;
uint dst_b = flat_idx;
return (uint8)(dst_b, dst_f, dst_v, dst_u, dst_w, dst_z, dst_y, dst_x);
}
// ======================================================================
// Reshape dimensions from indices in `src_prefix` DataTensor to indices
// in `dst_prefix` DataTensor.
// Reshaping assumes dimension order: b, f, reverse spatial - w, z, y, x.
#define RESHAPE_DIMS(src_prefix, dst_prefix, o, i, w, z, y, x) \
#define RESHAPE_DIMS(src_prefix, dst_prefix, b, f, v, u, w, z, y, x) \
FUNC_CALL(reshape_dims)( \
o, i, w, z, y, x, \
b, f, v, u, w, z, y, x, \
CAT(src_prefix, _FEATURE_NUM), \
CAT(src_prefix, _SIZE_V), \
CAT(src_prefix, _SIZE_U), \
CAT(src_prefix, _SIZE_W), \
CAT(src_prefix, _SIZE_Z), \
CAT(src_prefix, _SIZE_Y), \
CAT(src_prefix, _SIZE_X), \
CAT(dst_prefix, _FEATURE_NUM), \
CAT(dst_prefix, _SIZE_V), \
CAT(dst_prefix, _SIZE_U), \
CAT(dst_prefix, _SIZE_W), \
CAT(dst_prefix, _SIZE_Z), \
CAT(dst_prefix, _SIZE_Y), \
CAT(dst_prefix, _SIZE_X), \
CAT(src_prefix, _DIMS), \
CAT(dst_prefix, _DIMS))
// ======================================================================
// Reshape indices between WeightTensors.
// Weight tensor has _IFM in place of _FEATURE_NUM and does not support 6D,
// so macro for DataTensors cannot be used directly.
#define RESHAPE_WEIGHT_DIMS(src_prefix, dst_prefix, o, i, w, z, y, x) \
FUNC_CALL(reshape_dims)( \
o, i, w, z, y, x, \
CAT(src_prefix, _IFM_NUM), \
1, \
CAT(src_prefix, _SIZE_Z), \
CAT(src_prefix, _SIZE_Y), \
CAT(src_prefix, _SIZE_X), \
CAT(dst_prefix, _IFM_NUM), \
1, \
CAT(dst_prefix, _SIZE_Z), \
CAT(dst_prefix, _SIZE_Y), \
CAT(dst_prefix, _SIZE_X), \
CAT(src_prefix, _DIMS), \
CAT(dst_prefix, _DIMS))
#define RESHAPE_WEIGHT_DIMS_WITH_GROUPS(src_prefix, dst_prefix, g, o, i, w, z, y, x)\
FUNC_CALL(reshape_dims_with_groups)( \
g, o, i, w, z, y, x, \
CAT(src_prefix, _OFM_NUM), \
CAT(src_prefix, _IFM_NUM), \
1, \
CAT(src_prefix, _SIZE_Z), \
CAT(src_prefix, _SIZE_Y), \
CAT(src_prefix, _SIZE_X), \
CAT(dst_prefix, _OFM_NUM), \
CAT(dst_prefix, _IFM_NUM), \
1, \
CAT(dst_prefix, _SIZE_Z), \
CAT(dst_prefix, _SIZE_Y), \
CAT(dst_prefix, _SIZE_X), \
CAT(src_prefix, _DIMS), \
CAT(dst_prefix, _DIMS), \
CAT(src_prefix, _GROUPS_NUM), \
CAT(dst_prefix, _GROUPS_NUM))
inline uint8 FUNC(reshape_dims3d)(uint o, uint i, uint z, uint y, uint x, uint src_size_z, uint src_size_y, uint src_size_x, uint dst_size_z, uint dst_size_y, uint dst_size_x, uint src_dims, uint dst_dims)
{
if (src_dims == 4 && dst_dims == 5)
{
return (uint8)(0,o,i,1,y,x,0,0);
}
else if (src_dims == 5 && dst_dims == 4)
{
uint _y = z*src_size_y + y;
return (uint8)(0,o,i,0,_y,x,0,0);
}
return (uint8)(0,o,i,z,y,x,0,0);
}
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/permute_ref.cl
+25 -3
View File
@@ -14,7 +14,7 @@ KERNEL (permute_ref)(
)
{
#ifdef F_FIRST
//gws(f, x * y, z * w * b)
//gws(f, x * y, z * w * u * v * b)
const uint gid_0 = get_global_id(0);
const uint gid_1 = get_global_id(1);
const uint gid_2 = get_global_id(2);
@@ -26,10 +26,21 @@ KERNEL (permute_ref)(
#elif INPUT0_DIMS == 5
const uint b = gid_2 / INPUT0_SIZE_Z;
const uint z = gid_2 % INPUT0_SIZE_Z;
#else
#elif INPUT0_DIMS == 6
const uint b = gid_2 / (INPUT0_SIZE_W * INPUT0_SIZE_Z) % INPUT0_BATCH_NUM;
const uint z = gid_2 / INPUT0_SIZE_W % INPUT0_SIZE_Z;
const uint w = gid_2 % INPUT0_SIZE_W;
#elif INPUT0_DIMS == 7
const uint b = gid_2 / (INPUT0_SIZE_U * INPUT0_SIZE_W * INPUT0_SIZE_Z) % INPUT0_BATCH_NUM;
const uint z = gid_2 / (INPUT0_SIZE_U * INPUT0_SIZE_W) % INPUT0_SIZE_Z;
const uint w = gid_2 / INPUT0_SIZE_U % INPUT0_SIZE_W;
const uint u = gid_2 % INPUT0_SIZE_U;
#elif INPUT0_DIMS == 8
const uint b = gid_2 / (INPUT0_SIZE_V * INPUT0_SIZE_U * INPUT0_SIZE_W * INPUT0_SIZE_Z) % INPUT0_BATCH_NUM;
const uint z = gid_2 / (INPUT0_SIZE_V * INPUT0_SIZE_U * INPUT0_SIZE_W) % INPUT0_SIZE_Z;
const uint w = gid_2 / (INPUT0_SIZE_V * INPUT0_SIZE_U) % INPUT0_SIZE_W;
const uint u = gid_2 / INPUT0_SIZE_V % INPUT0_SIZE_U;
const uint v = gid_2 % INPUT0_SIZE_V;
#endif
#else
//gws(x, y * z * w, b*f)
@@ -39,7 +50,18 @@ KERNEL (permute_ref)(
#elif INPUT0_DIMS == 5
const uint z = gid_0 / INPUT0_SIZE_Y;
const uint y = gid_0 % INPUT0_SIZE_Y;
#else
#elif INPUT0_DIMS == 6
const uint w = gid_0 / (INPUT0_SIZE_Y * INPUT0_SIZE_Z) % INPUT0_SIZE_W;
const uint z = gid_0 / INPUT0_SIZE_Y % INPUT0_SIZE_Z;
const uint y = gid_0 % INPUT0_SIZE_Y;
#elif INPUT0_DIMS == 7
const uint u = gid_0 / (INPUT0_SIZE_Y * INPUT0_SIZE_Z * INPUT0_SIZE_W) % INPUT0_SIZE_U;
const uint w = gid_0 / (INPUT0_SIZE_Y * INPUT0_SIZE_Z) % INPUT0_SIZE_W;
const uint z = gid_0 / INPUT0_SIZE_Y % INPUT0_SIZE_Z;
const uint y = gid_0 % INPUT0_SIZE_Y;
#elif INPUT0_DIMS == 8
const uint v = gid_0 / (INPUT0_SIZE_Y * INPUT0_SIZE_Z * INPUT0_SIZE_W * INPUT0_SIZE_U) % INPUT0_SIZE_V;
const uint u = gid_0 / (INPUT0_SIZE_Y * INPUT0_SIZE_Z * INPUT0_SIZE_W) % INPUT0_SIZE_U;
const uint w = gid_0 / (INPUT0_SIZE_Y * INPUT0_SIZE_Z) % INPUT0_SIZE_W;
const uint z = gid_0 / INPUT0_SIZE_Y % INPUT0_SIZE_Z;
const uint y = gid_0 % INPUT0_SIZE_Y;
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reduce_ref.cl
+114 -58
View File
@@ -4,10 +4,12 @@
#include "include/batch_headers/fetch_data.cl"
inline uint FUNC(calc_linear_offset)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint w, uint z, uint y, uint x)
inline uint FUNC(calc_linear_offset)(OPTIONAL_SHAPE_INFO_ARG uint b, uint f, uint v, uint u, uint w, uint z, uint y, uint x)
{
uint index = b * OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z * OUTPUT_SIZE_W * OUTPUT_FEATURE_NUM +
f * OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z * OUTPUT_SIZE_W +
uint index = b * OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z * OUTPUT_SIZE_W * OUTPUT_SIZE_U * OUTPUT_SIZE_V * OUTPUT_FEATURE_NUM +
f * OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z * OUTPUT_SIZE_W * OUTPUT_SIZE_U * OUTPUT_SIZE_V +
v * OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z * OUTPUT_SIZE_W * OUTPUT_SIZE_U +
u * OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z * OUTPUT_SIZE_W +
w * OUTPUT_SIZE_X * OUTPUT_SIZE_Y * OUTPUT_SIZE_Z +
z * OUTPUT_SIZE_X * OUTPUT_SIZE_Y +
y * OUTPUT_SIZE_X +
@@ -25,23 +27,39 @@ KERNEL(reduce_ref)(
#endif
)
{
const uint xy = (uint)get_global_id(0);
const uint wz = (uint)get_global_id(1);
const uint bf = (uint)get_global_id(2);
const uint x = xy % OUTPUT_SIZE_X;
const uint y = xy / OUTPUT_SIZE_X;
const uint xy = (uint)get_global_id(0);
const uint wzuv = (uint)get_global_id(1);
const uint bf = (uint)get_global_id(2);
const uint x = xy % OUTPUT_SIZE_X;
const uint y = xy / OUTPUT_SIZE_X;
const uint b = bf / OUTPUT_FEATURE_NUM;
const uint f = bf % OUTPUT_FEATURE_NUM;
#if INPUT0_DIMS == 4
const uint w = 0;
const uint z = 0;
const uint u = 0;
const uint v = 0;
#elif INPUT0_DIMS == 5
const uint z = wz % OUTPUT_SIZE_Z;
const uint z = wzuv % OUTPUT_SIZE_Z;
const uint w = 0;
const uint u = 0;
const uint v = 0;
#elif INPUT0_DIMS == 6
const uint z = wz % OUTPUT_SIZE_Z;
const uint w = wz / OUTPUT_SIZE_Z;
const uint z = wzuv % OUTPUT_SIZE_Z;
const uint w = wzuv / OUTPUT_SIZE_Z;
const uint u = 0;
const uint v = 0;
#elif INPUT0_DIMS == 7
const uint z = wzuv % OUTPUT_SIZE_Z;
const uint w = wzuv / OUTPUT_SIZE_Z % OUTPUT_SIZE_W;
const uint u = wzuv / (OUTPUT_SIZE_Z * OUTPUT_SIZE_W) % OUTPUT_SIZE_U;
const uint v = 0;
#elif INPUT0_DIMS == 8
const uint z = wzuv % OUTPUT_SIZE_Z;
const uint w = wzuv / OUTPUT_SIZE_Z % OUTPUT_SIZE_W;
const uint u = wzuv / (OUTPUT_SIZE_Z * OUTPUT_SIZE_W) % OUTPUT_SIZE_U;
const uint v = wzuv / (OUTPUT_SIZE_Z * OUTPUT_SIZE_W * OUTPUT_SIZE_U);
#endif
#if OUTPUT_DIMS == 4
@@ -50,9 +68,13 @@ KERNEL(reduce_ref)(
const uint out_idx = OUTPUT_GET_INDEX(b, f, z, y, x);
#elif OUTPUT_DIMS == 6
const uint out_idx = OUTPUT_GET_INDEX(b, f, w, z, y, x);
#elif OUTPUT_DIMS == 7
const uint out_idx = OUTPUT_GET_INDEX(b, f, u, w, z, y, x);
#elif OUTPUT_DIMS == 8
const uint out_idx = OUTPUT_GET_INDEX(b, f, v, u, w, z, y, x);
#endif
const uint linear_idx = FUNC_CALL(calc_linear_offset)(OPTIONAL_SHAPE_INFO_TENSOR b, f, w, z, y, x);
const uint linear_idx = FUNC_CALL(calc_linear_offset)(OPTIONAL_SHAPE_INFO_TENSOR b, f, v, u, w, z, y, x);
if (linear_idx >= COMPUTATIONAL_OPERATIONS_NUMBER)
return;
@@ -72,7 +94,33 @@ KERNEL(reduce_ref)(
const uint feature_max_val = feature_out + 1;
#endif
#if INPUT0_DIMS == 6
#if INPUT0_DIMS >= 8
#ifdef REDUCE_V
const uint v_out = 0;
const uint v_max_val = INPUT0_SIZE_V;
#else
const uint v_out = SIZE_V_IDX_COMP(linear_idx);
const uint v_max_val = v_out + 1;
#endif
#else
const uint v_out = 0;
const uint v_max_val = 1;
#endif
#if INPUT0_DIMS >= 7
#ifdef REDUCE_U
const uint u_out = 0;
const uint u_max_val = INPUT0_SIZE_U;
#else
const uint u_out = SIZE_U_IDX_COMP(linear_idx);
const uint u_max_val = u_out + 1;
#endif
#else
const uint u_out = 0;
const uint u_max_val = 1;
#endif
#if INPUT0_DIMS >= 6
#ifdef REDUCE_W
const uint w_out = 0;
const uint w_max_val = INPUT0_SIZE_W;
@@ -85,7 +133,7 @@ KERNEL(reduce_ref)(
const uint w_max_val = 1;
#endif
#if INPUT0_DIMS == 6 || INPUT0_DIMS == 5
#if INPUT0_DIMS >= 5
#ifdef REDUCE_Z
const uint z_out = 0;
const uint z_max_val = INPUT0_SIZE_Z;
@@ -117,67 +165,75 @@ KERNEL(reduce_ref)(
uint counter = 0;
for (uint bi = batch_out; bi < batch_max_val; ++bi) {
for (uint fi = feature_out; fi < feature_max_val; ++fi) {
for (uint wi = w_out; wi < w_max_val; ++wi) {
for (uint zi = z_out; zi < z_max_val; ++zi) {
for (uint yi = y_out; yi < y_max_val; ++yi) {
for (uint xi = x_out; xi < x_max_val; ++xi) {
#if INPUT0_DIMS == 6
const uint input_idx = INPUT0_GET_INDEX(bi, fi, wi, zi, yi, xi);
for (uint vi = v_out; vi < v_max_val; ++vi) {
for (uint ui = u_out; ui < u_max_val; ++ui) {
for (uint wi = w_out; wi < w_max_val; ++wi) {
for (uint zi = z_out; zi < z_max_val; ++zi) {
for (uint yi = y_out; yi < y_max_val; ++yi) {
for (uint xi = x_out; xi < x_max_val; ++xi) {
#if INPUT0_DIMS == 8
const uint input_idx = INPUT0_GET_INDEX(bi, fi, vi, ui, wi, zi, yi, xi);
#elif INPUT0_DIMS == 7
const uint input_idx = INPUT0_GET_INDEX(bi, fi, ui, wi, zi, yi, xi);
#elif INPUT0_DIMS == 6
const uint input_idx = INPUT0_GET_INDEX(bi, fi, wi, zi, yi, xi);
#elif INPUT0_DIMS == 5
const uint input_idx = INPUT0_GET_INDEX(bi, fi, zi, yi, xi);
const uint input_idx = INPUT0_GET_INDEX(bi, fi, zi, yi, xi);
#else
const uint input_idx = INPUT0_GET_INDEX(bi, fi, yi, xi);
const uint input_idx = INPUT0_GET_INDEX(bi, fi, yi, xi);
#endif
#ifdef REDUCE_SUM_MODE
acc += data[input_idx];
acc += data[input_idx];
#elif REDUCE_MAX_MODE
if (counter == 0)
acc = data[input_idx];
else
acc = data[input_idx] > acc ? data[input_idx] : acc;
if (counter == 0)
acc = data[input_idx];
else
acc = data[input_idx] > acc ? data[input_idx] : acc;
#elif REDUCE_MIN_MODE
if (counter == 0)
acc = data[input_idx];
else
acc = data[input_idx] < acc ? data[input_idx] : acc;
if (counter == 0)
acc = data[input_idx];
else
acc = data[input_idx] < acc ? data[input_idx] : acc;
#elif REDUCE_MEAN_MODE
acc += data[input_idx];
acc += data[input_idx];
#elif REDUCE_PROD_MODE
if (counter == 0)
acc = data[input_idx];
else
acc *= data[input_idx];
if (counter == 0)
acc = data[input_idx];
else
acc *= data[input_idx];
#elif REDUCE_AND_MODE
if (counter == 0)
acc = data[input_idx];
else
acc = acc && data[input_idx];
if (counter == 0)
acc = data[input_idx];
else
acc = acc && data[input_idx];
#elif REDUCE_OR_MODE
if (counter == 0)
acc = data[input_idx];
else
acc = acc || data[input_idx];
if (counter == 0)
acc = data[input_idx];
else
acc = acc || data[input_idx];
#elif REDUCE_SUM_SQUARE_MODE
acc += data[input_idx] * data[input_idx];
acc += data[input_idx] * data[input_idx];
#elif REDUCE_L1_MODE
#if !INPUT0_IS_FP
acc += TO_ACCUMULATOR_TYPE(fabs(TO_FINAL_ACCUMULATOR_TYPE(data[input_idx])));
#else
acc += fabs(data[input_idx]);
#endif
#if !INPUT0_IS_FP
acc += TO_ACCUMULATOR_TYPE(fabs(TO_FINAL_ACCUMULATOR_TYPE(data[input_idx])));
#else
acc += fabs(data[input_idx]);
#endif
#elif REDUCE_L2_MODE
acc += data[input_idx] * data[input_idx];
acc += data[input_idx] * data[input_idx];
#elif REDUCE_LOG_SUM_MODE
acc += data[input_idx];
acc += data[input_idx];
#elif REDUCE_LOG_SUM_EXP_MODE
#if !INPUT0_IS_FP
acc += TO_ACCUMULATOR_TYPE(exp(TO_FINAL_ACCUMULATOR_TYPE(data[input_idx])));
#else
acc += exp(data[input_idx]);
#endif
#if !INPUT0_IS_FP
acc += TO_ACCUMULATOR_TYPE(exp(TO_FINAL_ACCUMULATOR_TYPE(data[input_idx])));
#else
acc += exp(data[input_idx]);
#endif
#endif
counter++;
counter++;
}
}
}
}
}
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_biplanar_nv12.cl
+12 -12
View File
@@ -63,26 +63,26 @@ KERNEL(reorder_biplanar_nv12)(
G -= VALUE_TO_SUBTRACT[1];
B -= VALUE_TO_SUBTRACT[2];
#elif defined MEAN_SUBTRACT_IN_BUFFER
uint8 msv = RESHAPE_DIMS(INPUT0, MEAN_SUBTRACT, b, 0, w, z, y, x);
R -= mean_subtract[GET_DATA_INDEX_SAFE(MEAN_SUBTRACT, msv.s1, msv.s2, msv.s5, msv.s6)];
uint8 msv = RESHAPE_DIMS(INPUT0, MEAN_SUBTRACT, b, 0, 0, 0, w, z, y, x);
R -= mean_subtract[GET_DATA_INDEX_SAFE(MEAN_SUBTRACT, msv.s0, msv.s1, msv.s6, msv.s7)];
msv = RESHAPE_DIMS(INPUT0, MEAN_SUBTRACT, b, 1, w, z, y, x);
G -= mean_subtract[GET_DATA_INDEX_SAFE(MEAN_SUBTRACT, msv.s1, msv.s2, msv.s5, msv.s6)];
msv = RESHAPE_DIMS(INPUT0, MEAN_SUBTRACT, b, 1, 0, 0, w, z, y, x);
G -= mean_subtract[GET_DATA_INDEX_SAFE(MEAN_SUBTRACT, msv.s0, msv.s1, msv.s6, msv.s7)];
msv = RESHAPE_DIMS(INPUT0, MEAN_SUBTRACT, b, 2, w, z, y, x);
B -= mean_subtract[GET_DATA_INDEX_SAFE(MEAN_SUBTRACT, msv.s1, msv.s2, msv.s5, msv.s6)];
msv = RESHAPE_DIMS(INPUT0, MEAN_SUBTRACT, b, 2, 0, 0, w, z, y, x);
B -= mean_subtract[GET_DATA_INDEX_SAFE(MEAN_SUBTRACT, msv.s0, msv.s1, msv.s6, msv.s7)];
#endif
uint8 ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 0, w, z, y, x);
uint output_idx = FUNC_CALL(get_output_index)(ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6);
uint8 ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 0, 0, 0, w, z, y, x);
uint output_idx = FUNC_CALL(get_output_index)(ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
output[output_idx] = ACTIVATION_FUNC_TYPED(OUTPUT_REORDER, TO_OUTPUT_REORDER_TYPE(R), NL_M, NL_N);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 1, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 1, 0, 0, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
output[output_idx] = ACTIVATION_FUNC_TYPED(OUTPUT_REORDER, TO_OUTPUT_REORDER_TYPE(G), NL_M, NL_N);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 2, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 2, 0, 0, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
output[output_idx] = ACTIVATION_FUNC_TYPED(OUTPUT_REORDER, TO_OUTPUT_REORDER_TYPE(B), NL_M, NL_N);
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data.cl
+48 -24
View File
@@ -34,11 +34,15 @@ KERNEL (reorder_data)(
const uint x = 0;
const uint z = 0;
const uint w = 0;
const uint u = 0;
const uint v = 0;
#elif INPUT0_DIMS == 4
const uint y = ((uint)(get_global_id(GWS_YX))) / INPUT0_SIZE_X;
const uint x = ((uint)(get_global_id(GWS_YX))) % INPUT0_SIZE_X;
const uint z = 0;
const uint w = 0;
const uint u = 0;
const uint v = 0;
#elif INPUT0_DIMS == 5
uint data_idx = get_global_id(GWS_YX);
uint tmp_data_idx = data_idx / INPUT0_SIZE_X;
@@ -52,12 +56,32 @@ KERNEL (reorder_data)(
tmp_data_idx = data_idx / INPUT0_SIZE_Z;
const uint z = data_idx - tmp_data_idx * INPUT0_SIZE_Z;
const uint w = 0;
const uint u = 0;
const uint v = 0;
#elif INPUT0_DIMS == 6
const uint gid_yx = (uint)(get_global_id(GWS_YX));
const uint x = gid_yx % INPUT0_SIZE_X;
const uint y = gid_yx / INPUT0_SIZE_X % INPUT0_SIZE_Y;
const uint z = gid_yx / INPUT0_SIZE_X / INPUT0_SIZE_Y % INPUT0_SIZE_Z;
const uint w = gid_yx / INPUT0_SIZE_X / INPUT0_SIZE_Y / INPUT0_SIZE_Z % INPUT0_SIZE_W;
const uint u = 0;
const uint v = 0;
#elif INPUT0_DIMS == 7
const uint gid_yx = (uint)(get_global_id(GWS_YX));
const uint x = gid_yx % INPUT0_SIZE_X;
const uint y = gid_yx / INPUT0_SIZE_X % INPUT0_SIZE_Y;
const uint z = gid_yx / INPUT0_SIZE_X / INPUT0_SIZE_Y % INPUT0_SIZE_Z;
const uint w = gid_yx / INPUT0_SIZE_X / INPUT0_SIZE_Y / INPUT0_SIZE_Z % INPUT0_SIZE_W;
const uint u = gid_yx / INPUT0_SIZE_X / INPUT0_SIZE_Y / INPUT0_SIZE_Z / INPUT0_SIZE_W % INPUT0_SIZE_U;
const uint v = 0;
#elif INPUT0_DIMS == 8
const uint gid_yx = (uint)(get_global_id(GWS_YX));
const uint x = gid_yx % INPUT0_SIZE_X;
const uint y = gid_yx / INPUT0_SIZE_X % INPUT0_SIZE_Y;
const uint z = gid_yx / INPUT0_SIZE_X / INPUT0_SIZE_Y % INPUT0_SIZE_Z;
const uint w = gid_yx / INPUT0_SIZE_X / INPUT0_SIZE_Y / INPUT0_SIZE_Z % INPUT0_SIZE_W;
const uint u = gid_yx / INPUT0_SIZE_X / INPUT0_SIZE_Y / INPUT0_SIZE_Z / INPUT0_SIZE_W % INPUT0_SIZE_U;
const uint v = gid_yx / INPUT0_SIZE_X / INPUT0_SIZE_Y / INPUT0_SIZE_Z / INPUT0_SIZE_W / INPUT0_SIZE_U % INPUT0_SIZE_V;
#endif
#if defined INPUT0_LAYOUT_NV12 && !SURFACE_INPUT
@@ -75,20 +99,20 @@ KERNEL (reorder_data)(
const sampler_t imageSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST | CLK_ADDRESS_CLAMP;
OUTPUT_TYPE4 colorRGBA = IMAGE_READ(input, (int2)(x, y));
#elif defined OUTPUT_LAYOUT_IMAGE_2D_RGBA
uint8 ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, f, w, z, y, x);
const uint input_idx_R = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR b, 0, w, z, y, x);
const uint input_idx_G = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR b, 1, w, z, y, x);
const uint input_idx_B = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR b, 2, w, z, y, x);
uint8 ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, f, v, u, w, z, y, x);
const uint input_idx_R = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR b, 0, v, u, w, z, y, x);
const uint input_idx_G = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR b, 1, v, u, w, z, y, x);
const uint input_idx_B = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR b, 2, v, u, w, z, y, x);
#if OUTPUT_FEATURE_NUM == 3
INPUT_TYPE4 colorRGBA = { TO_INPUT_REORDER_TYPE(input[input_idx_R]), TO_INPUT_REORDER_TYPE(input[input_idx_G]), TO_INPUT_REORDER_TYPE(input[input_idx_B]), TO_INPUT_REORDER_TYPE(0.f) };
#else
const uint input_idx_A = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR b, 3, w, z, y, x);
const uint input_idx_A = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR b, 3, v, u, w, z, y, x);
INPUT_TYPE4 colorRGBA = { TO_INPUT_REORDER_TYPE(input[input_idx_R]), TO_INPUT_REORDER_TYPE(input[input_idx_G]), TO_INPUT_REORDER_TYPE(input[input_idx_B]), TO_INPUT_REORDER_TYPE(input[input_idx_A]) };
#endif
#else
uint8 ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, f, w, z, y, x);
const uint input_idx = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR b, f, w, z, y, x);
const uint output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s1,ov.s2,ov.s3,ov.s4,ov.s5,ov.s6);
uint8 ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, f, v, u, w, z, y, x);
const uint input_idx = FUNC_CALL(get_input_index)(OPTIONAL_SHAPE_INFO_TENSOR b, f, v, u, w, z, y, x);
const uint output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
#if defined MEAN_SUBTRACT_INSIDE_PARAMS
float res = TO_MEAN_TYPE(input[input_idx]);
@@ -100,8 +124,8 @@ KERNEL (reorder_data)(
#else
// TODO Add support for 6D mean
MEAN_SUBTRACT_TYPE res = TO_MEAN_TYPE(input[input_idx]);
uint8 msv = RESHAPE_DIMS(INPUT0, MEAN_SUBTRACT, b, f, w, z, y, x);
res = MEAN_OP(res, mean_subtract[GET_DATA_INDEX_SAFE(MEAN_SUBTRACT, msv.s1, msv.s2, /*msv.s3, msv.s4,*/ msv.s5, msv.s6)]);
uint8 msv = RESHAPE_DIMS(INPUT0, MEAN_SUBTRACT, b, f, v, u, w, z, y, x);
res = MEAN_OP(res, mean_subtract[GET_DATA_INDEX_SAFE(MEAN_SUBTRACT, msv.s0, msv.s1, /*msv.s2, msv.s3, msv.s4,msv.s5,*/ msv.s6, msv.s7)]);
#endif
#elif SURFACE_INPUT
float4 Y = read_imagef(input, (int2)(x, y));
@@ -113,28 +137,28 @@ KERNEL (reorder_data)(
#endif
#if defined INPUT0_LAYOUT_NV12 && !SURFACE_INPUT
uint8 ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 0, w, z, y, x);
uint output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6);
uint8 ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 0, v, u, w, z, y, x);
uint output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
output[output_idx] = ACTIVATION_FUNC_TYPED(OUTPUT_REORDER, TO_OUTPUT_REORDER_TYPE(R), NL_M, NL_N);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 1, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 1, v, u, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
output[output_idx] = ACTIVATION_FUNC_TYPED(OUTPUT_REORDER, TO_OUTPUT_REORDER_TYPE(G), NL_M, NL_N);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 2, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 2, v, u, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
output[output_idx] = ACTIVATION_FUNC_TYPED(OUTPUT_REORDER, TO_OUTPUT_REORDER_TYPE(B), NL_M, NL_N);
#elif INPUT0_LAYOUT_IMAGE_2D_RGBA
uint8 ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 0, w, z, y, x);
uint output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6);
uint8 ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 0, v, u, w, z, y, x);
uint output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
output[output_idx] = ACTIVATION_FUNC_TYPED(OUTPUT_REORDER, TO_OUTPUT_REORDER_TYPE(colorRGBA.s0), NL_M, NL_N);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 1, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 1, v, u, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
output[output_idx] = ACTIVATION_FUNC_TYPED(OUTPUT_REORDER, TO_OUTPUT_REORDER_TYPE(colorRGBA.s1), NL_M, NL_N);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 2, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 2, v, u, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
output[output_idx] = ACTIVATION_FUNC_TYPED(OUTPUT_REORDER, TO_OUTPUT_REORDER_TYPE(colorRGBA.s2), NL_M, NL_N);
#if INPUT0_FEATURE_NUM == 4
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 3, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6);
ov = RESHAPE_DIMS(INPUT0, OUTPUT, b, 3, v, u, w, z, y, x);
output_idx = FUNC_CALL(get_output_index)(OPTIONAL_SHAPE_INFO_TENSOR ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
output[output_idx] = ACTIVATION_FUNC_TYPED(OUTPUT_REORDER, TO_OUTPUT_REORDER_TYPE(colorRGBA.s3), NL_M, NL_N);
#endif
#elif OUTPUT_LAYOUT_IMAGE_2D_RGBA
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data_fast_b1.cl
+5 -5
View File
@@ -184,9 +184,9 @@ KERNEL (reorder_data_fast_b1)(
const uint input_idx = data_idx;
const uint output_idx = data_idx;
#else
uint8 ov = RESHAPE_DIMS(OUTPUT, INPUT0, b, f, w, z, y, x);
const uint input_idx = FUNC_CALL(get_input_index)(ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6);
const uint output_idx = FUNC_CALL(get_output_index)(b, f, w, z, y, x);
uint8 ov = RESHAPE_DIMS(OUTPUT, INPUT0, b, f, 0, 0, w, z, y, x);
const uint input_idx = FUNC_CALL(get_input_index)(ov.s0, ov.s1, ov.s2, ov.s3, ov.s4, ov.s5, ov.s6, ov.s7);
const uint output_idx = FUNC_CALL(get_output_index)(b, f, 0, 0, w, z, y, x);
#endif
#if defined MEAN_SUBTRACT_INSIDE_PARAMS
@@ -194,8 +194,8 @@ KERNEL (reorder_data_fast_b1)(
res -= VALUE_TO_SUBTRACT[f % VALUE_TO_SUBTRACT_SIZE];
#elif defined MEAN_SUBTRACT_IN_BUFFER
MEAN_SUBTRACT_TYPE res = TO_MEAN_TYPE(input[input_idx]);
uint8 msv = RESHAPE_DIMS(INPUT0, MEAN_SUBTRACT, b, f, w, z, y, x);
res -= mean_subtract[GET_DATA_INDEX_SAFE(MEAN_SUBTRACT, msv.s1, msv.s2, msv.s5, msv.s6)];
uint8 msv = RESHAPE_DIMS(INPUT0, MEAN_SUBTRACT, b, f, 0, 0, w, z, y, x);
res -= mean_subtract[GET_DATA_INDEX_SAFE(MEAN_SUBTRACT, msv.s0, msv.s1, msv.s6, msv.s7)];
#else
CALC_TYPE res = TO_CALC_TYPE(input[input_idx]);
#endif
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_weights.cl
+246 -1
View File
@@ -3,7 +3,6 @@
//
#include "include/batch_headers/fetch_weights.cl"
#include "include/reshape_dims.cl"
#include "include/image_data.cl"
#define DECLARE_SAMPLER const sampler_t imageSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST
@@ -16,6 +15,252 @@
#define IMAGE_WRITE(image, coord, val) write_imagef((image), (coord), (val))
#endif
inline uint8 FUNC(reshape_2_to_4)(uint o, uint i, uint y, uint x, uint dst_size_y, uint dst_size_x)
{
uint _i = i / (dst_size_y*dst_size_x);
uint _yx = i % (dst_size_y*dst_size_x);
uint _y = _yx / dst_size_x;
uint _x = _yx % dst_size_x;
return (uint8)(0, o, _i, 0, 0, _y,_x, 0);
}
inline uint8 FUNC(reshape_4_to_2)(uint o, uint i, uint y, uint x, uint src_size_y, uint src_size_x)
{
uint _i = i*src_size_y*src_size_x + y*src_size_x + x;
return (uint8)(0, o, _i, 0, 0, 0, 0, 0);
}
inline uint8 FUNC(reshape_4_to_5)(uint o, uint i, uint y, uint x,
uint src_size_f, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_z, uint dst_size_y, uint dst_size_x)
{
const uint src_pitch_x = 1;
const uint src_pitch_y = src_pitch_x * src_size_x;
const uint src_pitch_f = src_pitch_y * src_size_y;
const uint src_pitch_b = src_pitch_f * src_size_f;
uint flat_idx = x * src_pitch_x + y * src_pitch_y + i * src_pitch_f + o * src_pitch_b;
uint dst_x = flat_idx % dst_size_x;
flat_idx /= dst_size_x;
uint dst_y = flat_idx % dst_size_y;
flat_idx /= dst_size_y;
uint dst_z = flat_idx % dst_size_z;
flat_idx /= dst_size_z;
uint dst_f = flat_idx % dst_size_f;
flat_idx /= dst_size_f;
uint dst_b = flat_idx;
return (uint8)(0, dst_b, dst_f, 0, dst_z, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_5_to_4)(uint o, uint i, uint z, uint y, uint x,
uint src_size_f, uint src_size_z, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_y, uint dst_size_x)
{
const uint src_pitch_x = 1;
const uint src_pitch_y = src_pitch_x * src_size_x;
const uint src_pitch_z = src_pitch_y * src_size_y;
const uint src_pitch_f = src_pitch_z * src_size_z;
const uint src_pitch_b = src_pitch_f * src_size_f;
uint flat_idx = x * src_pitch_x + y * src_pitch_y + z * src_pitch_z + i * src_pitch_f + o * src_pitch_b;
uint dst_x = flat_idx % dst_size_x;
flat_idx /= dst_size_x;
uint dst_y = flat_idx % dst_size_y;
flat_idx /= dst_size_y;
uint dst_f = flat_idx % dst_size_f;
flat_idx /= dst_size_f;
uint dst_b = flat_idx;
return (uint8)(0, dst_b, dst_f, 0, 0, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_4_to_6)(uint o, uint i, uint y, uint x,
uint src_size_f, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_w, uint dst_size_z, uint dst_size_y, uint dst_size_x)
{
const uint src_pitch_x = 1;
const uint src_pitch_y = src_pitch_x * src_size_x;
const uint src_pitch_f = src_pitch_y * src_size_y;
const uint src_pitch_b = src_pitch_f * src_size_f;
uint flat_idx = x * src_pitch_x + y * src_pitch_y + i * src_pitch_f + o * src_pitch_b;
uint dst_x = flat_idx % dst_size_x;
flat_idx /= dst_size_x;
uint dst_y = flat_idx % dst_size_y;
flat_idx /= dst_size_y;
uint dst_z = flat_idx % dst_size_z;
flat_idx /= dst_size_z;
uint dst_w = flat_idx % dst_size_w;
flat_idx /= dst_size_w;
uint dst_f = flat_idx % dst_size_f;
flat_idx /= dst_size_f;
uint dst_b = flat_idx;
return (uint8)(0, dst_b, dst_f, dst_w, dst_z, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_6_to_4)(uint o, uint i, uint w, uint z, uint y, uint x,
uint src_size_f, uint src_size_w, uint src_size_z, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_y, uint dst_size_x)
{
const uint src_pitch_x = 1;
const uint src_pitch_y = src_pitch_x * src_size_x;
const uint src_pitch_z = src_pitch_y * src_size_y;
const uint src_pitch_w = src_pitch_z * src_size_z;
const uint src_pitch_f = src_pitch_w * src_size_w;
const uint src_pitch_b = src_pitch_f * src_size_f;
uint flat_idx = x * src_pitch_x + y * src_pitch_y + z * src_pitch_z + w * src_pitch_w + i * src_pitch_f + o * src_pitch_b;
uint dst_x = flat_idx % dst_size_x;
flat_idx /= dst_size_x;
uint dst_y = flat_idx % dst_size_y;
flat_idx /= dst_size_y;
uint dst_f = flat_idx % dst_size_f;
flat_idx /= dst_size_f;
uint dst_b = flat_idx;
return (uint8)(0, dst_b, dst_f, 0, 0, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_6_to_5)(uint o, uint i, uint w, uint z, uint y, uint x,
uint src_size_f, uint src_size_w, uint src_size_z, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_z, uint dst_size_y, uint dst_size_x)
{
const uint src_pitch_x = 1;
const uint src_pitch_y = src_pitch_x * src_size_x;
const uint src_pitch_z = src_pitch_y * src_size_y;
const uint src_pitch_w = src_pitch_z * src_size_z;
const uint src_pitch_f = src_pitch_w * src_size_w;
const uint src_pitch_b = src_pitch_f * src_size_f;
uint flat_idx = x * src_pitch_x + y * src_pitch_y + z * src_pitch_z + w * src_pitch_w + i * src_pitch_f + o * src_pitch_b;
uint dst_x = flat_idx % dst_size_x;
flat_idx /= dst_size_x;
uint dst_y = flat_idx % dst_size_y;
flat_idx /= dst_size_y;
uint dst_z = flat_idx % dst_size_z;
flat_idx /= dst_size_z;
uint dst_f = flat_idx % dst_size_f;
flat_idx /= dst_size_f;
uint dst_b = flat_idx;
return (uint8)(0, dst_b, dst_f, 0, dst_z, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_grouped)(uint g, uint o, uint i, uint z, uint y, uint x, uint src_size_ofm, uint dst_size_ofm)
{
const uint flat_ofm = g * src_size_ofm + o;
const uint dst_ofm = flat_ofm % dst_size_ofm;
const uint dst_g = flat_ofm / dst_size_ofm;
const uint dst_ifm = i;
const uint dst_z = z;
const uint dst_y = y;
const uint dst_x = x;
return (uint8)(dst_g, dst_ofm, dst_ifm, 0, dst_z, dst_y, dst_x, 0);
}
inline uint8 FUNC(reshape_weights)(
uint o, uint i, uint w, uint z, uint y, uint x,
uint src_size_f, uint src_size_w, uint src_size_z, uint src_size_y, uint src_size_x,
uint dst_size_f, uint dst_size_w, uint dst_size_z, uint dst_size_y, uint dst_size_x,
uint src_dims, uint dst_dims)
{
if (src_dims == 4 && dst_dims == 2)
{
return FUNC_CALL(reshape_4_to_2)(o,i,y,x,src_size_y,src_size_x);
}
else if (src_dims == 2 && dst_dims == 4)
{
return FUNC_CALL(reshape_2_to_4)(o,i,y,x,dst_size_y,dst_size_x);
}
else if (src_dims == 4 && dst_dims == 6)
{
return FUNC_CALL(reshape_4_to_6)(o, i, y, x, src_size_f, src_size_y, src_size_x, dst_size_f, dst_size_w, dst_size_z, dst_size_y, dst_size_x);
}
else if (src_dims == 6 && dst_dims == 4)
{
return FUNC_CALL(reshape_6_to_4)(o, i, w, z, y, x, src_size_f, src_size_w, src_size_z, src_size_y, src_size_x, dst_size_f, dst_size_y, dst_size_x);
}
else if (src_dims == 4 && dst_dims == 5)
{
return FUNC_CALL(reshape_4_to_5)(o, i, y, x, src_size_f, src_size_y, src_size_x, dst_size_f, dst_size_z, dst_size_y, dst_size_x);
}
else if (src_dims == 5 && dst_dims == 4)
{
return FUNC_CALL(reshape_5_to_4)(o, i, z, y, x, src_size_f, src_size_z, src_size_y, src_size_x, dst_size_f, dst_size_y, dst_size_x);
}
else if (src_dims == 6 && dst_dims == 5)
{
return FUNC_CALL(reshape_6_to_5)(o, i, w, z, y, x, src_size_f, src_size_w, src_size_z, src_size_y, src_size_x, dst_size_f, dst_size_z, dst_size_y, dst_size_x);
}
return (uint8)(0, o, i, w, z, y, x, 0);
}
inline uint8 FUNC(reshape_dims_with_groups)(
uint g, uint o, uint i, uint w, uint z, uint y, uint x,
uint src_size_ofm, uint src_size_ifm, uint src_size_w, uint src_size_z, uint src_size_y, uint src_size_x,
uint dst_size_ofm, uint dst_size_ifm, uint dst_size_w, uint dst_size_z, uint dst_size_y, uint dst_size_x,
uint src_dims, uint dst_dims, uint src_size_groups, uint dst_size_groups)
{
if (src_dims == 5 && dst_dims == 4) // goiyx -> oiyx
{
return FUNC_CALL(reshape_grouped)(g, o, i, 0, y, x, src_size_ofm, dst_size_ofm);
}
else if (src_dims == 6 && dst_dims == 5) // goizyx -> oizyx or goizyx -> goiyx
{
return FUNC_CALL(reshape_grouped)(g, o, i, z, y, x, src_size_ofm, dst_size_ofm);
}
else if (src_dims == 6 && dst_dims == 4) // goizyx -> oiyx
{
return FUNC_CALL(reshape_grouped)(g, o, i, 0, y, x, src_size_ofm, dst_size_ofm);
}
return (uint8)(g, o, i, w, z, y, x, 0);
}
// ======================================================================
// Reshape indices between WeightTensors.
// Weight tensor has _IFM in place of _FEATURE_NUM and does not support 6D,
// so macro for DataTensors cannot be used directly.
#define RESHAPE_WEIGHT_DIMS(src_prefix, dst_prefix, o, i, w, z, y, x) \
FUNC_CALL(reshape_weights)( \
o, i, w, z, y, x, \
CAT(src_prefix, _IFM_NUM), \
1, \
CAT(src_prefix, _SIZE_Z), \
CAT(src_prefix, _SIZE_Y), \
CAT(src_prefix, _SIZE_X), \
CAT(dst_prefix, _IFM_NUM), \
1, \
CAT(dst_prefix, _SIZE_Z), \
CAT(dst_prefix, _SIZE_Y), \
CAT(dst_prefix, _SIZE_X), \
CAT(src_prefix, _DIMS), \
CAT(dst_prefix, _DIMS))
#define RESHAPE_WEIGHT_DIMS_WITH_GROUPS(src_prefix, dst_prefix, g, o, i, w, z, y, x)\
FUNC_CALL(reshape_dims_with_groups)( \
g, o, i, w, z, y, x, \
CAT(src_prefix, _OFM_NUM), \
CAT(src_prefix, _IFM_NUM), \
1, \
CAT(src_prefix, _SIZE_Z), \
CAT(src_prefix, _SIZE_Y), \
CAT(src_prefix, _SIZE_X), \
CAT(dst_prefix, _OFM_NUM), \
CAT(dst_prefix, _IFM_NUM), \
1, \
CAT(dst_prefix, _SIZE_Z), \
CAT(dst_prefix, _SIZE_Y), \
CAT(dst_prefix, _SIZE_X), \
CAT(src_prefix, _DIMS), \
CAT(dst_prefix, _DIMS), \
CAT(src_prefix, _GROUPS_NUM), \
CAT(dst_prefix, _GROUPS_NUM))
///////////////////////// Input Index /////////////////////////
inline uint FUNC(get_input_index)(uint g, uint o, uint i, uint z, uint y, uint x)
{
src/plugins/intel_gpu/src/kernel_selector/jitter.cpp
+64 -14
View File
@@ -318,7 +318,6 @@ public:
JitDefinitions DataTensorJitConstant::GetDefinitions() const {
JitDefinitions baseDefinitions = TensorBaseTJitConstant::GetDefinitions(_tensor);
JitDefinitions definitions{};
DimensionAccessHelper dims(_tensor, _dyn_array_index);
@@ -329,18 +328,24 @@ JitDefinitions DataTensorJitConstant::GetDefinitions() const {
{_name + "_SIZE_Y", dims.y},
{_name + "_SIZE_Z", dims.z},
{_name + "_SIZE_W", dims.w},
{_name + "_SIZE_U", dims.u},
{_name + "_SIZE_V", dims.v},
{_name + "_FEATURE_NUM", dims.f},
{_name + "_BATCH_NUM", dims.b},
{_name + "_PAD_BEFORE_SIZE_X", toCodeString(_tensor.X().pad.before)},
{_name + "_PAD_BEFORE_SIZE_Y", toCodeString(_tensor.Y().pad.before)},
{_name + "_PAD_BEFORE_SIZE_Z", toCodeString(_tensor.Z().pad.before)},
{_name + "_PAD_BEFORE_SIZE_W", toCodeString(_tensor.W().pad.before)},
{_name + "_PAD_BEFORE_SIZE_U", toCodeString(_tensor.U().pad.before)},
{_name + "_PAD_BEFORE_SIZE_V", toCodeString(_tensor.V().pad.before)},
{_name + "_PAD_BEFORE_FEATURE_NUM", toCodeString(_tensor.Feature().pad.before)},
{_name + "_PAD_BEFORE_BATCH_NUM", toCodeString(_tensor.Batch().pad.before)},
{_name + "_PAD_AFTER_SIZE_X", toCodeString(_tensor.X().pad.after)},
{_name + "_PAD_AFTER_SIZE_Y", toCodeString(_tensor.Y().pad.after)},
{_name + "_PAD_AFTER_SIZE_Z", toCodeString(_tensor.Z().pad.after)},
{_name + "_PAD_AFTER_SIZE_W", toCodeString(_tensor.W().pad.after)},
{_name + "_PAD_AFTER_SIZE_U", toCodeString(_tensor.U().pad.after)},
{_name + "_PAD_AFTER_SIZE_V", toCodeString(_tensor.V().pad.after)},
{_name + "_PAD_AFTER_FEATURE_NUM", toCodeString(_tensor.Feature().pad.after)},
{_name + "_PAD_AFTER_BATCH_NUM", toCodeString(_tensor.Batch().pad.after)},
};
@@ -354,8 +359,12 @@ JitDefinitions DataTensorJitConstant::GetDefinitions() const {
definitions.push_back({_name + "_Y_PITCH", dims_padded.x});
definitions.push_back({_name + "_Z_PITCH", toVectorMulString({dims_padded.x, dims_padded.y})});
definitions.push_back({_name + "_W_PITCH", toVectorMulString({dims_padded.x, dims_padded.y, dims_padded.z})});
definitions.push_back({_name + "_FEATURE_PITCH", toVectorMulString({dims_padded.x, dims_padded.y, dims_padded.z, dims_padded.w})});
definitions.push_back({_name + "_BATCH_PITCH", toVectorMulString({dims_padded.x, dims_padded.y, dims_padded.z, dims_padded.w, dims_padded.f})});
definitions.push_back({_name + "_U_PITCH", toVectorMulString({dims_padded.x, dims_padded.y, dims_padded.z, dims_padded.w})});
definitions.push_back({_name + "_V_PITCH", toVectorMulString({dims_padded.x, dims_padded.y, dims_padded.z, dims_padded.w, dims_padded.u})});
definitions.push_back({_name + "_FEATURE_PITCH",
toVectorMulString({dims_padded.x, dims_padded.y, dims_padded.z, dims_padded.w, dims_padded.u, dims_padded.v})});
definitions.push_back({_name + "_BATCH_PITCH",
toVectorMulString({dims_padded.x, dims_padded.y, dims_padded.z, dims_padded.w, dims_padded.u, dims_padded.v, dims_padded.f})});
} else {
OPENVINO_ASSERT(false, "[GPU] Jitter couldn't generate dynamic pitches for given layout");
}
@@ -364,6 +373,8 @@ JitDefinitions DataTensorJitConstant::GetDefinitions() const {
definitions.push_back({_name + "_Y_PITCH", toCodeString(_tensor.Y().pitch)});
definitions.push_back({_name + "_Z_PITCH", toCodeString(_tensor.Z().pitch)});
definitions.push_back({_name + "_W_PITCH", toCodeString(_tensor.W().pitch)});
definitions.push_back({_name + "_U_PITCH", toCodeString(_tensor.U().pitch)});
definitions.push_back({_name + "_V_PITCH", toCodeString(_tensor.V().pitch)});
definitions.push_back({_name + "_FEATURE_PITCH", toCodeString(_tensor.Feature().pitch)});
definitions.push_back({_name + "_BATCH_PITCH", toCodeString(_tensor.Batch().pitch)});
}
@@ -525,18 +536,53 @@ JitDefinitions DataTensorJitConstant::GetDefinitions() const {
Tensor::DataChannelName::X,
};
bool is_common_6d_layout = is_common_nd_layout(base_6d_channels, layout);
if (is_common_6d_layout) {
index_func_name = _name + "_GET_INDEX(b, f, w, z, y, x)";
safe_index_func_name = _name + "_GET_INDEX_SAFE(b, f, w, z, y, x)";
raw_index_func_name = _name + "_GET_INDEX_RAW(b, f, w, z, y, x)";
OPENVINO_ASSERT(is_common_6d_layout, "[GPU] Unhandled 6d format in jitter: ", toString(layout));
index_func_name = _name + "_GET_INDEX(b, f, w, z, y, x)";
safe_index_func_name = _name + "_GET_INDEX_SAFE(b, f, w, z, y, x)";
raw_index_func_name = _name + "_GET_INDEX_RAW(b, f, w, z, y, x)";
index_func_val = "GET_DATA_INDEX_6D(" + _name + ", b, f, w, z, y, x)";
safe_index_func_val = "GET_DATA_INDEX_6D_SAFE(" + _name + ", b, f, w, z, y, x)";
raw_index_func_val = "GET_DATA_INDEX_6D_RAW(" + _name + ", b, f, w, z, y, x)";
} else {
// TODO: implement support of non-default layouts with 6 channels
assert(0);
}
index_func_val = "GET_DATA_INDEX_6D(" + _name + ", b, f, w, z, y, x)";
safe_index_func_val = "GET_DATA_INDEX_6D_SAFE(" + _name + ", b, f, w, z, y, x)";
raw_index_func_val = "GET_DATA_INDEX_6D_RAW(" + _name + ", b, f, w, z, y, x)";
} else if (DataTensor::ChannelsCount(layout) == 7) {
std::vector<Tensor::DataChannelName> base_7d_channels = {
Tensor::DataChannelName::BATCH,
Tensor::DataChannelName::FEATURE,
Tensor::DataChannelName::U,
Tensor::DataChannelName::W,
Tensor::DataChannelName::Z,
Tensor::DataChannelName::Y,
Tensor::DataChannelName::X,
};
bool is_common_7d_layout = is_common_nd_layout(base_7d_channels, layout);
OPENVINO_ASSERT(is_common_7d_layout, "[GPU] Unhandled 7d format in jitter: ", toString(layout));
index_func_name = _name + "_GET_INDEX(b, f, u, w, z, y, x)";
safe_index_func_name = _name + "_GET_INDEX_SAFE(b, f, u, w, z, y, x)";
raw_index_func_name = _name + "_GET_INDEX_RAW(b, f, u, w, z, y, x)";
index_func_val = "GET_DATA_INDEX_7D(" + _name + ", b, f, u, w, z, y, x)";
safe_index_func_val = "GET_DATA_INDEX_7D_SAFE(" + _name + ", b, f, u, w, z, y, x)";
raw_index_func_val = "GET_DATA_INDEX_7D_RAW(" + _name + ", b, f, u, w, z, y, x)";
} else if (DataTensor::ChannelsCount(layout) == 8) {
std::vector<Tensor::DataChannelName> base_8d_channels = {
Tensor::DataChannelName::BATCH,
Tensor::DataChannelName::FEATURE,
Tensor::DataChannelName::V,
Tensor::DataChannelName::U,
Tensor::DataChannelName::W,
Tensor::DataChannelName::Z,
Tensor::DataChannelName::Y,
Tensor::DataChannelName::X,
};
bool is_common_8d_layout = is_common_nd_layout(base_8d_channels, layout);
OPENVINO_ASSERT(is_common_8d_layout, "[GPU] Unhandled 8d format in jitter: ", toString(layout));
index_func_name = _name + "_GET_INDEX(b, f, v, u, w, z, y, x)";
safe_index_func_name = _name + "_GET_INDEX_SAFE(b, f, v, u, w, z, y, x)";
raw_index_func_name = _name + "_GET_INDEX_RAW(b, f, v, u, w, z, y, x)";
index_func_val = "GET_DATA_INDEX_8D(" + _name + ", b, f, v, u, w, z, y, x)";
safe_index_func_val = "GET_DATA_INDEX_8D_SAFE(" + _name + ", b, f, v, u, w, z, y, x)";
raw_index_func_val = "GET_DATA_INDEX_8D_RAW(" + _name + ", b, f, v, u, w, z, y, x)";
} else {
throw std::runtime_error("Unsupported channels count(" + std::to_string(DataTensor::ChannelsCount(layout)) +
") in layout: " + toString(layout));
@@ -1902,6 +1948,10 @@ std::string FusedOpsCodeGenerator::GetIdx(size_t input_id, idx_desc idx, bool sh
idx_order = idx.b + "," + idx.f + "," + idx.z + "," + idx.y + "," + idx.x;
} else if (DataTensor::ChannelsCount(desc.tensors[input_id].GetLayout()) == 6) {
idx_order = idx.b + "," + idx.f + "," + idx.w + "," + idx.z + "," + idx.y + "," + idx.x;
} else if (DataTensor::ChannelsCount(desc.tensors[input_id].GetLayout()) == 7) {
idx_order = idx.b + "," + idx.f + "," + idx.u + "," + idx.w + "," + idx.z + "," + idx.y + "," + idx.x;
} else if (DataTensor::ChannelsCount(desc.tensors[input_id].GetLayout()) == 8) {
idx_order = idx.b + "," + idx.f + "," + idx.v + "," + idx.u + "," + idx.w + "," + idx.z + "," + idx.y + "," + idx.x;
}
if (should_be_safe) {
src/plugins/intel_gpu/src/kernel_selector/jitter.h
+12 -2
View File
@@ -331,13 +331,15 @@ public:
struct idx_desc {
std::string b;
std::string f;
std::string v;
std::string u;
std::string w;
std::string z;
std::string y;
std::string x;
size_t dims;
explicit idx_desc(std::vector<std::string> idx, DataTensor t)
: b("0"), f("0"), w("0"), z("0"), y("0"), x("0"), dims(0) {
: b("0"), f("0"), v("0"), u("0"), w("0"), z("0"), y("0"), x("0"), dims(0) {
dims = idx.size();
switch (dims) {
case 1: f = idx[0]; break;
@@ -346,7 +348,9 @@ public:
case 4: b = idx[0]; f = idx[1]; y = idx[2]; x = idx[3]; break;
case 5: b = idx[0]; f = idx[1]; z = idx[2]; y = idx[3]; x = idx[4]; break;
case 6: b = idx[0]; f = idx[1]; w = idx[2]; z = idx[3]; y = idx[4]; x = idx[5]; break;
default: throw std::runtime_error("More than 6 dimenstions is not supported in fused op generator");
case 7: b = idx[0]; f = idx[1]; u = idx[2]; w = idx[3]; z = idx[4]; y = idx[5]; x = idx[6]; break;
case 8: b = idx[0]; f = idx[1]; v = idx[2]; u = idx[3]; w = idx[4]; z = idx[5]; y = idx[6]; x = idx[7]; break;
default: throw std::runtime_error("More than 8 dimenstions is not supported in fused op generator");
}
if (t.Batch().v == 1) {
@@ -355,6 +359,12 @@ public:
if (t.Feature().v == 1) {
f = "0";
}
if (t.V().v == 1) {
v = "0";
}
if (t.U().v == 1) {
u = "0";
}
if (t.W().v == 1) {
w = "0";
}
src/plugins/intel_gpu/src/kernel_selector/kernel_selector_common.cpp
+2
View File
@@ -104,6 +104,8 @@ std::string toString(DataLayout l) {
case kernel_selector::DataLayout::bzyxf: return "BZYXF";
case kernel_selector::DataLayout::fs_b_yx_fsv32: return "FS_B_YX_FSV32";
case kernel_selector::DataLayout::bfwzyx: return "BFWZYX";
case kernel_selector::DataLayout::bfuwzyx: return "BFUWZYX";
case kernel_selector::DataLayout::bfvuwzyx: return "BFVUWZYX";
case kernel_selector::DataLayout::b_fs_zyx_fsv16: return "B_FS_ZYX_FSV16";
case kernel_selector::DataLayout::bs_fs_yx_bsv16_fsv16: return "BS_FS_YX_BSV16_FSV16";
case kernel_selector::DataLayout::bs_fs_yx_bsv16_fsv32: return "BS_FS_YX_BSV16_FSV32";
src/plugins/intel_gpu/src/kernel_selector/kernel_selector_utils.cpp
+31 -19
View File
@@ -184,18 +184,24 @@ std::vector<size_t> GetTensorFriendlyWorkGroups(const DataTensor& t) {
auto y = DataTensor::Channelndex(t.GetLayout(), Tensor::DataChannelName::Y);
auto z = DataTensor::Channelndex(t.GetLayout(), Tensor::DataChannelName::Z);
auto w = DataTensor::Channelndex(t.GetLayout(), Tensor::DataChannelName::W);
auto u = DataTensor::Channelndex(t.GetLayout(), Tensor::DataChannelName::U);
auto v = DataTensor::Channelndex(t.GetLayout(), Tensor::DataChannelName::V);
auto primary_spatial_axis = x;
if (y < primary_spatial_axis && y != -1) primary_spatial_axis = y;
if (z < primary_spatial_axis && z != -1) primary_spatial_axis = z;
if (w < primary_spatial_axis && w != -1) primary_spatial_axis = w;
if (u < primary_spatial_axis && u != -1) primary_spatial_axis = u;
if (v < primary_spatial_axis && v != -1) primary_spatial_axis = v;
for (size_t i = 0; i < t.GetDims().size(); i++) {
const auto& o = t.GetDims()[i];
auto cur_axis_is_spatial = x == static_cast<int>(i) ||
y == static_cast<int>(i) ||
z == static_cast<int>(i) ||
w == static_cast<int>(i);
w == static_cast<int>(i) ||
u == static_cast<int>(i) ||
v == static_cast<int>(i);
if (cur_axis_is_spatial && primary_spatial_axis != static_cast<int>(i)) {
sizes.back() *= o.v;
} else {
@@ -213,18 +219,18 @@ std::vector<size_t> GetTensorFriendlyWorkGroups(const DataTensor& t) {
std::vector<size_t> GetOptimalLocalWorkGroupSizes(std::vector<size_t> gws, const EngineInfo& info,
DataLayout input_layout, DataLayout output_layout,
std::vector<std::vector<Tensor::DataChannelName>> dims_by_gws) {
enum axis { x, y, z, w, f, b, unused_axis };
enum axis { x, y, z, w, u, v, f, b, unused_axis };
// GWS/LWS priority order should be considered for better local WGS setting
// and as a result more optimized data reading/writing inside kernels
std::vector<size_t> priority_order = { 0, 1, 2 };
std::vector<size_t> layout_order = { x, y, z, w, f, b };
std::vector<size_t> layout_order = { x, y, z, w, u, v, f, b };
const size_t gws_dims_num = priority_order.size();
const size_t axis_num = layout_order.size();
size_t first_axis_idx = 0;
std::vector<size_t> axis_by_gws = { unused_axis, unused_axis, unused_axis, unused_axis, unused_axis, unused_axis };
std::vector<size_t> axis_by_gws = { unused_axis, unused_axis, unused_axis, unused_axis, unused_axis, unused_axis, unused_axis, unused_axis };
for (size_t gws_idx = 0; gws_idx < gws_dims_num; gws_idx++) {
for (size_t axis_idx = 0; axis_idx < dims_by_gws[gws_idx].size(); axis_idx++) {
axis_by_gws[static_cast<size_t>(dims_by_gws[gws_idx][axis_idx])] = gws_idx;
@@ -280,51 +286,57 @@ std::vector<size_t> GetOptimalLocalWorkGroupSizes(std::vector<size_t> gws, const
if (simple_planar_layout) {
switch (output_layout) {
case DataLayout::bf:
layout_order = { f, b, x, y, z, w };
layout_order = { f, b, x, y, z, w, u, v };
break;
case DataLayout::fb:
layout_order = { b, f, x, y, z, w };
layout_order = { b, f, x, y, z, w, u, v };
break;
case DataLayout::bfyx:
layout_order = { x, y, f, b, z, w };
layout_order = { x, y, f, b, z, w, u, v };
break;
case DataLayout::yxfb:
layout_order = { b, f, x, y, z, w };
layout_order = { b, f, x, y, z, w, u, v };
break;
case DataLayout::byxf:
layout_order = { f, x, y, b, z, w };
layout_order = { f, x, y, b, z, w, u, v };
break;
case DataLayout::fyxb:
layout_order = { b, x, y, f, z, w };
layout_order = { b, x, y, f, z, w, u, v };
break;
case DataLayout::bfxy:
layout_order = { y, x, f, b, z, w };
layout_order = { y, x, f, b, z, w, u, v };
break;
case DataLayout::bfzyx:
layout_order = { x, y, z, f, b, w };
layout_order = { x, y, z, f, b, w, u, v };
break;
case DataLayout::bzyxf:
layout_order = { f, x, y, z, b, w };
layout_order = { f, x, y, z, b, w, u, v };
break;
case DataLayout::bfwzyx:
layout_order = { x, y, z, w, f, b };
layout_order = { x, y, z, w, f, b, u, v };
break;
case DataLayout::bfuwzyx:
layout_order = { x, y, z, w, u, f, b, v };
break;
case DataLayout::bfvuwzyx:
layout_order = { x, y, z, w, u, v , f, b };
break;
default:
layout_order = { x, y, z, w, f, b };
layout_order = { x, y, z, w, u, v, f, b };
break;
}
} else if (blocked_fsv_layout) {
if (output_layout == DataLayout::b_fs_yx_fsv2 || output_layout == DataLayout::b_fs_yx_fsv4 ||
output_layout == DataLayout::b_fs_yx_fsv16 || output_layout == DataLayout::b_fs_yx_fsv32) {
layout_order = { f, x, y, b, z, w };
layout_order = { f, x, y, b, z, w, u, v };
} else if (output_layout == DataLayout::b_fs_zyx_fsv2 || output_layout == DataLayout::b_fs_zyx_fsv4 ||
output_layout == DataLayout::b_fs_zyx_fsv16 || output_layout == DataLayout::b_fs_zyx_fsv32) {
layout_order = { f, x, y, z, b, w };
layout_order = { f, x, y, z, b, w, u, v };
} else { // output_layout == DataLayout::fs_b_yx_fsv32
layout_order = { f, x, y, b, z, w };
layout_order = { f, x, y, b, z, w, u, v };
}
} else if (blocked_bsv_fsv_layout) {
layout_order = { f, b, x, y, z, w };
layout_order = { f, b, x, y, z, w, u, v };
}
calculate_optimized_priority_order();
src/plugins/intel_gpu/src/kernel_selector/kernel_selector_utils.h
+4 -2
View File
@@ -14,7 +14,7 @@ struct WeightsReorderParams;
struct DimensionAccessHelper {
static size_t dynamic_dimension_offset(Tensor::DataChannelName c, size_t tensor_id = 0) {
const auto shape_info_layout = DataLayout::bfwzyx;
const auto shape_info_layout = DataLayout::bfvuwzyx;
auto idx = DataTensor::max_rank() - DataTensor::Channelndex(shape_info_layout, c) - 1;
return DataTensor::max_rank() * tensor_id + idx;
}
@@ -24,10 +24,12 @@ struct DimensionAccessHelper {
, y(toCodeString(t.Y(), dynamic_dimension_offset(Tensor::DataChannelName::Y, dyn_tensor_id), padded))
, z(toCodeString(t.Z(), dynamic_dimension_offset(Tensor::DataChannelName::Z, dyn_tensor_id), padded))
, w(toCodeString(t.W(), dynamic_dimension_offset(Tensor::DataChannelName::W, dyn_tensor_id), padded))
, u(toCodeString(t.U(), dynamic_dimension_offset(Tensor::DataChannelName::U, dyn_tensor_id), padded))
, v(toCodeString(t.V(), dynamic_dimension_offset(Tensor::DataChannelName::V, dyn_tensor_id), padded))
, f(toCodeString(t.Feature(), dynamic_dimension_offset(Tensor::DataChannelName::FEATURE, dyn_tensor_id), padded))
, b(toCodeString(t.Batch(), dynamic_dimension_offset(Tensor::DataChannelName::BATCH, dyn_tensor_id), padded)) { }
std::string x, y, z, w, f, b;
std::string x, y, z, w, u, v, f, b;
};
std::vector<size_t> GetImageSizes(const kernel_selector::WeightsTensor& dimensions, const WeightsLayout layout);
src/plugins/intel_gpu/src/kernel_selector/kernels/eltwise/eltwise_kernel_base.cpp
+29 -1
View File
@@ -517,6 +517,28 @@ JitConstants EltwiseKernelBase::MakeIndexJitConstants(const eltwise_params& para
} else {
jit.AddConstant(MakeJitConstant(idx_order, "d6,d5,d4,d3,d2,d1"));
}
} else if (out_c == 7) {
if (in_c < 5) {
jit.AddConstant(MakeJitConstant(idx_order, "d7,d6,d2,d1"));
} else if (in_c == 5) {
jit.AddConstant(MakeJitConstant(idx_order, "d7,d6,d3,d2,d1"));
} else if (in_c == 6) {
jit.AddConstant(MakeJitConstant(idx_order, "d7,d6,d4,d3,d2,d1"));
} else {
jit.AddConstant(MakeJitConstant(idx_order, "d7,d6,d5,d4,d3,d2,d1"));
}
} else if (out_c == 8) {
if (in_c < 5) {
jit.AddConstant(MakeJitConstant(idx_order, "d8,d7,d2,d1"));
} else if (in_c == 5) {
jit.AddConstant(MakeJitConstant(idx_order, "d8,d7,d3,d2,d1"));
} else if (in_c == 6) {
jit.AddConstant(MakeJitConstant(idx_order, "d8,d7,d4,d3,d2,d1"));
} else if (in_c == 7) {
jit.AddConstant(MakeJitConstant(idx_order, "d8,d7,d5,d4,d3,d2,d1"));
} else {
jit.AddConstant(MakeJitConstant(idx_order, "d8,d7,d6,d5,d4,d3,d2,d1"));
}
} else {
assert(0);
}
@@ -601,7 +623,13 @@ EltwiseKernelBase::DispatchData EltwiseKernelBase::SetDefault(const eltwise_para
}
dispatchData.gws[0] = gws[0];
if (n_dims == 6) {
if (n_dims == 8) {
dispatchData.gws[1] = gws[1] * gws[2] * gws[3] * gws[4] * gws[5]; // y*z*w*u*v
dispatchData.gws[2] = gws[6] * gws[7];
} else if (n_dims == 7) {
dispatchData.gws[1] = gws[1] * gws[2] * gws[3] * gws[4]; // y*z*w*u
dispatchData.gws[2] = gws[5] * gws[6];
} else if (n_dims == 6) {
dispatchData.gws[1] = gws[1] * gws[2] * gws[3]; // y*z*w
dispatchData.gws[2] = gws[4] * gws[5];
} else if (n_dims == 5) {
src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bs_f_bsv16_b1.cpp
+1 -1
View File
@@ -64,7 +64,7 @@ FullyConnected_bs_f_bsv16_b1::DispatchData FullyConnected_bs_f_bsv16_b1::SetDefa
const uint32_t units_per_sg_read = sub_group_size * units_per_chunk;
// Properties of primitive responses.
constexpr uint32_t responses_per_sg_exec =
16; // Must match batch slice size of weights format (bs_x_bsv16).
16; // Must match batch slice size of weights format (bs_f_bsv16).
// Number of response groups. Each group (except last) writes responses_per_sg_exec responses
// for at least one input data set from batch.
const auto response_size = arg.outputs[0].Feature().v;
src/plugins/intel_gpu/src/kernel_selector/kernels/gather/gather_kernel_ref.cpp
+1 -1
View File
@@ -65,7 +65,7 @@ static size_t GetNonEmptyDimsNumber(const DataTensor& data_tensor) {
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++) {
for (size_t i = 0; i < DataTensor::max_rank(); i++) {
int shape_raw_idx =
data_tensor.Channelndex(data_tensor.GetLayout(), static_cast<Tensor::DataChannelName>(i));
if (shape_raw_idx >= 0)
src/plugins/intel_gpu/src/kernel_selector/kernels/permute/permute_kernel_ref.cpp
+14 -12
View File
@@ -134,14 +134,16 @@ CommonDispatchData PermuteKernelRef::SetDefault(const permute_params& params) co
std::vector<std::vector<Tensor::DataChannelName>> dims_by_gws = {{Tensor::DataChannelName::FEATURE},
{Tensor::DataChannelName::X, Tensor::DataChannelName::Y},
{Tensor::DataChannelName::Z, Tensor::DataChannelName::W,
Tensor::DataChannelName::U, Tensor::DataChannelName::V,
Tensor::DataChannelName::BATCH}};
dispatchData.gws = {in.Feature().v, in.X().v * in.Y().v, in.Z().v * in.W().v * in.Batch().v};
dispatchData.gws = {in.Feature().v, in.X().v * in.Y().v, in.Z().v * in.W().v * in.U().v * in.V().v * in.Batch().v};
dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo, in_layout, out_layout, dims_by_gws);
} else {
std::vector<std::vector<Tensor::DataChannelName>> dims_by_gws = {{Tensor::DataChannelName::X},
{Tensor::DataChannelName::Y, Tensor::DataChannelName::Z, Tensor::DataChannelName::W},
{Tensor::DataChannelName::Y, Tensor::DataChannelName::Z, Tensor::DataChannelName::W,
Tensor::DataChannelName::U, Tensor::DataChannelName::V},
{Tensor::DataChannelName::FEATURE, Tensor::DataChannelName::BATCH}};
dispatchData.gws = {in.X().v, in.Y().v * in.Z().v * in.W().v, in.Feature().v * in.Batch().v};
dispatchData.gws = {in.X().v, in.Y().v * in.Z().v * in.W().v * in.U().v * in.V().v , in.Feature().v * in.Batch().v};
dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo, in_layout, out_layout, dims_by_gws);
}
@@ -151,6 +153,13 @@ CommonDispatchData PermuteKernelRef::SetDefault(const permute_params& params) co
bool PermuteKernelRef::Validate(const Params& p, const optional_params& o) const {
if (!Parent::Validate(p, o)) return false;
const permute_params& params = static_cast<const permute_params&>(p);
auto in_rank = params.inputs[0].GetDims().size();
auto out_rank = params.outputs[0].GetDims().size();
if (in_rank != out_rank && (in_rank > 6 || out_rank > 6))
return false;
return true;
}
@@ -159,15 +168,6 @@ JitConstants PermuteKernelRef::GetJitConstants(const permute_params& params, con
std::vector<std::string> in_idx;
std::vector<std::string> permute_out_idx;
std::map<std::string, std::string> size_str_map = {
{"b", "INPUT0_BATCH_NUM"},
{"f", "INPUT0_FEATURE_NUM"},
{"w", "INPUT0_SIZE_W"},
{"z", "INPUT0_SIZE_Z"},
{"y", "INPUT0_SIZE_Y"},
{"x", "INPUT0_SIZE_X"}
};
std::pair<size_t, size_t> dim_change;
bool reorder_to_different_dim = false;
std::vector<std::string> reordered_out_idx;
@@ -178,6 +178,8 @@ JitConstants PermuteKernelRef::GetJitConstants(const permute_params& params, con
}
switch (DataTensor::ChannelsCount(params.inputs[0].GetLayout())) {
case 8: in_idx = {"b", "f", "x", "y", "z", "w", "u", "v" }; break;
case 7: in_idx = {"b", "f", "x", "y", "z", "w", "u" }; break;
case 6: in_idx = {"b", "f", "x", "y", "z", "w" }; break;
case 5: in_idx = {"b", "f", "x", "y", "z" }; break;
default: in_idx = {"b", "f", "x", "y" }; break;
src/plugins/intel_gpu/src/kernel_selector/kernels/reduce/reduce_kernel_base.cpp
+51 -32
View File
@@ -32,7 +32,14 @@ JitConstants ReduceKernelBase::GetJitConstants(const reduce_params& params) cons
if (output.is_dynamic()) {
size_t output_tensor_offset = 1 + GetFusedPrimitiveInputsCount(params);
DimensionAccessHelper dims(output, output_tensor_offset);
jit.AddConstant(MakeJitConstant("COMPUTATIONAL_OPERATIONS_NUMBER", toVectorMulString({dims.x, dims.y, dims.z, dims.w, dims.f, dims.b})));
jit.AddConstant(MakeJitConstant("COMPUTATIONAL_OPERATIONS_NUMBER", toVectorMulString({dims.x,
dims.y,
dims.z,
dims.w,
dims.u,
dims.v,
dims.f,
dims.b})));
} else {
jit.AddConstant(MakeJitConstant("COMPUTATIONAL_OPERATIONS_NUMBER", params.outputs[0].LogicalSize()));
}
@@ -59,44 +66,50 @@ JitConstants ReduceKernelBase::GetJitConstants(const reduce_params& params) cons
};
auto getDimSizeNameByNum = [&](size_t dim) -> std::string {
if (params.inputs[0].Dimentions() == 6) {
if (params.inputs[0].Dimentions() == 8) {
switch (dim) {
case 0:
return "BATCH_NUM";
case 1:
return "FEATURE_NUM";
case 2:
return "SIZE_W";
case 3:
return "SIZE_Z";
case 4:
return "SIZE_Y";
case 5:
return "SIZE_X";
case 0: return "BATCH_NUM";
case 1: return "FEATURE_NUM";
case 2: return "SIZE_V";
case 3: return "SIZE_U";
case 4: return "SIZE_W";
case 5: return "SIZE_Z";
case 6: return "SIZE_Y";
case 7: return "SIZE_X";
}
} else if (params.inputs[0].Dimentions() == 7) {
switch (dim) {
case 0: return "BATCH_NUM";
case 1: return "FEATURE_NUM";
case 2: return "SIZE_U";
case 3: return "SIZE_W";
case 4: return "SIZE_Z";
case 5: return "SIZE_Y";
case 6: return "SIZE_X";
}
} else if (params.inputs[0].Dimentions() == 6) {
switch (dim) {
case 0: return "BATCH_NUM";
case 1: return "FEATURE_NUM";
case 2: return "SIZE_W";
case 3: return "SIZE_Z";
case 4: return "SIZE_Y";
case 5: return "SIZE_X";
}
} else if (params.inputs[0].Dimentions() == 5) {
switch (dim) {
case 0:
return "BATCH_NUM";
case 1:
return "FEATURE_NUM";
case 2:
return "SIZE_Z";
case 3:
return "SIZE_Y";
case 4:
return "SIZE_X";
case 0: return "BATCH_NUM";
case 1: return "FEATURE_NUM";
case 2: return "SIZE_Z";
case 3: return "SIZE_Y";
case 4: return "SIZE_X";
}
} else if (params.inputs[0].Dimentions() == 4) {
switch (dim) {
case 0:
return "BATCH_NUM";
case 1:
return "FEATURE_NUM";
case 2:
return "SIZE_Y";
case 3:
return "SIZE_X";
case 0: return "BATCH_NUM";
case 1: return "FEATURE_NUM";
case 2: return "SIZE_Y";
case 3: return "SIZE_X";
}
}
return "";
@@ -169,6 +182,12 @@ JitConstants ReduceKernelBase::GetJitConstants(const reduce_params& params) cons
case 5:
jit.AddConstant(MakeJitConstant("REDUCE_W", 1));
break;
case 6:
jit.AddConstant(MakeJitConstant("REDUCE_U", 1));
break;
case 7:
jit.AddConstant(MakeJitConstant("REDUCE_V", 1));
break;
}
}
src/plugins/intel_gpu/src/kernel_selector/kernels/reduce/reduce_kernel_ref.cpp
+5 -2
View File
@@ -36,11 +36,12 @@ CommonDispatchData ReduceKernelRef::SetDefault(const reduce_params& params) cons
auto in_layout = params.inputs[0].GetLayout();
auto out_layout = params.outputs[0].GetLayout();
std::vector<std::vector<Tensor::DataChannelName>> dims_by_gws = {{ Tensor::DataChannelName::X, Tensor::DataChannelName::Y },
{ Tensor::DataChannelName::Z, Tensor::DataChannelName::W },
{ Tensor::DataChannelName::Z, Tensor::DataChannelName::W,
Tensor::DataChannelName::U, Tensor::DataChannelName::V },
{ Tensor::DataChannelName::FEATURE, Tensor::DataChannelName::BATCH }};
dispatchData.gws = { params.outputs[0].X().v * params.outputs[0].Y().v,
params.outputs[0].Z().v * params.outputs[0].W().v,
params.outputs[0].Z().v * params.outputs[0].W().v * params.outputs[0].U().v * params.outputs[0].V().v,
params.outputs[0].Batch().v * params.outputs[0].Feature().v };
dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo, in_layout, out_layout, dims_by_gws);
@@ -59,6 +60,8 @@ JitConstants ReduceKernelRef::GetJitConstants(const reduce_params& params) const
std::vector<std::string> idx_order;
switch (DataTensor::ChannelsCount(params.inputs[0].GetLayout())) {
case 8: idx_order = {"b", "f", "v", "u", "w", "z", "y", "x" }; break;
case 7: idx_order = {"b", "f", "u", "w", "z", "y", "x" }; break;
case 6: idx_order = {"b", "f", "w", "z", "y", "x" }; break;
case 5: idx_order = {"b", "f", "z", "y", "x" }; break;
default: idx_order = {"b", "f", "y", "x" }; break;
src/plugins/intel_gpu/src/kernel_selector/tensor_type.cpp
+46 -44
View File
@@ -14,50 +14,52 @@ DataTensor::DataChannelArray DataTensor::dataChannelArray {{
// explanation:
// 0, 1, 2, 3, 4, 5, 6 means the ordering starts from X, then Y, then Z, then W, then F, then B
// -1 means it's not used
// X, Y, Z, W, F, B
{ DataLayout::f, { -1, -1, -1, -1, 0, -1 } },
{ DataLayout::bf, { -1, -1, -1, -1, 0, 1 } },
{ DataLayout::fb, { -1, -1, -1, -1, 1, 0 } },
{ DataLayout::bfyx, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::yxfb, { 2, 3, -1, -1, 1, 0 } },
{ DataLayout::byxf, { 1, 2, -1, -1, 0, 3 } },
{ DataLayout::fyxb, { 1, 2, -1, -1, 3, 0 } },
{ DataLayout::b_fs_yx_fsv2, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::b_fs_yx_fsv4, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::b_fs_yx_fsv16, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::b_fs_yx_fsv32, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::b_fs_zyx_fsv2, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::b_fs_zyx_fsv4, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::b_fs_zyx_fsv16, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::b_fs_zyx_fsv32, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv16_fsv32, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv16_fsv32, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::bs_fs_zyx_bsv16_fsv16, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv16_fsv16, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_yx_bsv4_fsv4, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_yx_bsv8_fsv4, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv8_fsv4, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv16_fsv4, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv16_fsv4, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv16_fsv2, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv16_fsv2, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv8_fsv2, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv8_fsv2, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv4_fsv2, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv32_fsv32, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv32_fsv32, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv32_fsv16, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv32_fsv16, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::bs_f_bsv8__af8, { -1, -1, -1, -1, 0, 1 } },
{ DataLayout::bs_f_bsv16__af8, { -1, -1, -1, -1, 0, 1 } },
{ DataLayout::winograd_2x3_s1_data, { 2, 1, -1, -1, 0, 3 } },
{ DataLayout::bfzyx, { 0, 1, 2, -1, 3, 4 } },
{ DataLayout::bzyxf, { 1, 2, 3, -1, 0, 4 } },
{ DataLayout::fs_b_yx_fsv32, { 0, 1, -1, -1, 3, 2 } },
{ DataLayout::b_fs_yx_32fp, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::bfwzyx, { 0, 1, 2, 3, 4, 5 } },
{ DataLayout::nv12, { 0, 1, -1, -1, 2, 3 } },
{ DataLayout::image_2d_rgba, { 0, 1, -1, -1, 2, 3 } },
// X, Y, Z, W, U, V, F, B
{ DataLayout::f, { -1, -1, -1, -1, -1, -1, 0, -1 } },
{ DataLayout::bf, { -1, -1, -1, -1, -1, -1, 0, 1 } },
{ DataLayout::fb, { -1, -1, -1, -1, -1, -1, 1, 0 } },
{ DataLayout::bfyx, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::yxfb, { 2, 3, -1, -1, -1, -1, 1, 0 } },
{ DataLayout::byxf, { 1, 2, -1, -1, -1, -1, 0, 3 } },
{ DataLayout::fyxb, { 1, 2, -1, -1, -1, -1, 3, 0 } },
{ DataLayout::b_fs_yx_fsv2, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::b_fs_yx_fsv4, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::b_fs_yx_fsv16, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::b_fs_yx_fsv32, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::b_fs_zyx_fsv2, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::b_fs_zyx_fsv4, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::b_fs_zyx_fsv16, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::b_fs_zyx_fsv32, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv16_fsv32, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv16_fsv32, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::bs_fs_zyx_bsv16_fsv16, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv16_fsv16, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_yx_bsv4_fsv4, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_yx_bsv8_fsv4, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv8_fsv4, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv16_fsv4, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv16_fsv4, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv16_fsv2, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv16_fsv2, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv8_fsv2, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv8_fsv2, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv4_fsv2, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv32_fsv32, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv32_fsv32, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::bs_fs_zyx_bsv32_fsv16, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::bs_fs_yx_bsv32_fsv16, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::bs_f_bsv8__af8, { -1, -1, -1, -1, -1, -1, 0, 1 } },
{ DataLayout::bs_f_bsv16__af8, { -1, -1, -1, -1, -1, -1, 0, 1 } },
{ DataLayout::winograd_2x3_s1_data, { 2, 1, -1, -1, -1, -1, 0, 3 } },
{ DataLayout::bfzyx, { 0, 1, 2, -1, -1, -1, 3, 4 } },
{ DataLayout::bzyxf, { 1, 2, 3, -1, -1, -1, 0, 4 } },
{ DataLayout::fs_b_yx_fsv32, { 0, 1, -1, -1, -1, -1, 3, 2 } },
{ DataLayout::b_fs_yx_32fp, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::bfwzyx, { 0, 1, 2, 3, -1, -1, 4, 5 } },
{ DataLayout::bfuwzyx, { 0, 1, 2, 3, 4, -1, 5, 6 } },
{ DataLayout::bfvuwzyx, { 0, 1, 2, 3, 4, 5, 6, 7 } },
{ DataLayout::nv12, { 0, 1, -1, -1, -1, -1, 2, 3 } },
{ DataLayout::image_2d_rgba, { 0, 1, -1, -1, -1, -1, 2, 3 } },
}};
WeightsTensor::WeightsChannelArray WeightsTensor::weightsChannelArray {{
src/plugins/intel_gpu/src/kernel_selector/tensor_type.h
+8 -1
View File
@@ -4,6 +4,7 @@
#pragma once
#include "openvino/core/except.hpp"
#include "common_types.h"
#include "common_tools.h"
#include <vector>
@@ -66,6 +67,8 @@ enum DataLayout {
fs_b_yx_fsv32, // for FP16 kernels, 32 features to avoid partial writes
b_fs_yx_32fp, // bfyx with blocks of 16 packed binary input channels
bfwzyx, // batch, feature, 4D spatial
bfuwzyx, // batch, feature, 5D spatial
bfvuwzyx, // batch, feature, 6D spatial
nv12, // media nv12 layout
image_2d_rgba, // image2d RGBA
DataLayoutCount // NUMBER OF ELEMENTS IN ENUM
@@ -258,7 +261,7 @@ using NDims = std::vector<Dim>;
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
// extract code
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
enum class DataChannelName { X = 0, Y = 1, Z = 2, W = 3, FEATURE = 4, BATCH = 5, COUNT = 6 };
enum class DataChannelName { X = 0, Y = 1, Z = 2, W = 3, U = 4, V = 5, FEATURE = 6, BATCH = 7, COUNT = 8 };
enum class WeightsChannelName { X = 0, Y = 1, Z = 2, IFM = 3, OFM = 4, G = 5, COUNT = 6 };
@@ -292,6 +295,8 @@ inline bool SimpleLayout(DataLayout l) {
case DataLayout::bfzyx:
case DataLayout::bzyxf:
case DataLayout::bfwzyx:
case DataLayout::bfuwzyx:
case DataLayout::bfvuwzyx:
return true;
default:
return false;
@@ -611,6 +616,8 @@ struct DataTensor : public TensorBaseT<Datatype, DataLayout> {
Dim Y() const { return Extract(layout, DataChannelName::Y, dims); }
Dim Z() const { return Extract(layout, DataChannelName::Z, dims); }
Dim W() const { return Extract(layout, DataChannelName::W, dims); }
Dim U() const { return Extract(layout, DataChannelName::U, dims); }
Dim V() const { return Extract(layout, DataChannelName::V, dims); }
Dim Feature() const { return Extract(layout, DataChannelName::FEATURE, dims); }
Dim Batch() const { return Extract(layout, DataChannelName::BATCH, dims); }
src/plugins/intel_gpu/src/plugin/ops/constant.cpp
+10
View File
@@ -24,8 +24,18 @@ namespace ov {
namespace intel_gpu {
static cldnn::tensor getConstTensor(const ngraph::Shape constDims) {
std::vector<cldnn::tensor::value_type> shuffled_dims(constDims.size());
// cldnn tensor c-tor expects constants be in a reversed order (x, y, z, w, u, v)
for (size_t i = 0; i < constDims.size(); i++) {
shuffled_dims[i] = TensorValue(constDims[i < 2 ? i : (constDims.size() - 1 - i)]);
}
cldnn::tensor constTensor;
switch (constDims.size()) {
case 8:
case 7:
constTensor = cldnn::tensor(shuffled_dims);
break;
case 6: constTensor = cldnn::tensor(TensorValue(constDims[0]), TensorValue(constDims[1]),
TensorValue(constDims[5]), TensorValue(constDims[4]),
TensorValue(constDims[3]), TensorValue(constDims[2]));
src/plugins/intel_gpu/src/plugin/plugin.cpp
+1 -1
View File
@@ -319,7 +319,7 @@ QueryNetworkResult Plugin::QueryNetwork(const CNNNetwork& network,
return false;
auto pshape = node->get_output_partial_shape(0);
if (pshape.rank().is_dynamic())
if (pshape.rank().is_dynamic() || pshape.size() > cldnn::layout::max_rank())
return false;
int dynCount = 0;
src/plugins/intel_gpu/src/plugin/program.cpp
+12 -1
View File
@@ -358,8 +358,19 @@ std::shared_ptr<cldnn::program> Program::BuildProgram(const std::vector<std::sha
for (const auto& op : ops) {
if (op->is_dynamic()) {
allow_new_shape_infer = true;
break;
}
for (size_t i = 0; i < op->get_output_size(); i++) {
if (op->get_output_partial_shape(i).size() > 6)
allow_new_shape_infer = true;
}
for (size_t i = 0; i < op->get_input_size(); i++) {
if (op->get_input_partial_shape(i).size() > 6)
allow_new_shape_infer = true;
}
if (allow_new_shape_infer)
break;
}
m_config.set_property(ov::intel_gpu::partial_build_program(partialBuild));
src/plugins/intel_gpu/src/runtime/format.cpp
+103 -97
View File
@@ -21,50 +21,52 @@ static const std::map<format::type, format_traits> format_traits_map {
// Order - dims changing order from rare to often
// Inner order - dims order for internal storage in _sizes array
// Block sizes - vector of pairs of dimension number (by inner order) and block size ordered from rare to often
// Format B F S G Dims order Order Inner order Block sizes
FMT_TRAITS(yxfb, 1, 1, 2, 0, {2, 3, 1, 0}, "yxfb", "bfxy?", {}),
FMT_TRAITS(byxf, 1, 1, 2, 0, {0, 2, 3, 1}, "byxf", "bfxy?", {}),
FMT_TRAITS(bfyx, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {}),
FMT_TRAITS(fyxb, 1, 1, 2, 0, {1, 2, 3, 0}, "fyxb", "bfxy?", {}),
FMT_TRAITS(b_fs_yx_fsv2, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{1, 2}}),
FMT_TRAITS(b_fs_yx_fsv4, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{1, 4}}),
FMT_TRAITS(b_fs_yx_fsv16, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy", {{1, 16}}),
FMT_TRAITS(b_fs_yx_fsv32, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy", {{1, 32}}),
FMT_TRAITS(b_fs_zyx_fsv2, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{1, 2}}),
FMT_TRAITS(b_fs_zyx_fsv4, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{1, 4}}),
FMT_TRAITS(b_fs_zyx_fsv32, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{1, 32}}),
FMT_TRAITS(bs_xs_xsv8_bsv8, 1, 0, 1, 0, {0, 1}, "bx", "b?x??", {{2, 8}, {0, 8}}),
FMT_TRAITS(bs_xs_xsv8_bsv16, 1, 0, 1, 0, {0, 1}, "bx", "b?x??", {{2, 8}, {0, 16}}),
FMT_TRAITS(bs_x_bsv16, 1, 0, 1, 0, {0, 1}, "bx", "b?x??", {{0, 16}}),
FMT_TRAITS(winograd_2x3_s1_data, 1, 1, 2, 0, {0, 2, 3, 1}, "bxyf", "bfxy?", {}),
FMT_TRAITS(bzyxf, 1, 1, 3, 0, {0, 2, 3, 4, 1}, "bzyxf", "bfxyz", {}),
FMT_TRAITS(bfzyx, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {}),
FMT_TRAITS(bfwzyx, 1, 1, 4, 0, {0, 1, 2, 3, 4, 5}, "bfwzyx", "bfxyzw", {}),
FMT_TRAITS(fs_b_yx_fsv32, 1, 1, 2, 0, {1, 0, 2, 3}, "fbyx", "bfxy?", {{1, 32}}),
FMT_TRAITS(b_fs_yx_32fp, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {}),
FMT_TRAITS(b_fs_zyx_fsv16, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{1, 16}}),
FMT_TRAITS(bs_fs_zyx_bsv16_fsv32, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 16 }, {1, 32}}),
FMT_TRAITS(bs_fs_zyx_bsv16_fsv16, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 16 }, {1, 16}}),
FMT_TRAITS(bs_fs_yx_bsv16_fsv16, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 16 }, {1, 16}}),
FMT_TRAITS(bs_fs_yx_bsv16_fsv32, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 16 }, {1, 32}}),
FMT_TRAITS(bs_fs_yx_bsv4_fsv4, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 4 }, {1, 4}}),
FMT_TRAITS(bs_fs_yx_bsv8_fsv4, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 8 }, {1, 4}}),
FMT_TRAITS(bs_fs_zyx_bsv8_fsv4, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 8 }, {1, 4}}),
FMT_TRAITS(bs_fs_yx_bsv16_fsv4, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 16 }, {1, 4}}),
FMT_TRAITS(bs_fs_zyx_bsv16_fsv4, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 16 }, {1, 4}}),
FMT_TRAITS(bs_fs_yx_bsv16_fsv2, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 16 }, {1, 2}}),
FMT_TRAITS(bs_fs_zyx_bsv16_fsv2, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 16 }, {1, 2}}),
FMT_TRAITS(bs_fs_yx_bsv8_fsv2, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 8 }, {1, 2}}),
FMT_TRAITS(bs_fs_zyx_bsv8_fsv2, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 8 }, {1, 2}}),
FMT_TRAITS(bs_fs_yx_bsv4_fsv2, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 4 }, {1, 2}}),
FMT_TRAITS(bs_fs_zyx_bsv4_fsv4, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 4 }, {1, 4}}),
FMT_TRAITS(bs_fs_zyx_bsv4_fsv2, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 4 }, {1, 2}}),
FMT_TRAITS(bs_fs_zyx_bsv32_fsv32, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 32 }, {1, 32}}),
FMT_TRAITS(bs_fs_zyx_bsv32_fsv16, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 32 }, {1, 16}}),
FMT_TRAITS(bs_fs_yx_bsv32_fsv32, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 32 }, {1, 32}}),
FMT_TRAITS(bs_fs_yx_bsv32_fsv16, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 32 }, {1, 16}}),
FMT_TRAITS(nv12, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {}),
FMT_TRAITS(image_2d_rgba, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {}),
// Format B F S G Dims order Order Inner order Block sizes
FMT_TRAITS(yxfb, 1, 1, 2, 0, {2, 3, 1, 0}, "yxfb", "bfxy?", {}),
FMT_TRAITS(byxf, 1, 1, 2, 0, {0, 2, 3, 1}, "byxf", "bfxy?", {}),
FMT_TRAITS(bfyx, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {}),
FMT_TRAITS(fyxb, 1, 1, 2, 0, {1, 2, 3, 0}, "fyxb", "bfxy?", {}),
FMT_TRAITS(b_fs_yx_fsv2, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{1, 2}}),
FMT_TRAITS(b_fs_yx_fsv4, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{1, 4}}),
FMT_TRAITS(b_fs_yx_fsv16, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy", {{1, 16}}),
FMT_TRAITS(b_fs_yx_fsv32, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy", {{1, 32}}),
FMT_TRAITS(b_fs_zyx_fsv2, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{1, 2}}),
FMT_TRAITS(b_fs_zyx_fsv4, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{1, 4}}),
FMT_TRAITS(b_fs_zyx_fsv32, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{1, 32}}),
FMT_TRAITS(bs_fs_fsv8_bsv8, 1, 1, 0, 0, {0, 1}, "bf", "bf??", {{0, 8}, {1, 8}}),
FMT_TRAITS(bs_fs_fsv8_bsv16, 1, 1, 0, 0, {0, 1}, "bf", "bf??", {{0, 16}, {1, 8}}),
FMT_TRAITS(bs_f_bsv16, 1, 1, 0, 0, {0, 1}, "bf", "bf??", {{0, 16}}),
FMT_TRAITS(winograd_2x3_s1_data, 1, 1, 2, 0, {0, 2, 3, 1}, "bxyf", "bfxy?", {}),
FMT_TRAITS(bzyxf, 1, 1, 3, 0, {0, 2, 3, 4, 1}, "bzyxf", "bfxyz", {}),
FMT_TRAITS(bfzyx, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {}),
FMT_TRAITS(bfwzyx, 1, 1, 4, 0, {0, 1, 2, 3, 4, 5}, "bfwzyx", "bfxyzw", {}),
FMT_TRAITS(bfuwzyx, 1, 1, 5, 0, {0, 1, 2, 3, 4, 5, 6}, "bfuwzyx", "bfxyzwu", {}),
FMT_TRAITS(bfvuwzyx, 1, 1, 6, 0, {0, 1, 2, 3, 4, 5, 6, 7}, "bfvuwzyx", "bfxyzwuv", {}),
FMT_TRAITS(fs_b_yx_fsv32, 1, 1, 2, 0, {1, 0, 2, 3}, "fbyx", "bfxy?", {{1, 32}}),
FMT_TRAITS(b_fs_yx_32fp, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {}),
FMT_TRAITS(b_fs_zyx_fsv16, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{1, 16}}),
FMT_TRAITS(bs_fs_zyx_bsv16_fsv32, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 16 }, {1, 32}}),
FMT_TRAITS(bs_fs_zyx_bsv16_fsv16, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 16 }, {1, 16}}),
FMT_TRAITS(bs_fs_yx_bsv16_fsv16, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 16 }, {1, 16}}),
FMT_TRAITS(bs_fs_yx_bsv16_fsv32, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 16 }, {1, 32}}),
FMT_TRAITS(bs_fs_yx_bsv4_fsv4, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 4 }, {1, 4}}),
FMT_TRAITS(bs_fs_yx_bsv8_fsv4, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 8 }, {1, 4}}),
FMT_TRAITS(bs_fs_zyx_bsv8_fsv4, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 8 }, {1, 4}}),
FMT_TRAITS(bs_fs_yx_bsv16_fsv4, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 16 }, {1, 4}}),
FMT_TRAITS(bs_fs_zyx_bsv16_fsv4, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 16 }, {1, 4}}),
FMT_TRAITS(bs_fs_yx_bsv16_fsv2, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 16 }, {1, 2}}),
FMT_TRAITS(bs_fs_zyx_bsv16_fsv2, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 16 }, {1, 2}}),
FMT_TRAITS(bs_fs_yx_bsv8_fsv2, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 8 }, {1, 2}}),
FMT_TRAITS(bs_fs_zyx_bsv8_fsv2, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 8 }, {1, 2}}),
FMT_TRAITS(bs_fs_yx_bsv4_fsv2, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 4 }, {1, 2}}),
FMT_TRAITS(bs_fs_zyx_bsv4_fsv4, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 4 }, {1, 4}}),
FMT_TRAITS(bs_fs_zyx_bsv4_fsv2, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 4 }, {1, 2}}),
FMT_TRAITS(bs_fs_zyx_bsv32_fsv32, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 32 }, {1, 32}}),
FMT_TRAITS(bs_fs_zyx_bsv32_fsv16, 1, 1, 3, 0, {0, 1, 2, 3, 4}, "bfzyx", "bfxyz", {{0, 32 }, {1, 16}}),
FMT_TRAITS(bs_fs_yx_bsv32_fsv32, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 32 }, {1, 32}}),
FMT_TRAITS(bs_fs_yx_bsv32_fsv16, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {{0, 32 }, {1, 16}}),
FMT_TRAITS(nv12, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {}),
FMT_TRAITS(image_2d_rgba, 1, 1, 2, 0, {0, 1, 2, 3}, "bfyx", "bfxy?", {}),
FMT_TRAITS(oiyx, 1, 1, 2, 0, {0, 1, 2, 3}, "oiyx", "oixy", {}),
FMT_TRAITS(ioyx, 1, 1, 2, 0, {1, 0, 2, 3}, "ioyx", "oixy", {}),
@@ -156,59 +158,59 @@ static const std::map<format::type, format_traits> format_traits_map {
FMT_TRAITS(os_zy_is_x_osv8_isv2, 1, 1, 3, 0, {0, 2, 3, 1, 4}, "ozyix", "oixyz", {{0, 8}, {1, 2}}),
FMT_TRAITS(os_zy_is_x_osv8_isv4, 1, 1, 3, 0, {0, 2, 3, 1, 4}, "ozyix", "oixyz", {{0, 8}, {1, 4}}),
FMT_TRAITS(goiyx, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {}),
FMT_TRAITS(gioyx, 1, 1, 2, 1, {0, 2, 1, 3, 4}, "gioyx", "oixy??g", {}),
FMT_TRAITS(goizyx, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz?g", {}),
FMT_TRAITS(giozyx, 1, 1, 3, 1, {0, 2, 1, 3, 4, 5}, "giozyx", "oixyz?g", {}),
FMT_TRAITS(g_os_iyx_osv8, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{0, 8}}),
FMT_TRAITS(g_os_iyx_osv16, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{0, 16}}),
FMT_TRAITS(g_os_iyx_osv32, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{0, 32}}),
FMT_TRAITS(gs_oiyx_gsv16, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{6, 16}}),
FMT_TRAITS(gs_oizyx_gsv16, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz?g", {{6, 16}}),
FMT_TRAITS(gs_oiyx_gsv32, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{6, 32}}),
FMT_TRAITS(gs_oizyx_gsv32, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz?g", {{6, 32}}),
FMT_TRAITS(gyxio, 1, 1, 2, 1, {0, 3, 4, 2, 1}, "gyxio", "oixy??g", {}),
FMT_TRAITS(g_is_os_zyx_isv16_osv16, 1, 1, 3, 1, {0, 2, 1, 3, 4, 5}, "giozyx", "oixyz?g", {{1, 16}, {0, 16}}),
FMT_TRAITS(g_is_os_yx_isv16_osv16, 1, 1, 2, 1, {0, 2, 1, 3, 4}, "gioyx", "oixy??g", {{1, 16}, {0, 16}}),
FMT_TRAITS(g_os_is_zyx_isv8_osv16_isv2, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz?g", {{1, 8}, {0, 16}, {1, 2}}),
FMT_TRAITS(g_os_is_yx_isv8_osv16_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{1, 8}, {0, 16}, {1, 2}}),
FMT_TRAITS(g_os_is_zyx_isv16_osv16, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz?g", {{0, 16}, {1, 16}}),
FMT_TRAITS(g_os_zy_is_x_osv8_isv2, 1, 1, 3, 1, {0, 1, 3, 4, 2, 5}, "gozyix", "oixyz?g", {{0, 8}, {1, 2}}),
FMT_TRAITS(g_os_zy_is_x_osv8_isv4, 1, 1, 3, 1, {0, 1, 3, 4, 2, 5}, "gozyix", "oixyz?g", {{0, 8}, {1, 4}}),
FMT_TRAITS(g_os_zyx_is_osv8_isv2, 1, 1, 3, 1, {0, 1, 3, 4, 5, 2}, "gozyxi", "oixyz?g", {{0, 8}, {1, 2}}),
FMT_TRAITS(g_os_zyx_is_osv8_isv4, 1, 1, 3, 1, {0, 1, 3, 4, 5, 2}, "gozyxi", "oixyz?g", {{0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_yx_osv8_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_yx_osv8_isv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_yx_osv16_isv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{0, 16}, {1, 4}}),
FMT_TRAITS(g_os_is_zyx_osv16_isv16, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz?g", {{0, 16}, {1, 16}}),
FMT_TRAITS(g_os_zyx_is_osv16_isv4, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz?g", {{0, 16}, {1, 4}}),
FMT_TRAITS(g_os_zyx_is_osv16_isv16, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz?g", {{0, 16}, {1, 16}}),
FMT_TRAITS(g_os_zyx_is_osv16_isv32, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz?g", {{0, 16}, {1, 32}}),
FMT_TRAITS(g_os_zyx_is_osv32_isv4, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz?g", {{0, 32}, {1, 4}}),
FMT_TRAITS(g_os_zyx_is_osv32_isv16, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz?g", {{0, 32}, {1, 16}}),
FMT_TRAITS(g_os_zyx_is_osv32_isv32, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz?g", {{0, 32}, {1, 32}}),
FMT_TRAITS(g_os_is_yx_isa8_osv8_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{1, 8}, {0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_yx_isa8_osv8_isv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{1, 8}, {0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_zyx_isa8_osv8_isv2, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz?g", {{1, 8}, {0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_zyx_isa8_osv8_isv4, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz?g", {{1, 8}, {0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_yx_osa2_isa8_osv8_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{0, 2}, {1, 8}, {0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_yx_osa4_isa8_osv8_isv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{0, 4}, {1, 8}, {0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_zyx_osa4_isa8_osv8_isv4, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz?g", {{0, 4}, {1, 8}, {0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_yx_osa4_isa8_osv8_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{0, 4}, {1, 8}, {0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_zyx_osa4_isa8_osv8_isv2, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz?g", {{0, 4}, {1, 8}, {0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_yx_osa2_isa8_osv16_isv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{0, 2}, {1, 8}, {0, 16}, {1, 4}}),
FMT_TRAITS(g_os_is_yx_osa2_isa8_osv16_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{0, 2}, {1, 8}, {0, 16}, {1, 2}}),
FMT_TRAITS(gs_oi_yxs_gsv4_yxsv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{6, 4}}),
FMT_TRAITS(gs_oi_yxs_gsv16_yxsv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{6, 16}}),
FMT_TRAITS(gs_oi_yxs_gsv32_yxsv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{6, 32}}),
FMT_TRAITS(g_os_is_yx_isv16_osv16, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy??g", {{1, 16}, {0, 16}}),
FMT_TRAITS(gi_yxs_os_yxsv2_osv16, 1, 1, 2, 1, {0, 2, 3, 4, 1}, "giyxo", "oixy??g", {{0, 16}}),
FMT_TRAITS(giy_xs_os_xsv2_osv8__ao32, 1, 1, 2, 1, {0, 2, 3, 4, 1}, "giyxo", "oixy??g", {{2, 2}, {0, 8}}),
FMT_TRAITS(giy_xs_os_xsv2_osv16__ao32, 1, 1, 2, 1, {0, 2, 3, 4, 1}, "giyxo", "oixy??g", {{2, 2}, {0, 16}}),
FMT_TRAITS(g_os_yx_is_osv8_isv2, 1, 1, 2, 1, {0, 1, 3, 4, 2}, "goyxi", "oixy??g", {{0, 8}, {1, 2}}),
FMT_TRAITS(g_os_yx_is_osv8_isv4, 1, 1, 2, 1, {0, 1, 3, 4, 2}, "goyxi", "oixy??g", {{0, 8}, {1, 4}}),
FMT_TRAITS(g_os_y_is_x_osv8_isv2, 1, 1, 2, 1, {0, 1, 3, 2, 4}, "goyix", "oixy??g", {{0, 8}, {1, 2}}),
FMT_TRAITS(g_os_y_is_x_osv8_isv4, 1, 1, 2, 1, {0, 1, 3, 2, 4}, "goyix", "oixy??g", {{0, 8}, {1, 4}}),
FMT_TRAITS(goiyx, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {}),
FMT_TRAITS(gioyx, 1, 1, 2, 1, {0, 2, 1, 3, 4}, "gioyx", "oixy????g", {}),
FMT_TRAITS(goizyx, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz???g", {}),
FMT_TRAITS(giozyx, 1, 1, 3, 1, {0, 2, 1, 3, 4, 5}, "giozyx", "oixyz???g", {}),
FMT_TRAITS(g_os_iyx_osv8, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{0, 8}}),
FMT_TRAITS(g_os_iyx_osv16, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{0, 16}}),
FMT_TRAITS(g_os_iyx_osv32, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{0, 32}}),
FMT_TRAITS(gs_oiyx_gsv16, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{6, 16}}),
FMT_TRAITS(gs_oizyx_gsv16, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz???g", {{6, 16}}),
FMT_TRAITS(gs_oiyx_gsv32, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{6, 32}}),
FMT_TRAITS(gs_oizyx_gsv32, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz???g", {{6, 32}}),
FMT_TRAITS(gyxio, 1, 1, 2, 1, {0, 3, 4, 2, 1}, "gyxio", "oixy????g", {}),
FMT_TRAITS(g_is_os_zyx_isv16_osv16, 1, 1, 3, 1, {0, 2, 1, 3, 4, 5}, "giozyx", "oixyz???g", {{1, 16}, {0, 16}}),
FMT_TRAITS(g_is_os_yx_isv16_osv16, 1, 1, 2, 1, {0, 2, 1, 3, 4}, "gioyx", "oixy????g", {{1, 16}, {0, 16}}),
FMT_TRAITS(g_os_is_zyx_isv8_osv16_isv2, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz???g", {{1, 8}, {0, 16}, {1, 2}}),
FMT_TRAITS(g_os_is_yx_isv8_osv16_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{1, 8}, {0, 16}, {1, 2}}),
FMT_TRAITS(g_os_is_zyx_isv16_osv16, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz???g", {{0, 16}, {1, 16}}),
FMT_TRAITS(g_os_zy_is_x_osv8_isv2, 1, 1, 3, 1, {0, 1, 3, 4, 2, 5}, "gozyix", "oixyz???g", {{0, 8}, {1, 2}}),
FMT_TRAITS(g_os_zy_is_x_osv8_isv4, 1, 1, 3, 1, {0, 1, 3, 4, 2, 5}, "gozyix", "oixyz???g", {{0, 8}, {1, 4}}),
FMT_TRAITS(g_os_zyx_is_osv8_isv2, 1, 1, 3, 1, {0, 1, 3, 4, 5, 2}, "gozyxi", "oixyz???g", {{0, 8}, {1, 2}}),
FMT_TRAITS(g_os_zyx_is_osv8_isv4, 1, 1, 3, 1, {0, 1, 3, 4, 5, 2}, "gozyxi", "oixyz???g", {{0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_yx_osv8_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_yx_osv8_isv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_yx_osv16_isv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{0, 16}, {1, 4}}),
FMT_TRAITS(g_os_is_zyx_osv16_isv16, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz???g", {{0, 16}, {1, 16}}),
FMT_TRAITS(g_os_zyx_is_osv16_isv4, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz???g", {{0, 16}, {1, 4}}),
FMT_TRAITS(g_os_zyx_is_osv16_isv16, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz???g", {{0, 16}, {1, 16}}),
FMT_TRAITS(g_os_zyx_is_osv16_isv32, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz???g", {{0, 16}, {1, 32}}),
FMT_TRAITS(g_os_zyx_is_osv32_isv4, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz???g", {{0, 32}, {1, 4}}),
FMT_TRAITS(g_os_zyx_is_osv32_isv16, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz???g", {{0, 32}, {1, 16}}),
FMT_TRAITS(g_os_zyx_is_osv32_isv32, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "gozyxi", "oixyz???g", {{0, 32}, {1, 32}}),
FMT_TRAITS(g_os_is_yx_isa8_osv8_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{1, 8}, {0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_yx_isa8_osv8_isv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{1, 8}, {0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_zyx_isa8_osv8_isv2, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz???g", {{1, 8}, {0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_zyx_isa8_osv8_isv4, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz???g", {{1, 8}, {0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_yx_osa2_isa8_osv8_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{0, 2}, {1, 8}, {0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_yx_osa4_isa8_osv8_isv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{0, 4}, {1, 8}, {0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_zyx_osa4_isa8_osv8_isv4, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz???g", {{0, 4}, {1, 8}, {0, 8}, {1, 4}}),
FMT_TRAITS(g_os_is_yx_osa4_isa8_osv8_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{0, 4}, {1, 8}, {0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_zyx_osa4_isa8_osv8_isv2, 1, 1, 3, 1, {0, 1, 2, 3, 4, 5}, "goizyx", "oixyz???g", {{0, 4}, {1, 8}, {0, 8}, {1, 2}}),
FMT_TRAITS(g_os_is_yx_osa2_isa8_osv16_isv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{0, 2}, {1, 8}, {0, 16}, {1, 4}}),
FMT_TRAITS(g_os_is_yx_osa2_isa8_osv16_isv2, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{0, 2}, {1, 8}, {0, 16}, {1, 2}}),
FMT_TRAITS(gs_oi_yxs_gsv4_yxsv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{6, 4}}),
FMT_TRAITS(gs_oi_yxs_gsv16_yxsv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{6, 16}}),
FMT_TRAITS(gs_oi_yxs_gsv32_yxsv4, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{6, 32}}),
FMT_TRAITS(g_os_is_yx_isv16_osv16, 1, 1, 2, 1, {0, 1, 2, 3, 4}, "goiyx", "oixy????g", {{1, 16}, {0, 16}}),
FMT_TRAITS(gi_yxs_os_yxsv2_osv16, 1, 1, 2, 1, {0, 2, 3, 4, 1}, "giyxo", "oixy????g", {{0, 16}}),
FMT_TRAITS(giy_xs_os_xsv2_osv8__ao32, 1, 1, 2, 1, {0, 2, 3, 4, 1}, "giyxo", "oixy????g", {{2, 2}, {0, 8}}),
FMT_TRAITS(giy_xs_os_xsv2_osv16__ao32, 1, 1, 2, 1, {0, 2, 3, 4, 1}, "giyxo", "oixy????g", {{2, 2}, {0, 16}}),
FMT_TRAITS(g_os_yx_is_osv8_isv2, 1, 1, 2, 1, {0, 1, 3, 4, 2}, "goyxi", "oixy????g", {{0, 8}, {1, 2}}),
FMT_TRAITS(g_os_yx_is_osv8_isv4, 1, 1, 2, 1, {0, 1, 3, 4, 2}, "goyxi", "oixy????g", {{0, 8}, {1, 4}}),
FMT_TRAITS(g_os_y_is_x_osv8_isv2, 1, 1, 2, 1, {0, 1, 3, 2, 4}, "goyix", "oixy????g", {{0, 8}, {1, 2}}),
FMT_TRAITS(g_os_y_is_x_osv8_isv4, 1, 1, 2, 1, {0, 1, 3, 2, 4}, "goyix", "oixy????g", {{0, 8}, {1, 4}}),
};
const format_traits& format::traits(type fmt) {
@@ -244,6 +246,10 @@ format format::get_default_format(size_t rank, bool is_weights, bool is_grouped)
default_fmt = cldnn::format::bfzyx;
} else if (rank == 6) {
default_fmt = cldnn::format::bfwzyx;
} else if (rank == 7) {
default_fmt = cldnn::format::bfuwzyx;
} else if (rank == 8) {
default_fmt = cldnn::format::bfvuwzyx;
}
}
return default_fmt;
src/plugins/intel_gpu/src/runtime/layout.cpp
+31 -83
View File
@@ -12,8 +12,8 @@ namespace cldnn {
static inline bool check_redundant_1d_along_feature(layout const& l1, layout const& l2);
namespace {
std::vector<cldnn::tensor::value_type> convert_dimensions(const std::vector<cldnn::tensor::value_type>& sizes, std::string in_order, std::string out_order) {
std::vector<cldnn::tensor::value_type> new_sizes(out_order.size(), {-1});
std::vector<int32_t> convert_dimensions(const std::vector<int32_t>& sizes, std::string in_order, std::string out_order) {
std::vector<int32_t> new_sizes(out_order.size(), {-1});
for (size_t out_idx = 0; out_idx < out_order.size(); ++out_idx) {
auto channel = out_order[out_idx];
if (channel == '?')
@@ -153,7 +153,7 @@ static format to_weights_format(format f, bool is_grouped) {
}
case format::b_fs_yx_fsv16:
return format::o_is_yx_isv16;
case format::bs_xs_xsv8_bsv8:
case format::bs_fs_fsv8_bsv8:
return format::os_i_osv8__ai8;
case format::b_fs_yx_32fp:
return format::os_is_yx_osv32_isv32p;
@@ -518,97 +518,45 @@ ov::PartialShape layout::transform(const ov::PartialShape& pshape, cldnn::format
return pshape;
}
cldnn::tensor::value_type default_size = -1;
int32_t default_size = -1;
auto shape = pshape.to_shape();
std::vector<tensor::value_type> dims;
std::vector<int32_t> dims;
for (auto dim : shape) {
dims.push_back(static_cast<tensor::value_type>(dim));
dims.push_back(static_cast<int32_t>(dim));
}
const cldnn::format default_fmt = cldnn::format::bfwzyx;
auto old_sizes = convert_dimensions(dims, old_fmt.order(), default_fmt.internal_order()); // convert to internal order (bfxyzw)
const cldnn::format default_fmt = cldnn::format::bfvuwzyx;
auto old_sizes = convert_dimensions(dims, old_fmt.order(), default_fmt.internal_order()); // convert to internal order (bfxyzwuv)
auto val_order = default_fmt.internal_order();
auto new_order = new_fmt.internal_order();
const auto& new_traits = format::traits(new_fmt);
std::vector<tensor::value_type> new_sizes(old_sizes.size(), {default_size});
auto tmp = 1;
auto tmp_z = 1;
auto tmp_w = 1;
std::vector<int32_t> new_sizes(old_sizes.size(), {default_size});
static const std::map<char, char> flatten_mapping = {
{ 'v', 'u'},
{ 'u', 'w'},
{ 'w', 'z'},
{ 'z', 'y'}
};
for (size_t i = 0; i < default_fmt.order().size(); i++) {
auto c = val_order[i]; //bfxywz
// skip f, y, z, and w for the formats that do not have it
if (((new_fmt == format::bs_xs_xsv8_bsv8) ||
(new_fmt == format::bs_xs_xsv8_bsv16) ||
(new_fmt == format::os_i_osv8__ai8) ||
(new_fmt == format::os_i_osv16__ai8) ||
(new_fmt == format::bs_x_bsv16)) &&
((c == 'f') ||
(c == 'y') ||
(c == 'z') ||
(c == 'w'))) {
if (new_order[i] == '?')
new_sizes[i] = default_size;
tmp *= old_sizes[i]; //0f0ywz
continue;
}
// skip z for the formats that do not have it
if (((new_fmt != format::bfzyx && new_fmt != format::b_fs_zyx_fsv16 && new_fmt != format::b_fs_zyx_fsv32 &&
new_fmt != format::bfwzyx && new_fmt != format::bs_fs_zyx_bsv16_fsv16 && new_fmt != format::bs_fs_zyx_bsv16_fsv32 &&
new_fmt != format::bs_fs_zyx_bsv32_fsv16 && new_fmt != format::bs_fs_zyx_bsv32_fsv32 &&
new_fmt != format::b_fs_zyx_fsv2 && new_fmt != format::b_fs_zyx_fsv4 &&
new_fmt != format::bs_fs_zyx_bsv8_fsv2 && new_fmt != format::bs_fs_zyx_bsv8_fsv4)) && (c == 'z')) {
if (new_order[i] == '?')
new_sizes[i] = default_size;
tmp_z *= old_sizes[i]; //00000z
continue;
}
if (new_fmt != format::bfwzyx && c == 'w') {
if (new_order[i] == '?')
new_sizes[i] = default_size;
if (new_fmt == format::bfzyx || new_fmt == format::b_fs_zyx_fsv16 ||
new_fmt == format::bs_fs_zyx_bsv16_fsv16 || new_fmt == format::b_fs_zyx_fsv32 ||
new_fmt == format::bs_fs_zyx_bsv16_fsv32)
tmp_w *= old_sizes[i]; //0000w0
else
tmp_z *= old_sizes[i]; //0000w0
continue;
}
auto new_pos = new_order.find(c);
if (new_pos == std::string::npos)
throw std::invalid_argument("cannot convert to new format");
new_sizes[new_pos] = old_sizes[i];
}
// in case of formats with smaller number of dimensions than input, flatten is performed below
if (tmp != 1 || tmp_z != 1 || tmp_w != 1) {
for (size_t i = 0; i < default_fmt.order().size(); i++) {
auto c = val_order[i];
if (c == 'x') {
auto new_pos = new_order.find(c);
new_sizes[new_pos] *= tmp;
auto target_dim = val_order[i]; //bfxywzuv
while (!new_traits.has_dimension(target_dim)) {
if (flatten_mapping.find(target_dim) != flatten_mapping.end()) {
target_dim = flatten_mapping.at(target_dim);
} else {
target_dim = new_fmt.order().back();
}
}
if (c == 'y') {
auto new_pos = new_order.find(c);
if (new_pos != std::string::npos)
new_sizes[new_pos] *= tmp_z;
}
if (c == 'z') {
auto new_pos = new_order.find(c);
if (new_pos != std::string::npos)
new_sizes[new_pos] *= tmp_w;
auto new_pos = new_order.find(target_dim);
if (new_pos != std::string::npos) {
if (new_sizes[new_pos] == -1) {
new_sizes[new_pos] = old_sizes[i];
} else {
new_sizes[new_pos] *= old_sizes[i];
}
}
}
src/plugins/intel_gpu/tests/module_tests/layout_test.cpp
+64
View File
@@ -241,3 +241,67 @@ INSTANTIATE_TEST_SUITE_P(smoke, layout_cmp_test,
{layout{ov::PartialShape{9, 17, 3, 2, 5}, data_types::f16, format::is_os_zyx_isa8_osv8_isv2},
layout{ov::PartialShape{9, 17, 3, 2, 5}, data_types::f16, format::os_is_zyx_isa8_osv8_isv2}, false, false},
}));
struct layouts_transform_test_params {
format::type from;
format::type to;
ov::PartialShape shape;
ov::PartialShape expected;
};
class layout_transform_test : public testing::TestWithParam<layouts_transform_test_params> { };
TEST_P(layout_transform_test, basic) {
auto p = GetParam();
ASSERT_EQ(layout::transform(p.shape, p.from, p.to), p.expected)
<< "from=" << fmt_to_str(p.from) << " to=" << fmt_to_str(p.to) << " shape=" << p.shape;
}
INSTANTIATE_TEST_SUITE_P(smoke, layout_transform_test,
testing::ValuesIn(std::vector<layouts_transform_test_params>{
{format::yxfb, format::bfyx, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{4, 3, 1, 2}},
{format::bfyx, format::yxfb, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{3, 4, 2, 1}},
{format::bfyx, format::bs_f_bsv16, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{1, 2*3*4}},
{format::bs_f_bsv16, format::bfyx, ov::PartialShape{1, 2*3*4}, ov::PartialShape{1, 2*3*4, 1, 1}},
{format::bfyx, format::bs_fs_yx_bsv16_fsv16, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{1, 2, 3, 4}},
{format::bfyx, format::bfzyx, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{1, 2, 1, 3, 4}},
{format::bfyx, format::bfwzyx, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{1, 2, 1, 1, 3, 4}},
{format::bfyx, format::bfuwzyx, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{1, 2, 1, 1, 1, 3, 4}},
{format::bfyx, format::bfvuwzyx, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{1, 2, 1, 1, 1, 1, 3, 4}},
{format::b_fs_yx_fsv16, format::bfyx, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{1, 2, 3, 4}},
{format::b_fs_yx_fsv16, format::bfzyx, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{1, 2, 1, 3, 4}},
{format::b_fs_yx_fsv16, format::bfwzyx, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{1, 2, 1, 1, 3, 4}},
{format::b_fs_yx_fsv16, format::bfuwzyx, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{1, 2, 1, 1, 1, 3, 4}},
{format::b_fs_yx_fsv16, format::bfvuwzyx, ov::PartialShape{1, 2, 3, 4}, ov::PartialShape{1, 2, 1, 1, 1, 1, 3, 4}},
{format::bfzyx, format::b_fs_zyx_fsv16, ov::PartialShape{1, 2, 3, 4, 5}, ov::PartialShape{1, 2, 3, 4, 5}},
{format::bfzyx, format::bfwzyx, ov::PartialShape{1, 2, 3, 4, 5}, ov::PartialShape{1, 2, 1, 3, 4, 5}},
{format::bfzyx, format::bfuwzyx, ov::PartialShape{1, 2, 3, 4, 5}, ov::PartialShape{1, 2, 1, 1, 3, 4, 5}},
{format::bfzyx, format::bfvuwzyx, ov::PartialShape{1, 2, 3, 4, 5}, ov::PartialShape{1, 2, 1, 1, 1, 3, 4, 5}},
{format::b_fs_zyx_fsv16, format::bfzyx, ov::PartialShape{1, 2, 3, 4, 5}, ov::PartialShape{1, 2, 3, 4, 5}},
{format::b_fs_zyx_fsv16, format::bfwzyx, ov::PartialShape{1, 2, 3, 4, 5}, ov::PartialShape{1, 2, 1, 3, 4, 5}},
{format::b_fs_zyx_fsv16, format::bfuwzyx, ov::PartialShape{1, 2, 3, 4, 5}, ov::PartialShape{1, 2, 1, 1, 3, 4, 5}},
{format::b_fs_zyx_fsv16, format::bfvuwzyx, ov::PartialShape{1, 2, 3, 4, 5}, ov::PartialShape{1, 2, 1, 1, 1, 3, 4, 5}},
{format::bfwzyx, format::bfuwzyx, ov::PartialShape{1, 2, 3, 4, 5, 6}, ov::PartialShape{1, 2, 1, 3, 4, 5, 6}},
{format::bfwzyx, format::bfvuwzyx, ov::PartialShape{1, 2, 3, 4, 5, 6}, ov::PartialShape{1, 2, 1, 1, 3, 4, 5, 6}},
{format::bfuwzyx, format::bfvuwzyx, ov::PartialShape{1, 2, 3, 4, 5, 6, 7}, ov::PartialShape{1, 2, 1, 3, 4, 5, 6, 7}},
{format::bfvuwzyx, format::bfuwzyx, ov::PartialShape{1, 2, 3, 4, 5, 6, 7, 8}, ov::PartialShape{1, 2, 3*4, 5, 6, 7, 8}},
{format::bfvuwzyx, format::bfuwzyx, ov::PartialShape{1, 2, 3, 4, 5, 6, 7, 8}, ov::PartialShape{1, 2, 3*4, 5, 6, 7, 8}},
{format::bfvuwzyx, format::bfzyx, ov::PartialShape{1, 2, 3, 4, 5, 6, 7, 8}, ov::PartialShape{1, 2, 3*4*5*6, 7, 8}},
{format::bfvuwzyx, format::bfyx, ov::PartialShape{1, 2, 3, 4, 5, 6, 7, 8}, ov::PartialShape{1, 2, 3*4*5*6*7, 8}},
{format::bfuwzyx, format::bfwzyx, ov::PartialShape{1, 2, 3, 4, 5, 6, 7}, ov::PartialShape{1, 2, 3*4, 5, 6, 7}},
{format::bfuwzyx, format::bfzyx, ov::PartialShape{1, 2, 3, 4, 5, 6, 7}, ov::PartialShape{1, 2, 3*4*5, 6, 7}},
{format::bfuwzyx, format::bfyx, ov::PartialShape{1, 2, 3, 4, 5, 6, 7}, ov::PartialShape{1, 2, 3*4*5*6, 7}},
{format::bfwzyx, format::bfzyx, ov::PartialShape{1, 2, 3, 4, 5, 6}, ov::PartialShape{1, 2, 3*4, 5, 6}},
{format::bfwzyx, format::bfyx, ov::PartialShape{1, 2, 3, 4, 5, 6}, ov::PartialShape{1, 2, 3*4*5, 6}},
{format::bfzyx, format::bfyx, ov::PartialShape{1, 2, 3, 4, 5}, ov::PartialShape{1, 2, 3*4, 5}},
}));
src/plugins/intel_gpu/tests/test_cases/canonicalization_gpu_test.cpp
+3 -3
View File
@@ -60,9 +60,9 @@ layout create_default_layout(const ov::PartialShape& pshape) {
}
std::vector<Shapes> shape_of_shapes {
{{{3}}, {{3, 1, 1, 1, 1, 1}}, {{1, 1, 1, 1, 1, 1}}},
{{{1, 2, 3}}, {{1, 2, 3, 1, 1, 1}}, {{3, 1, 1, 1, 1, 1}}},
{{{1, 2, 3, 4, 5}}, {{1, 2, 3, 4, 5, 1}}, {{5, 1, 1, 1, 1, 1}}}
{{{3}}, {{3, 1, 1, 1, 1, 1, 1, 1}}, {{1, 1, 1, 1, 1, 1, 1, 1}}},
{{{1, 2, 3}}, {{1, 2, 3, 1, 1, 1, 1, 1}}, {{3, 1, 1, 1, 1, 1, 1, 1}}},
{{{1, 2, 3, 4, 5}}, {{1, 2, 3, 4, 5, 1, 1, 1}}, {{5, 1, 1, 1, 1, 1, 1, 1}}}
};
TEST(canonicalization, shape_of) {
src/plugins/intel_gpu/tests/test_cases/eltwise_gpu_test.cpp
+41
View File
@@ -1474,6 +1474,47 @@ TEST(eltwise_gpu_f32, dynamic_kernel_broadcast_mixed_ranks_5d_2d) {
}
}
TEST(eltwise_gpu_f32, add_basic_8d) {
auto& engine = get_test_engine();
auto input1 = engine.allocate_memory({{1, 3, 2, 2, 2, 3, 2, 3}, data_types::f32, format::bfvuwzyx });
auto input2 = engine.allocate_memory({{1, 3, 2, 2, 2, 3, 2, 3}, data_types::f32, format::bfvuwzyx });
topology topology;
topology.add(input_layout("input1", input1->get_layout()));
topology.add(input_layout("input2", input2->get_layout()));
topology.add(eltwise("eltwise", { input_info("input1"), input_info("input2") }, eltwise_mode::sum));
{
mem_lock<float> lock1(input1, get_test_stream());
mem_lock<float> lock2(input2, get_test_stream());
std::iota(lock1.begin(), lock1.end(), 0);
std::iota(lock2.begin(), lock2.end(), 0);
}
auto config = get_test_default_config(engine);
config.set_property(ov::intel_gpu::allow_new_shape_infer(true));
network network(engine, topology, config);
network.set_input_data("input1", input1);
network.set_input_data("input2", input2);
auto outputs = network.execute();
ASSERT_EQ(outputs.size(), size_t(1));
ASSERT_EQ(outputs.begin()->first, "eltwise");
auto output = outputs.at("eltwise").get_memory();
ov::PartialShape expected_shape = {1, 3, 2, 2, 2, 3, 2, 3};
ASSERT_EQ(output->get_layout().get_partial_shape(), expected_shape);
cldnn::mem_lock<float> output_ptr(output, get_test_stream());
for (size_t i = 0; i < output->count(); i++) {
ASSERT_EQ(2.f*i, output_ptr[i]) << " i = " << i;
}
}
TEST(eltwise_gpu_f32, add_basic_in4x4x2x2) {
// Input2 : 2x2x2
// Input : 2x2x2x2
src/plugins/intel_gpu/tests/test_cases/hash_key_gpu_test.cpp
+28 -28
View File
@@ -26,8 +26,8 @@ using namespace tests;
TEST(check_hash_value, eltwise_basic) {
auto& engine = get_test_engine();
auto input1 = engine.allocate_memory({ data_types::f32, format::bfyx, { 2, 2, 2, 2 } });
auto input2 = engine.allocate_memory({ data_types::f32, format::bfyx, { 2, 2, 2, 2 } });
auto input1 = engine.allocate_memory({ { 2, 2, 2, 2 }, data_types::f32, format::bfyx });
auto input2 = engine.allocate_memory({ { 2, 2, 2, 2 }, data_types::f32, format::bfyx });
auto key_prim_id = "eltwise";
topology topology;
@@ -45,7 +45,7 @@ TEST(check_hash_value, eltwise_basic) {
const auto params_hash = prog_node.get_kernel_impl_params()->hash();
ASSERT_EQ(primitive_hash, 11385140218618178073UL);
ASSERT_EQ(params_hash, 10460622021476296271UL);
ASSERT_EQ(params_hash, 1027235386189301125UL);
}
TEST(check_hash_value, fc_basic) {
@@ -53,9 +53,9 @@ TEST(check_hash_value, fc_basic) {
const int32_t b = 1, in_f = 128, in_x = 1, in_y = 1, out_f = 65;
auto input_prim = engine.allocate_memory({ data_types::f32, format::bfyx, { b, in_f, in_y, in_x } });
auto weights_prim = engine.allocate_memory({ data_types::f32, format::bfyx, { out_f, in_f, in_y, in_x } });
auto bias_prim = engine.allocate_memory({ data_types::f32, format::bfyx, { 1, 1, out_f, 1 } });
auto input_prim = engine.allocate_memory({ { b, in_f, in_y, in_x }, data_types::f32, format::bfyx });
auto weights_prim = engine.allocate_memory({ { out_f, in_f, in_y, in_x }, data_types::f32, format::bfyx });
auto bias_prim = engine.allocate_memory({ { 1, 1, out_f, 1 }, data_types::f32, format::bfyx });
const auto key_prim_id = "fc";
topology topology(
@@ -75,14 +75,14 @@ TEST(check_hash_value, fc_basic) {
const auto params_hash = prog_node.type()->get_fake_aligned_params(*prog_node.get_kernel_impl_params()).hash();
ASSERT_EQ(primitive_hash, 2197080758510296176UL);
ASSERT_EQ(params_hash, 5241462399408562393UL);
ASSERT_EQ(params_hash, 11955559181519986341UL);
}
TEST(check_hash_value, gather_basic) {
auto& engine = get_test_engine();
auto input1 = engine.allocate_memory({ data_types::f16, format::bfzyx, tensor{ 3, 2, 2, 4, 3} }); // Dictionary
auto input2 = engine.allocate_memory({ data_types::f32, format::bfyx, tensor{ 3, 2, 1, 3 } }); // Indexes
auto input1 = engine.allocate_memory({ { 3, 2, 2, 4, 3}, data_types::f16, format::bfzyx }); // Dictionary
auto input2 = engine.allocate_memory({ { 3, 2, 1, 3 }, data_types::f32, format::bfyx }); // Indexes
int64_t axis = 3;
int64_t batch_dim = -1;
@@ -106,7 +106,7 @@ TEST(check_hash_value, gather_basic) {
const auto params_hash = prog_node.get_kernel_impl_params()->hash();
ASSERT_EQ(primitive_hash, 93320679543770233UL);
ASSERT_EQ(params_hash, 18126277300376770566UL);
ASSERT_EQ(params_hash, 10172614448743302593UL);
}
TEST(check_hash_value, gemm_basic) {
@@ -132,13 +132,13 @@ TEST(check_hash_value, gemm_basic) {
const auto params_hash = prog_node.get_kernel_impl_params()->hash();
ASSERT_EQ(primitive_hash, 8009877756431655269UL);
ASSERT_EQ(params_hash, 2966249915421110547UL);
ASSERT_EQ(params_hash, 5997261677234251325UL);
}
TEST(check_hash_value, permute_basic) {
auto& engine = get_test_engine();
auto input = engine.allocate_memory({ data_types::f32, format::bfyx,{ 2, 2, 3, 2 } });
auto input = engine.allocate_memory({ { 2, 2, 3, 2 }, data_types::f32, format::bfyx });
auto key_prim_id = "permute";
topology topology(
@@ -155,7 +155,7 @@ TEST(check_hash_value, permute_basic) {
const auto params_hash = prog_node.get_kernel_impl_params()->hash();
ASSERT_EQ(primitive_hash, 4658575237077439700UL);
ASSERT_EQ(params_hash, 4319508487906266226UL);
ASSERT_EQ(params_hash, 17614582924449795856UL);
}
TEST(check_hash_value, reorder_basic) {
@@ -166,8 +166,8 @@ TEST(check_hash_value, reorder_basic) {
const int32_t y_in = 4;
const int32_t x_in = 8 * 2;
auto input = engine.allocate_memory({ data_types::f32, format::bfyx, { b_in,f_in,x_in,y_in } });
layout output_layout(data_types::f32, format::b_fs_yx_fsv16, { b_in,f_in,x_in,y_in });
auto input = engine.allocate_memory({ { b_in,f_in,y_in,x_in }, data_types::f32, format::bfyx });
layout output_layout({ b_in,f_in,y_in,x_in }, data_types::f32, format::b_fs_yx_fsv16);
auto key_prim_id = "reorder";
topology topology(
@@ -184,13 +184,13 @@ TEST(check_hash_value, reorder_basic) {
const auto params_hash = prog_node.get_kernel_impl_params()->hash();
ASSERT_EQ(primitive_hash, 16293979194373117693UL);
ASSERT_EQ(params_hash, 1719378641386629286UL);
ASSERT_EQ(params_hash, 3697547102322623168UL);
}
TEST(check_hash_value, reshape_basic) {
auto& engine = get_test_engine();
auto input = engine.allocate_memory({ data_types::f32, format::bfyx, { 1, 1, 2, 2 } });
auto input = engine.allocate_memory({ { 1, 1, 2, 2 }, data_types::f32, format::bfyx });
auto key_prim_id = "reshape";
topology topology;
@@ -210,15 +210,15 @@ TEST(check_hash_value, reshape_basic) {
const auto params_hash = prog_node.get_kernel_impl_params()->hash();
ASSERT_EQ(primitive_hash, 1534749073560581535UL);
ASSERT_EQ(params_hash, 1686780870642992006UL);
ASSERT_EQ(params_hash, 11227242497540193830UL);
}
TEST(check_hash_value, conv_basic) {
auto& engine = get_test_engine();
auto input = engine.allocate_memory({ data_types::f32, format::bfzyx, { 1, 1, 4, 4, 4 } });
auto weights = engine.allocate_memory({ data_types::f32, format::bfzyx, { 1, 1, 2, 2, 2 } });
auto biases = engine.allocate_memory({ data_types::f32, format::bfyx, { 1, 1, 1, 1, 1 } });
auto input = engine.allocate_memory({ { 1, 1, 4, 4, 4 }, data_types::f32, format::bfzyx });
auto weights = engine.allocate_memory({ { 1, 1, 2, 2, 2 }, data_types::f32, format::bfzyx });
auto biases = engine.allocate_memory({ { 1, 1, 1, 1 }, data_types::f32, format::bfyx });
auto key_prim_id = "convolution";
topology topology(
@@ -237,17 +237,17 @@ TEST(check_hash_value, conv_basic) {
const auto params_hash = prog_node.get_kernel_impl_params()->hash();
ASSERT_EQ(primitive_hash, 14591385718963138714UL);
ASSERT_EQ(params_hash, 6876197578014654797UL);
ASSERT_EQ(params_hash, 14648650631759496785UL);
}
TEST(check_hash_value, quantize_basic) {
auto& engine = get_test_engine();
auto input = engine.allocate_memory({data_types::f32, format::bfyx, {1, 16, 2, 2}});
auto input_low = engine.allocate_memory({ data_types::f32,format::bfyx,{ 1, 16, 1, 1 } });
auto input_high = engine.allocate_memory({ data_types::f32,format::bfyx,{ 1, 16, 1, 1 } });
auto output_low = engine.allocate_memory({ data_types::f32,format::bfyx,{ 1, 1, 1, 1 } });
auto output_high = engine.allocate_memory({ data_types::f32,format::bfyx,{ 1, 1, 1, 1 } });
auto input = engine.allocate_memory({ {1, 16, 2, 2}, data_types::f32, format::bfyx });
auto input_low = engine.allocate_memory({ { 1, 16, 1, 1 }, data_types::f32,format::bfyx });
auto input_high = engine.allocate_memory({ { 1, 16, 1, 1 }, data_types::f32,format::bfyx });
auto output_low = engine.allocate_memory({ { 1, 1, 1, 1 }, data_types::f32,format::bfyx });
auto output_high = engine.allocate_memory({ { 1, 1, 1, 1 }, data_types::f32,format::bfyx });
auto key_prim_id = "quantize";
topology topology;
@@ -270,5 +270,5 @@ TEST(check_hash_value, quantize_basic) {
const auto params_hash = prog_node.get_kernel_impl_params()->hash();
ASSERT_EQ(primitive_hash, 4135863035456568493UL);
ASSERT_EQ(params_hash, 13898649554943348250UL);
ASSERT_EQ(params_hash, 11586366692713421904UL);
}
src/plugins/intel_gpu/tests/test_cases/lru_caches_gpu_test.cpp
+2 -1
View File
@@ -134,7 +134,8 @@ struct ImplHasher {
TEST(lru_cache, collisions) {
auto l1 = layout{{8, 3, 131, 384}, data_types::f32, format::bfyx};
auto l2 = layout{{8, 5, 1, 384}, data_types::f32, format::bfyx};
auto l2 = layout{{6, 136, 80, 321}, data_types::f32, format::bfyx};
auto input1_prim = std::make_shared<input_layout>("input1", l1);
auto input2_prim = std::make_shared<input_layout>("input2", l2);
auto shape_of1_prim = std::make_shared<shape_of>("shape_of1", input_info("input1"), 4, data_types::i64);
src/plugins/intel_gpu/tests/test_cases/shape_of_gpu_test.cpp
+2 -2
View File
@@ -149,7 +149,7 @@ TEST(shape_of_gpu, dynamic) {
std::vector<int32_t> expected_results = {1, 2, 3, 4};
for (size_t i = 0; i < expected_results.size(); ++i) {
ASSERT_TRUE(are_equal(expected_results[i], output_ptr[i]));
ASSERT_EQ(expected_results[i], output_ptr[i]);
}
}
@@ -164,7 +164,7 @@ TEST(shape_of_gpu, dynamic) {
std::vector<int32_t> expected_results = {4, 3, 2, 1};
for (size_t i = 0; i < expected_results.size(); ++i) {
ASSERT_TRUE(are_equal(expected_results[i], output_ptr[i]));
ASSERT_EQ(expected_results[i], output_ptr[i]);
}
}
}
src/plugins/intel_gpu/tests/test_cases/tensor_test.cpp
+55
View File
@@ -93,3 +93,58 @@ TEST(tensor_api, linear_offsets) {
test_tensor_offset({ 2, 5, 4, 3 }, { 1, 3, 1, 2 }, cldnn::format::fs_b_yx_fsv32, 675);
test_tensor_offset({ 2, 37, 4, 3 }, { 1, 35, 3, 2 }, cldnn::format::fs_b_yx_fsv32, 1507);
}
TEST(tensor_api, transform) {
cldnn::tensor t4d{{1, 2, 3, 4}};
ASSERT_EQ(t4d.transform(cldnn::format::bfvuwzyx, 1), cldnn::tensor({1, 2, 3, 4, 1, 1, 1, 1}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::bfuwzyx, 1), cldnn::tensor({1, 2, 3, 4, 1, 1, 1}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::bfwzyx, 1), cldnn::tensor({1, 2, 3, 4, 1, 1}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::bfzyx, 1), cldnn::tensor({1, 2, 3, 4, 1}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::bfyx, 1), cldnn::tensor({1, 2, 3, 4}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::bs_fs_fsv8_bsv8, 1), cldnn::tensor({1, 2*3*4}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::os_i_osv8__ai8, 1), cldnn::tensor({1, 2*3*4}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::os_i_osv16__ai8, 1), cldnn::tensor({1, 2*3*4}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::bs_f_bsv16, 1), cldnn::tensor({1, 2*3*4}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::b_fs_zyx_fsv16, 1), cldnn::tensor({1, 2, 3, 4, 1}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::bs_fs_zyx_bsv16_fsv16, 1), cldnn::tensor({1, 2, 3, 4, 1}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::bs_fs_zyx_bsv16_fsv32, 1), cldnn::tensor({1, 2, 3, 4, 1}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::b_fs_zyx_fsv32, 1), cldnn::tensor({1, 2, 3, 4, 1}, 1));
ASSERT_EQ(t4d.transform(cldnn::format::bs_fs_yx_bsv16_fsv16, 1), cldnn::tensor({1, 2, 3, 4}, 1));
cldnn::tensor t5d{{1, 2, 3, 4, 5}};
ASSERT_EQ(t5d.transform(cldnn::format::bfvuwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 1, 1, 1}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::bfuwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 1, 1}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::bfwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 1}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::bfzyx, 1), cldnn::tensor({1, 2, 3, 4, 5}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::bfyx, 1), cldnn::tensor({1, 2, 3, 4*5}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::bs_fs_fsv8_bsv8, 1), cldnn::tensor({1, 2*3*4*5}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::os_i_osv8__ai8, 1), cldnn::tensor({1, 2*3*4*5}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::os_i_osv16__ai8, 1), cldnn::tensor({1, 2*3*4*5}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::bs_f_bsv16, 1), cldnn::tensor({1, 2*3*4*5}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::b_fs_zyx_fsv16, 1), cldnn::tensor({1, 2, 3, 4, 5}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::bs_fs_zyx_bsv16_fsv16, 1), cldnn::tensor({1, 2, 3, 4, 5}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::bs_fs_zyx_bsv16_fsv32, 1), cldnn::tensor({1, 2, 3, 4, 5}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::b_fs_zyx_fsv32, 1), cldnn::tensor({1, 2, 3, 4, 5}, 1));
ASSERT_EQ(t5d.transform(cldnn::format::bs_fs_yx_bsv16_fsv16, 1), cldnn::tensor({1, 2, 3, 4*5}, 1));
cldnn::tensor t6d{{1, 2, 3, 4, 5, 6}};
ASSERT_EQ(t6d.transform(cldnn::format::bfvuwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 6, 1, 1}, 1));
ASSERT_EQ(t6d.transform(cldnn::format::bfuwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 6, 1}, 1));
ASSERT_EQ(t6d.transform(cldnn::format::bfwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 6}, 1));
ASSERT_EQ(t6d.transform(cldnn::format::bfzyx, 1), cldnn::tensor({1, 2, 3, 4, 5*6}, 1));
ASSERT_EQ(t6d.transform(cldnn::format::bfyx, 1), cldnn::tensor({1, 2, 3, 4*5*6}, 1));
cldnn::tensor t7d{{1, 2, 3, 4, 5, 6, 7}};
ASSERT_EQ(t7d.transform(cldnn::format::bfvuwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 6, 7, 1}, 1));
ASSERT_EQ(t7d.transform(cldnn::format::bfuwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 6, 7}, 1));
ASSERT_EQ(t7d.transform(cldnn::format::bfwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 6*7}, 1));
ASSERT_EQ(t7d.transform(cldnn::format::bfzyx, 1), cldnn::tensor({1, 2, 3, 4, 5*6*7}, 1));
ASSERT_EQ(t7d.transform(cldnn::format::bfyx, 1), cldnn::tensor({1, 2, 3, 4*5*6*7}, 1));
cldnn::tensor t8d{{1, 2, 3, 4, 5, 6, 7, 8}};
ASSERT_EQ(t8d.transform(cldnn::format::bfvuwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 6, 7, 8}, 1));
ASSERT_EQ(t8d.transform(cldnn::format::bfuwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 6, 7*8}, 1));
ASSERT_EQ(t8d.transform(cldnn::format::bfwzyx, 1), cldnn::tensor({1, 2, 3, 4, 5, 6*7*8}, 1));
ASSERT_EQ(t8d.transform(cldnn::format::bfzyx, 1), cldnn::tensor({1, 2, 3, 4, 5*6*7*8}, 1));
ASSERT_EQ(t8d.transform(cldnn::format::bfyx, 1), cldnn::tensor({1, 2, 3, 4*5*6*7*8}, 1));
}
src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/eltwise.cpp
+23 -11
View File
@@ -24,6 +24,10 @@ std::vector<std::vector<ov::Shape>> inShapes = {
{{1, 4, 4, 1}},
{{1, 4, 3, 2, 1, 3}},
{{1, 3, 1, 1, 1, 3}, {1, 3, 1, 1, 1, 1}},
{{1, 3, 2, 2, 2, 3, 2, 3}, {1, 3, 1, 1, 1, 1, 1, 1}},
{{1, 3, 2, 2, 2, 3, 2, 3}, {3}},
{{1, 3, 2, 2, 2, 3, 2, 3}, {1, 3, 2, 2, 2, 3, 2, 3}},
{{1, 3, 2, 2, 2, 3, 2}, {1, 3, 2, 2, 2, 3, 2}},
};
std::vector<ov::test::ElementType> netPrecisions = {
@@ -42,6 +46,11 @@ std::vector<CommonTestUtils::OpType> opTypes = {
CommonTestUtils::OpType::VECTOR,
};
std::vector<ngraph::helpers::EltwiseTypes> smoke_eltwiseOpTypes = {
ngraph::helpers::EltwiseTypes::ADD,
ngraph::helpers::EltwiseTypes::MULTIPLY,
};
std::vector<ngraph::helpers::EltwiseTypes> eltwiseOpTypes = {
ngraph::helpers::EltwiseTypes::ADD,
ngraph::helpers::EltwiseTypes::MULTIPLY,
@@ -58,6 +67,20 @@ ov::AnyMap additional_config = {};
INSTANTIATE_TEST_SUITE_P(
smoke_CompareWithRefs,
EltwiseLayerTest,
::testing::Combine(::testing::ValuesIn(ov::test::static_shapes_to_test_representation(inShapes)),
::testing::ValuesIn(smoke_eltwiseOpTypes),
::testing::ValuesIn(secondaryInputTypes),
::testing::ValuesIn(opTypes),
::testing::ValuesIn(netPrecisions),
::testing::Values(ov::element::undefined),
::testing::Values(ov::element::undefined),
::testing::Values(CommonTestUtils::DEVICE_GPU),
::testing::Values(additional_config)),
EltwiseLayerTest::getTestCaseName);
INSTANTIATE_TEST_SUITE_P(
CompareWithRefs,
EltwiseLayerTest,
::testing::Combine(::testing::ValuesIn(ov::test::static_shapes_to_test_representation(inShapes)),
::testing::ValuesIn(eltwiseOpTypes),
::testing::ValuesIn(secondaryInputTypes),
@@ -69,15 +92,4 @@ INSTANTIATE_TEST_SUITE_P(
::testing::Values(additional_config)),
EltwiseLayerTest::getTestCaseName);
std::vector<std::vector<std::vector<size_t>>> inShapesSingleThread = {
{{1, 2, 3, 4}},
{{2, 2, 2, 2}},
{{2, 1, 2, 1, 2, 2}}
};
std::vector<ngraph::helpers::EltwiseTypes> eltwiseOpTypesSingleThread = {
ngraph::helpers::EltwiseTypes::ADD,
ngraph::helpers::EltwiseTypes::POWER,
};
} // namespace
src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/reduce_ops.cpp
+18
View File
@@ -28,6 +28,7 @@ const std::vector<std::vector<size_t>> inputShapes = {
};
const std::vector<std::vector<size_t>> inputShapesOneAxis = {
std::vector<size_t>{2, 3, 4, 5, 4, 3, 2, 3},
std::vector<size_t>{10, 20, 30, 40},
std::vector<size_t>{3, 5, 7, 9},
std::vector<size_t>{10},
@@ -115,6 +116,23 @@ INSTANTIATE_TEST_SUITE_P(smoke_Reduce_Precisions,
testing::Values(CommonTestUtils::DEVICE_GPU)),
ReduceOpsLayerTest::getTestCaseName);
INSTANTIATE_TEST_SUITE_P(smoke_Reduce_Ranks,
ReduceOpsLayerTest,
testing::Combine(testing::Values(std::vector<int>{2, 3, 4},
std::vector<int>{-1, -2}),
testing::Values(opTypes[1]),
testing::ValuesIn(keepDims),
testing::Values(ngraph::helpers::ReductionType::Sum),
testing::Values(InferenceEngine::Precision::FP32,
InferenceEngine::Precision::I32),
testing::Values(InferenceEngine::Precision::UNSPECIFIED),
testing::Values(InferenceEngine::Precision::UNSPECIFIED),
testing::Values(InferenceEngine::Layout::ANY),
testing::Values(std::vector<size_t>{2, 3, 4, 5, 4, 3, 2, 3},
std::vector<size_t>{2, 3, 4, 5, 4, 3, 2}),
testing::Values(CommonTestUtils::DEVICE_GPU)),
ReduceOpsLayerTest::getTestCaseName);
INSTANTIATE_TEST_SUITE_P(smoke_Reduce_InputShapes,
ReduceOpsLayerTest,
testing::Combine(testing::Values(std::vector<int>{0}),
src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/reshape.cpp
+33
View File
@@ -57,4 +57,37 @@ INSTANTIATE_TEST_SUITE_P(smoke_ReshapeCheckNegative, ReshapeLayerTest,
::testing::Values(CommonTestUtils::DEVICE_GPU),
::testing::Values(std::map<std::string, std::string>({}))),
ReshapeLayerTest::getTestCaseName);
static std::vector<reshapeParams> generate_tests() {
std::vector<reshapeParams> res;
using TargetPattern = std::vector<int64_t>;
using InputShape = std::vector<size_t>;
std::vector<std::pair<InputShape, TargetPattern>> params = {
{{1, 2, 3, 4}, {-1}},
{{1, 2, 3, 4}, {1, 6, 4}},
{{1, 3, 4}, {1, 3, 2, 2}},
{{2, 2, 3, 4, 5}, {4, 3, 20}},
{{2, 2, 3, 4, 5}, {12, 2, 2, 5}},
{{2, 2, 3, 4, 5}, {2, 2, 3, 1, 4, 5}},
{{2, 2, 3, 1, 4, 5}, {2, 2, 3, 4, 5}},
{{2, 2, 3, 1, 4, 5}, {2, 1, -1, 5}},
{{2, 2, 3, 1, 4, 5}, {2, 1, -1}},
{{2, 2, 3, 1, 4, 5}, {2, -1}},
{{2, 2, 3, 1, 6, 8}, {2, 2, 3, 1, 3, 2, 4, 2}},
{{2, 2, 3, 1, 3, 2, 4, 2}, {2, 2, 3, 1, 6, 8}},
};
for (auto& p : params) {
reshapeParams test_case = std::make_tuple(false, InferenceEngine::Precision::FP16,
InferenceEngine::Precision::UNSPECIFIED, InferenceEngine::Precision::UNSPECIFIED,
InferenceEngine::Layout::ANY, InferenceEngine::Layout::ANY,
p.first, p.second,
CommonTestUtils::DEVICE_GPU, std::map<std::string, std::string>({}));
res.push_back(test_case);
}
return res;
}
INSTANTIATE_TEST_SUITE_P(smoke_ReshapeTests, ReshapeLayerTest, ::testing::ValuesIn(generate_tests()), ReshapeLayerTest::getTestCaseName);
} // namespace
src/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/transpose.cpp
+23
View File
@@ -109,4 +109,27 @@ INSTANTIATE_TEST_SUITE_P(smoke_Transpose_6D,
testing::Values(CommonTestUtils::DEVICE_GPU)),
TransposeLayerTest::getTestCaseName);
/**
* 8D permute tests
*/
const std::vector<std::vector<size_t>> inputShapes8D = {
std::vector<size_t>{1, 2, 3, 4, 5, 6, 7, 8},
};
const std::vector<std::vector<size_t>> inputOrder8D = {
std::vector<size_t>{1, 2, 4, 3, 6, 7, 5, 0},
};
INSTANTIATE_TEST_SUITE_P(smoke_Transpose_8D,
TransposeLayerTest,
testing::Combine(testing::ValuesIn(inputOrder8D),
testing::ValuesIn(netPrecisions),
testing::Values(InferenceEngine::Precision::UNSPECIFIED),
testing::Values(InferenceEngine::Precision::UNSPECIFIED),
testing::Values(InferenceEngine::Layout::ANY),
testing::Values(InferenceEngine::Layout::ANY),
testing::ValuesIn(inputShapes8D),
testing::Values(CommonTestUtils::DEVICE_GPU)),
TransposeLayerTest::getTestCaseName);
} // namespace