[GPU] New data and weights formats support (#6616)

inference-engine/thirdparty/clDNN/api/cldnn/runtime/layout.hpp

@@ -14,6 +14,7 @@
 #include <limits>
 #include <string>
 #include <functional>
+#include <set>
 
 namespace cldnn {
 /// @addtogroup cpp_api C++ API
@@ -121,6 +122,8 @@ struct data_type_traits {
                 return alignof(data_type_to_type<data_types::bin>::type);
             case data_types::i8:
                 return alignof(data_type_to_type<data_types::i8>::type);
+            case data_types::u8:
+                return alignof(data_type_to_type<data_types::u8>::type);
             case data_types::i32:
                 return alignof(data_type_to_type<data_types::i32>::type);
             case data_types::i64:
@@ -395,20 +398,33 @@ struct layout {
     size_t get_linear_size() const {
         auto sizes = get_buffer_size().sizes();
 
-        for (const auto& block : this->format.block_sizes()) {
-            auto block_axis = block.first;
-            auto block_size = block.second;
+        std::set<size_t> processed_dims;
+        const auto& blocks = format.block_sizes();
+        for (size_t i = 0; i < blocks.size(); i++) {
+            if (processed_dims.count(blocks[i].first))
+                continue;
+
+            auto block_axis = blocks[i].first;
+            auto block_size = blocks[i].second;
+
+            for (size_t j = i + 1; j < blocks.size(); j++) {
+                if (blocks[j].first != block_axis)
+                    continue;
+
+                block_size *= blocks[j].second;
+            }
 
             sizes[block_axis] = align_to(sizes[block_axis], block_size);
+            processed_dims.insert(block_axis);
         }
 
-        if (this->format == cldnn::format::os_is_yx_isa8_osv8_isv4 && !(is_aligned_to(sizes[0], 8)) && !(is_aligned_to(sizes[1], 32))) {
+        if (this->format == cldnn::format::os_is_yx_isa8_osv8_isv4 && (!(is_aligned_to(sizes[0], 8)) || !(is_aligned_to(sizes[1], 32)))) {
             sizes[0] = align_to(sizes[0], 8);
             sizes[1] = align_to(sizes[1], 32);
-        } else if (this->format == cldnn::format::os_is_yx_isa8_osv16_isv4 && !(is_aligned_to(sizes[0], 16)) && !(is_aligned_to(sizes[1], 32))) {
+        } else if (this->format == cldnn::format::os_is_yx_isa8_osv16_isv4 && (!(is_aligned_to(sizes[0], 16)) || !(is_aligned_to(sizes[1], 32)))) {
             sizes[0] = align_to(sizes[0], 16);
             sizes[1] = align_to(sizes[1], 32);
-        } else if (this->format == cldnn::format::os_is_yx_isa8_osv8_isv4_swizzled_by_4 && !(is_aligned_to(sizes[0], 32)) && !(is_aligned_to(sizes[1], 32))) {
+        } else if (this->format == cldnn::format::os_is_yx_isa8_osv8_isv4_swizzled_by_4 && (!(is_aligned_to(sizes[0], 32)) || !(is_aligned_to(sizes[1], 32)))) {
             sizes[0] = align_to(sizes[0], 32);
             sizes[1] = align_to(sizes[1], 32);
         } else if (this->format == cldnn::format::is_o32_yx_isv32_swizzled_by_4 && (!is_aligned_to(sizes[1], 32) || !(is_aligned_to(sizes[0], 32)))) {

inference-engine/thirdparty/clDNN/api/cldnn/runtime/tensor.hpp

@@ -90,6 +90,10 @@ struct format {
         b_fs_zyx_fsv32,                         ///< format used for blocked int8 3d convolution
         bs_fs_zyx_bsv16_fsv16,                  ///< format used for 3D blocked convolution (batch and features blocked by 16)
         bs_fs_yx_bsv16_fsv16,                   ///< format used for 2D blocked convolution (batch and features blocked by 16)
+        bs_fs_yx_bsv4_fsv4,                     ///< format used for 2D blocked convolution (batch and features blocked by 4)
+        bs_fs_yx_bsv4_fsv2,                     ///< format used for 2D blocked convolution (batch blocked by 4, features blocked by 2)
+        bs_fs_yx_bsv32_fsv32,                   ///< format used for big batches (batch and features blocked by 32)
+        bs_fs_yx_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
         b_fs_yx_fsv4,                           ///< format for input for IMAD convolutions
         bs_xs_xsv8_bsv8,                        ///< format used only for fully connected weights: bs - batch slice,
@@ -110,7 +114,9 @@ struct format {
         oizyx,                                        ///< the most common format for 3D convolution
         iozyx,                                        ///< 3D weights format for deconvolutions
         iyxo,
-        os_iyx_osv16,                                 ///< format used only for convolution weights:
+        os_iyx_osv16,                                 ///< format used only for convolution weights
+        o_is_yx_isv16,                                ///< format used only for convolution weights
+        os_yxi_osv16,                                 ///< format used only for convolution weights
         os_is_yx_osv16_isv16,                         ///< format used for convolution i8 weights
         os_is_zyx_osv32_isv16,
         os_is_zyx_osv64_isv16,
@@ -150,11 +156,17 @@ struct format {
         os_is_yx_isa8_osv8_isv4_swizzled_by_4,        ///< format for weights for MMAD convolution
         os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4,   ///< format for weights for MMAD fsv32 convolution
         os_is_zyx_osa4_isa8_osv8_isv4_swizzled_by_4,  ///< format for weights for MMAD fsv32 convolution
+        os_is_yx_osa4_isa8_osv8_isv4,                 ///< format for weights for MMAD fsv32 convolution
+        g_os_is_yx_osa4_isa8_osv8_isv4,               ///< format for weights for MMAD fsv32 convolution
+        g_os_is_yx_osa4_isa8_osv8_isv2,               ///< format for weights for MMAD fsv32 convolution
+        os_is_yx_osa4_isa8_osv8_isv2,                 ///< format for weights for MMAD fsv32 convolution
         is_o_yx_isv32,                                ///< format for weights for 1x1 MMAD convolutions
         is_o32_yx_isv32_swizzled_by_4,                ///< format for weights for 1x1 MMAD convolutions
         os_is_y_x8_osv8_isv4,                         ///< format for weights for 1x1 MMAD convolutions
         os_is_y_x8_osv8_isv4_swizzled_by_4,           ///< format for weights for 1x1 MMAD convolutions
         os_is_yx_osv16_isv4,                          ///< format for weights for IMAD convolutions
+        os_is_yx_osv8_isv4,                           ///< format used for convolution i8 weights
+        os_is_yx_osv8_isv2,                           ///< format used for convolution i8 weights
         os_is_zyx_osv16_isv16,                        ///< format for weights for IMAD convolutions
         os_is_yx_osv32_isv4_swizzled_by_2,            ///< format for weights for IMAD convolutions
         os_is_yx_osv32_isv4,                          ///< format for weights for IMAD convolutions
@@ -236,6 +248,10 @@ struct format {
                 { b_fs_zyx_fsv16,        { 1, 1, 3, 0, 0, "bfzyx",  "bfxyz",  {{1, 16}}}},
                 { bs_fs_zyx_bsv16_fsv16, { 1, 1, 3, 0, 0, "bfzyx",  "bfxyz",  {{0, 16 }, {1, 16}}}},
                 { bs_fs_yx_bsv16_fsv16,  { 1, 1, 2, 0, 0, "bfyx",   "bfxy?",  {{0, 16 }, {1, 16}}}},
+                { bs_fs_yx_bsv4_fsv4,    { 1, 1, 2, 0, 0, "bfyx",   "bfxy?",  {{0, 4 }, {1, 4}}}},
+                { bs_fs_yx_bsv4_fsv2,    { 1, 1, 2, 0, 0, "bfyx",   "bfxy?",  {{0, 4 }, {1, 2}}}},
+                { bs_fs_yx_bsv32_fsv32,  { 1, 1, 2, 0, 0, "bfyx",   "bfxy?",  {{0, 32 }, {1, 32}}}},
+                { bs_fs_yx_bsv32_fsv16,  { 1, 1, 2, 0, 0, "bfyx",   "bfxy?",  {{0, 32 }, {1, 16}}}},
                 { nv12,                  { 1, 1, 2, 0, 0, "bfyx",   "bfxy?",  {}}},
                 { image_2d_rgba,         { 1, 1, 2, 0, 0, "bfyx",   "bfxy?",  {}}},
 
@@ -246,6 +262,8 @@ struct format {
                 { oizyx,                                       { 1, 1, 3, 0, 0, "oizyx",  "oixyz",      {}}},
                 { iozyx,                                       { 1, 1, 3, 0, 0, "iozyx",  "oixyz",      {}}},
                 { os_is_yx_isv16_osv16,                        { 1, 1, 2, 0, 0, "oiyx",   "oixy",       {{1, 16}, {0, 16}}}},
+                { o_is_yx_isv16,                               { 1, 1, 2, 0, 0, "oiyx",   "oixy?",      {{1, 16}}}},
+                { os_yxi_osv16,                                { 1, 1, 2, 0, 0, "oyxi",   "oixy?",      {{0, 16}}}},
                 { os_iyx_osv16,                                { 1, 1, 2, 0, 0, "oiyx",   "oixy?",      {{0, 16}}}},
                 { os_iyx_osv32,                                { 1, 1, 2, 0, 0, "oiyx",   "oixy?",      {{0, 32}}}},
                 { os_iyx_osv64,                                { 1, 1, 2, 0, 0, "oiyx",   "oixy?",      {{0, 64}}}},
@@ -260,6 +278,8 @@ struct format {
                 { os_is_yx_isa8_osv8_isv4,                     { 1, 1, 2, 0, 0, "oiyx",   "oixy?",      {}}},
                 { os_is_yx_isa8_osv16_isv4,                    { 1, 1, 2, 0, 0, "oiyx",   "oixy?",      {}}},
                 { os_is_yx_isa8_osv8_isv4_swizzled_by_4,       { 1, 1, 2, 0, 0, "oiyx",   "oixy?",      {}}},
+                { os_is_yx_osa4_isa8_osv8_isv4,                { 1, 1, 2, 0, 0, "oiyx",   "oixy?",      {{0, 32}, {1, 32}}}},
+                { os_is_yx_osa4_isa8_osv8_isv2,                { 1, 1, 2, 0, 0, "oiyx",   "oixy?",      {{0, 32}, {1, 16}}}},
                 { os_is_zyx_isa8_osv8_isv4,                    { 1, 1, 3, 0, 0, "oizyx",  "oixyz",      {{1, 8}, {0, 8}, {1, 4}}}},
                 { os_is_zyx_isa8_osv16_isv4,                   { 1, 1, 3, 0, 0, "oizyx",  "oixyz",      {{1, 8}, {0, 16}, {1, 4}}}},
                 { os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4,  { 1, 1, 2, 0, 0, "oiyx",   "oixy?",      {{0, 32}, {1, 32}}}},
@@ -269,6 +289,8 @@ struct format {
                 { os_is_y_x8_osv8_isv4,                        { 1, 1, 2, 0, 0, "oyxi",   "oixy?",      {}}},
                 { os_is_y_x8_osv8_isv4_swizzled_by_4,          { 1, 1, 2, 0, 0, "oyxi",   "oixy?",      {}}},
                 { os_is_yx_osv16_isv4,                         { 1, 1, 2, 0, 0, "oixy",   "oixy?",      {{0, 16}, {1, 4}}}},
+                { os_is_yx_osv8_isv4,                          { 1, 1, 2, 0, 0, "oiyx",   "oixy",       {{1, 4}, {0, 8}}}},
+                { os_is_yx_osv8_isv2,                          { 1, 1, 2, 0, 0, "oiyx",   "oixy",       {{1, 2}, {0, 8}}}},
                 { os_is_zyx_osv16_isv16,                       { 1, 1, 3, 0, 0, "oizyx",  "oixyz",      {{0, 16}, {1, 16}}}},
                 { os_is_yx_osv32_isv4_swizzled_by_2,           { 1, 1, 2, 0, 0, "oixy",   "oixy?",      {{0, 32}, {1, 4}}}},
                 { os_is_yx_osv32_isv4,                         { 1, 1, 2, 0, 0, "oixy",   "oixy?",      {{0, 32}, {1, 4}}}},
@@ -313,6 +335,8 @@ struct format {
                 { g_os_zyx_is_osv32_isv4,                      { 1, 1, 3, 0, 1, "gozyxi", "oixyz???g",  {{0, 32}, {1, 4}}}},
                 { g_os_zyx_is_osv32_isv16,                     { 1, 1, 3, 0, 1, "gozyxi", "oixyz???g",  {{0, 32}, {1, 16}}}},
                 { g_os_zyx_is_osv32_isv32,                     { 1, 1, 3, 0, 1, "gozyxi", "oixyz???g",  {{0, 32}, {1, 32}}}},
+                { g_os_is_yx_osa4_isa8_osv8_isv4,              { 1, 1, 2, 0, 1, "goiyx",  "oixy????g",  {{0, 32}, {1, 32}}}},
+                { g_os_is_yx_osa4_isa8_osv8_isv2,              { 1, 1, 2, 0, 1, "goiyx",  "oixy????g",  {{0, 32}, {1, 16}}}},
                 { gs_oi_yxs_gsv4_yxsv4,                        { 1, 1, 2, 0, 1, "goiyx",  "oixy????g",  {{8, 4}}}},
                 { gs_oi_yxs_gsv16_yxsv4,                       { 1, 1, 2, 0, 1, "goiyx",  "oixy????g",  {{8, 16}}}},
                 { gs_oi_yxs_gsv32_yxsv4,                       { 1, 1, 2, 0, 1, "goiyx",  "oixy????g",  {{8, 32}}}},

inference-engine/thirdparty/clDNN/kernel_selector/common/tensor_type.cpp

@@ -28,6 +28,10 @@ DataTensor::DataChannelArray DataTensor::dataChannelArray {{
     { DataLayout::b_fs_zyx_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_bsv4_fsv2,    {  0,  1, -1, -1,  2,  3 } },
+    { DataLayout::bs_fs_yx_bsv32_fsv32,  {  0,  1, -1, -1,  2,  3 } },
+    { 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 } },
@@ -54,6 +58,8 @@ WeightsTensor::WeightsChannelArray WeightsTensor::weightsChannelArray {{
     { WeightsLayout::os_iyx_osv32__ai32,                          {  0,  1, -1,   2,   3, -1, -1, -1 } },
     { WeightsLayout::os_iyx_osv64,                                {  0,  1, -1,   2,   3, -1, -1, -1 } },
     { WeightsLayout::os_iyx_osv16_rotate_180,                     {  0,  1, -1,   2,   3, -1, -1, -1 } },
+    { WeightsLayout::o_is_yx_isv16,                               {  0,  1, -1,   2,   3, -1, -1, -1 } },
+    { WeightsLayout::os_yxi_osv16,                                {  1,  2, -1,   0,   3, -1, -1, -1 } },
     { WeightsLayout::os_i_osv8__ai8,                              { -1, -1, -1,   0,   1, -1, -1, -1 } },
     { WeightsLayout::os_i_osv16__ai8,                             { -1, -1, -1,   0,   1, -1, -1, -1 } },
     { WeightsLayout::os_i_osv16,                                  { -1, -1, -1,   0,   1, -1, -1, -1 } },
@@ -76,6 +82,10 @@ WeightsTensor::WeightsChannelArray WeightsTensor::weightsChannelArray {{
     { WeightsLayout::os_is_yx_isa8_osv8_isv4,                     {  0,  1, -1,   2,   3, -1, -1, -1 } },
     { WeightsLayout::os_is_yx_isa8_osv16_isv4,                    {  0,  1, -1,   2,   3, -1, -1, -1 } },
     { WeightsLayout::os_is_yx_isa8_osv8_isv4_swizzled_by_4,       {  0,  1, -1,   2,   3, -1, -1, -1 } },
+    { WeightsLayout::os_is_yx_osa4_isa8_osv8_isv4,                {  0,  1, -1,   2,   3, -1, -1, -1 } },
+    { WeightsLayout::g_os_is_yx_osa4_isa8_osv8_isv4,              {  0,  1, -1,   2,   3, -1, -1,  4 } },
+    { WeightsLayout::g_os_is_yx_osa4_isa8_osv8_isv2,              {  0,  1, -1,   2,   3, -1, -1,  4 } },
+    { WeightsLayout::os_is_yx_osa4_isa8_osv8_isv2,                {  0,  1, -1,   2,   3, -1, -1, -1 } },
     { WeightsLayout::os_is_zyx_isa8_osv8_isv4,                    {  0,  1,  2,   3,   4, -1, -1, -1 } },
     { WeightsLayout::os_is_zyx_isa8_osv16_isv4,                   {  0,  1,  2,   3,   4, -1, -1, -1 } },
     { WeightsLayout::os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4,  {  0,  1, -1,   2,   3, -1, -1, -1 } },
@@ -84,6 +94,8 @@ WeightsTensor::WeightsChannelArray WeightsTensor::weightsChannelArray {{
     { WeightsLayout::is_o32_yx_isv32_swizzled_by_4,               {  1,  2, -1,   0,   3, -1, -1, -1 } },
     { WeightsLayout::os_is_y_x8_osv8_isv4,                        {  0,  1, -1,   2,   3, -1, -1, -1 } },
     { WeightsLayout::os_is_y_x8_osv8_isv4_swizzled_by_4,          {  0,  1, -1,   2,   3, -1, -1, -1 } },
+    { WeightsLayout::os_is_yx_osv8_isv4,                          {  0,  1, -1,   2,   3, -1, -1, -1 } },
+    { WeightsLayout::os_is_yx_osv8_isv2,                          {  0,  1, -1,   2,   3, -1, -1, -1 } },
     { WeightsLayout::os_is_yx_osv16_isv4,                         {  0,  1, -1,   2,   3, -1, -1, -1 } },
     { WeightsLayout::os_is_yx_osv32_isv4_swizzled_by_2,           {  0,  1, -1,   2,   3, -1, -1, -1 } },
     { WeightsLayout::os_is_yx_osv32_isv4,                         {  0,  1, -1,   2,   3, -1, -1, -1 } },
@@ -180,6 +192,26 @@ NDims DataTensor::GetSimpleDims(const std::vector<size_t>& d, DataLayout l) {
             newDims[3] = RoundUp(newDims[3], 16);
             newDims[4] = RoundUp(newDims[4], 16);
             break;
+        case bs_fs_yx_bsv4_fsv4:
+            assert(newDims.size() == 4);
+            newDims[2] = RoundUp(newDims[2], 4);
+            newDims[3] = RoundUp(newDims[3], 4);
+            break;
+        case bs_fs_yx_bsv4_fsv2:
+            assert(newDims.size() == 4);
+            newDims[2] = RoundUp(newDims[2], 2);
+            newDims[3] = RoundUp(newDims[3], 4);
+            break;
+        case bs_fs_yx_bsv32_fsv32:
+            assert(newDims.size() == 4);
+            newDims[2] = RoundUp(newDims[2], 32);
+            newDims[3] = RoundUp(newDims[3], 32);
+            break;
+        case bs_fs_yx_bsv32_fsv16:
+            assert(newDims.size() == 4);
+            newDims[2] = RoundUp(newDims[2], 16);
+            newDims[3] = RoundUp(newDims[3], 32);
+            break;
         default:
             break;
     }
@@ -391,7 +423,12 @@ NDims WeightsTensor::GetSimpleDims(const std::vector<size_t>& d, WeightsLayout l
 
     // TODO: It's not the right pitches. it's here in order to calculate physical size
     switch (l) {
+        case o_is_yx_isv16:
+            assert(newDims.size() == 4);
+            newDims[2] = RoundUp(newDims[2], 16);
+            break;
         case os_iyx_osv16:
+        case os_yxi_osv16:
         case os_iyx_osv16_rotate_180:
             assert(newDims.size() == 4);
             newDims[3] = RoundUp(newDims[3], 16);
@@ -481,6 +518,16 @@ NDims WeightsTensor::GetSimpleDims(const std::vector<size_t>& d, WeightsLayout l
             newDims[0] = RoundUp(newDims[0], 32);
             newDims[3] = RoundUp(newDims[3], 32);
             break;
+        case os_is_yx_osa4_isa8_osv8_isv4:
+        case g_os_is_yx_osa4_isa8_osv8_isv4:
+            newDims[3] = RoundUp(newDims[3], 32);
+            newDims[2] = RoundUp(newDims[2], 32);
+            break;
+        case os_is_yx_osa4_isa8_osv8_isv2:
+        case g_os_is_yx_osa4_isa8_osv8_isv2:
+            newDims[3] = RoundUp(newDims[3], 32);
+            newDims[2] = RoundUp(newDims[2], 16);
+            break;
         case os_is_y_x8_osv8_isv4:
         case os_is_y_x8_osv8_isv4_swizzled_by_4:
             assert(newDims.size() == 4);
@@ -543,6 +590,11 @@ NDims WeightsTensor::GetSimpleDims(const std::vector<size_t>& d, WeightsLayout l
             assert(newDims.size() == 5);
             newDims[3] = RoundUp(newDims[0], 16);
             break;
+        case os_is_yx_osv8_isv4:
+            assert(newDims.size() == 4);
+            newDims[2] = RoundUp(newDims[2], 4);
+            newDims[3] = RoundUp(newDims[3], 8);
+            break;
         case os_i_yxs_osv4_yxsv4:
             newDims[3] = RoundUp(newDims[3], 4);
             break;

inference-engine/thirdparty/clDNN/kernel_selector/common/tensor_type.h

@@ -38,6 +38,10 @@ enum DataLayout {
     b_fs_zyx_fsv32,         // 4D+batch
     bs_fs_yx_bsv16_fsv16,   // batch, feature, 2D spatial. Blocks of 16 batch and channels
     bs_fs_zyx_bsv16_fsv16,  // batch, feature, 3D spatial. Blocks of 16 batch and channels
+    bs_fs_yx_bsv4_fsv4,     // batch, feature, 2D spatial. Blocks of 4 batch and 4 channels
+    bs_fs_yx_bsv4_fsv2,     // batch, feature, 2D spatial. Blocks of 4 batch and 2 channels
+    bs_fs_yx_bsv32_fsv32,   // batch, feature, 2D spatial. Blocks of 32 batch and 32 channels
+    bs_fs_yx_bsv32_fsv16,   // batch, feature, 2D spatial. Blocks of 32 batch and 16 channels
     bs_f_bsv8__af8,         // for optimized FC
     bs_f_bsv16__af8,        // for optimized FC
     winograd_2x3_s1_data,   // winograd convolution input, F(2,3) -- filter 3x3 with stride 1
@@ -62,6 +66,8 @@ enum WeightsLayout {
     oyxi,
     iyxo,
     yxio,
+    o_is_yx_isv16,
+    os_yxi_osv16,
     os_iyx_osv16,
     os_iyx_osv32,
     os_iyx_osv32__ai32,
@@ -98,6 +104,10 @@ enum WeightsLayout {
     os_is_zyx_isa8_osv8_isv4,                // for MMAD convolution
     os_is_yx_isa8_osv16_isv4,                // for fully connected MMAD
     os_is_zyx_isa8_osv16_isv4,               // for fully connected MMAD
+    os_is_yx_osa4_isa8_osv8_isv4,            // for MMAD convolution swizzled from ofm 0..7 to 0,4,8,12,16,20,24,28,
+    g_os_is_yx_osa4_isa8_osv8_isv4,          // for MMAD convolution swizzled from ofm 0..7 to 0,4,8,12,16,20,24,28,
+    g_os_is_yx_osa4_isa8_osv8_isv2,          // for MMAD convolution swizzled from ofm 0..7 to 0,4,8,12,16,20,24,28,
+    os_is_yx_osa4_isa8_osv8_isv2,            // for MMAD convolution swizzled from ofm 0..7 to 0,4,8,12,16,20,24,28,
     os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4,  // for MMAD convolution swizzled from ofm 0..7 to 0,4,8,12,16,20,24,28,
                                                  // 1,5...
     os_is_zyx_osa4_isa8_osv8_isv4_swizzled_by_4,  // for MMAD convolution swizzled from ofm 0..7 to 0,4,8,12,16,20,24,28,
@@ -110,6 +120,8 @@ enum WeightsLayout {
     os_is_y_x8_osv8_isv4_swizzled_by_4,  // for MMAD 1x1 convolutions swizzled from ofm 0..7 to 0,4,8,12,16,20,24,28,
                                          // 1,5...
     os_is_yx_osv16_isv4,                 // swizzled weights for convolution using IMAD
+    os_is_yx_osv8_isv4,                      // weights for int8 blocked conv
+    os_is_yx_osv8_isv2,                      // weights for int8 blocked conv
     os_is_yx_osv32_isv4_swizzled_by_2,   //  weights for bfyx -> b_fs_yx_fsv32 convolution using IMAD with swizzeled ofm (0, 2, 4..), (1, 3, 5...)
     os_is_yx_osv32_isv4,                 //  weights for bfyx -> b_fs_yx_fsv{32,16} convolution using IMAD
     os_is_zyx_osv32_isv4,                //  weights for bfzyx -> b_fs_zyx_fsv16 convolution using IMAD

inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/activation/activation_kernel_base.cpp

@@ -16,9 +16,12 @@ ActivationKernelBase::DispatchData ActivationKernelBase::SetDefault(const activa
     if (out.GetLayout() == DataLayout::yxfb) {
         dispatchData.gws = {out.Feature().v * out.Batch().v, out.X().v, out.Y().v};
         dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, arg.engineInfo);
-    } else if (out.GetLayout() == DataLayout::b_fs_yx_fsv16) {
+    } else if (out.GetLayout() == DataLayout::b_fs_yx_fsv16 || out.GetLayout() == DataLayout::b_fs_yx_fsv32) {
         dispatchData.gws = {Align(out.Feature().v, 16) * out.Batch().v, out.X().v, out.Y().v};
         dispatchData.lws = {16, 1, 1};
+    } else if (out.GetLayout() == DataLayout::bs_fs_yx_bsv32_fsv16 || out.GetLayout() == DataLayout::bs_fs_yx_bsv32_fsv32) {
+        dispatchData.gws = {out.X().v * out.Y().v, Align(out.Feature().v, 16), Align(out.Batch().v, 16)};
+        dispatchData.lws = {1, 16, 16};
     } else {
         dispatchData.gws = {out.X().v, out.Y().v * out.Z().v, out.Feature().v * out.Batch().v};
         dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, arg.engineInfo);

inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/activation/activation_kernel_opt.cpp

@@ -10,6 +10,7 @@ namespace kernel_selector {
 
 ParamsKey ActivationKernelOpt::GetSupportedKey() const {
     ParamsKey k;
+    k.EnableInputDataType(Datatype::UINT8);
     k.EnableInputDataType(Datatype::INT8);
     k.EnableInputDataType(Datatype::INT32);
     k.EnableInputDataType(Datatype::F16);

inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/concatenation/concatenation_kernel_simple_ref.cpp

@@ -40,6 +40,8 @@ ParamsKey ConcatenationKernel_simple_Ref::GetSupportedKey() const {
     k.EnableOutputLayout(DataLayout::bs_fs_zyx_bsv16_fsv16);
     k.EnableInputLayout(DataLayout::bs_fs_yx_bsv16_fsv16);
     k.EnableOutputLayout(DataLayout::bs_fs_yx_bsv16_fsv16);
+    k.EnableInputLayout(DataLayout::bs_fs_yx_bsv32_fsv32);
+    k.EnableOutputLayout(DataLayout::bs_fs_yx_bsv32_fsv32);
     k.EnableTensorOffset();
     k.EnableTensorPitches();
     k.EnableBatching();
@@ -66,7 +68,8 @@ bool ConcatenationKernel_simple_Ref::Validate(const Params& p, const optional_pa
     for (const auto& lt : params.inputs) {
         auto cur_layout = lt.GetLayout();
         if ((cur_layout == DataLayout::bfzyx || cur_layout == DataLayout::b_fs_zyx_fsv16 || cur_layout == DataLayout::bs_fs_zyx_bsv16_fsv16) &&
-            (same_layout == DataLayout::bfzyx || same_layout == DataLayout::b_fs_zyx_fsv16 || same_layout == DataLayout::bs_fs_zyx_bsv16_fsv16)) {
+            (same_layout == DataLayout::bfzyx || same_layout == DataLayout::b_fs_zyx_fsv16 || same_layout == DataLayout::bs_fs_zyx_bsv16_fsv16
+            || same_layout == DataLayout::bs_fs_yx_bsv32_fsv32)) {
             continue;
         } else if (cur_layout != same_layout) {
             return false;

inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/eltwise/eltwise_kernel_base.cpp

@@ -39,6 +39,25 @@ static uint32_t GetNumberOfInputs(EltwiseMode m) {
     }
 }
 
+std::vector<size_t> GetLimitedOptimalLocalWorkGroupSizes(std::vector<size_t> gws, const EngineInfo& info, std::vector<size_t> limited_size_lws) {
+    const size_t lws_max = info.maxWorkGroupSize;
+    const size_t optimal_lws_values[] = {256, 227, 224, 192, 160, 128, 96, 64, 32, 16, 8, 7, 6, 5, 4, 2, 1};
+    size_t total_lws = 1;
+    std::vector<size_t> lws;
+    for (size_t i = 0; i < gws.size(); ++i) {
+        auto rest_lws = lws_max / total_lws;
+        size_t lws_idx = 0;
+        while (rest_lws < optimal_lws_values[lws_idx] || optimal_lws_values[lws_idx] > limited_size_lws[i]) lws_idx++;
+
+        while (gws[i] % optimal_lws_values[lws_idx]) lws_idx++;
+
+        lws.push_back(optimal_lws_values[lws_idx]);
+        total_lws *= optimal_lws_values[lws_idx];
+    }
+
+    return lws;
+}
+
 ParamsKey eltwise_params::GetParamsKey() const {
     ParamsKey k = base_params::GetParamsKey();
 
@@ -569,6 +588,7 @@ EltwiseKernelBase::DispatchData EltwiseKernelBase::SetDefault(const eltwise_para
     const size_t optimal_lws_values[] = {256, 224, 192, 160, 128, 96, 64, 32, 16};
     if ((params.output.GetLayout() == DataLayout::b_fs_yx_fsv16 ||
          params.output.GetLayout() == DataLayout::b_fs_zyx_fsv16 ||
+         params.output.GetLayout() == DataLayout::bs_fs_yx_bsv32_fsv16 ||
          params.output.GetLayout() == DataLayout::bs_fs_yx_bsv16_fsv16) &&
         params.output.Feature().v % 16 == 0 && dispatchData.gws[1] % 16 == 0) {
         dispatchData.lws[0] = 1;
@@ -584,17 +604,21 @@ EltwiseKernelBase::DispatchData EltwiseKernelBase::SetDefault(const eltwise_para
         dispatchData.lws[0] = 1;
         dispatchData.lws[1] = 1;
         dispatchData.lws[2] = 32;
-    } else if (params.output.GetLayout() == DataLayout::b_fs_yx_fsv32 && params.output.Feature().v % 32 == 0) {
+    } else if ((params.output.GetLayout() == DataLayout::b_fs_yx_fsv32 ||
+                params.output.GetLayout() == DataLayout::bs_fs_yx_bsv32_fsv32)) {
         if (params.layoutBased || params.int8_quantization || params.broadcast) {
-            dispatchData.lws[0] = 1;
-            dispatchData.lws[1] = 32;
-            dispatchData.lws[2] = 1;
+            auto bs_fsv32_local = GetLimitedOptimalLocalWorkGroupSizes({dispatchData.gws[1], dispatchData.gws[2], dispatchData.gws[0]},
+                                                                        params.engineInfo, {32, 32, 1024});
+            dispatchData.lws[0] = bs_fsv32_local[2];
+            dispatchData.lws[1] = bs_fsv32_local[0];
+            dispatchData.lws[2] = bs_fsv32_local[1];
         } else if (dispatchData.gws[0] == params.output.LogicalSize()) {
             dispatchData.lws = local;
         } else {
-            dispatchData.lws[0] = 1;
-            dispatchData.lws[1] = 1;
-            dispatchData.lws[2] = 32;
+            auto bs_fsv32_local = GetOptimalLocalWorkGroupSizes({dispatchData.gws[2], dispatchData.gws[0], dispatchData.gws[1]}, params.engineInfo);
+            dispatchData.lws[0] = bs_fsv32_local[1];
+            dispatchData.lws[1] = bs_fsv32_local[2];
+            dispatchData.lws[2] = bs_fsv32_local[0];
         }
     } else {
         dispatchData.lws[0] = local[0];

inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/lrn/lrn_kernel_across_channel_multiple_features_fsv16.cpp

@@ -16,7 +16,13 @@ ParamsKey LRNKernelAcrossChannelMultipleFeaturesFSV16::GetSupportedKey() const {
     k.EnableOutputDataType(Datatype::INT8);
     k.EnableOutputDataType(Datatype::UINT8);
     k.EnableInputLayout(DataLayout::b_fs_yx_fsv16);
+    k.EnableInputLayout(DataLayout::b_fs_yx_fsv32);
+    k.EnableInputLayout(DataLayout::bs_fs_yx_bsv32_fsv16);
+    k.EnableInputLayout(DataLayout::bs_fs_yx_bsv32_fsv32);
     k.EnableOutputLayout(DataLayout::b_fs_yx_fsv16);
+    k.EnableOutputLayout(DataLayout::b_fs_yx_fsv32);
+    k.EnableOutputLayout(DataLayout::bs_fs_yx_bsv32_fsv16);
+    k.EnableOutputLayout(DataLayout::bs_fs_yx_bsv32_fsv32);
     k.EnableLRNMode(LRNMode::ACROSS_CHANNEL);
     k.EnableLRNKernelDividerMode(KernelDividerMode::FIXED);
     k.EnableTensorOffset();

inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/pooling/pooling_kernel_gpu_int8_ref.cpp

@@ -22,6 +22,7 @@ ParamsKey PoolingKernelGPUInt8Ref::GetSupportedKey() const {
     k.EnableInputLayout(DataLayout::b_fs_yx_fsv32);
     k.EnableInputLayout(DataLayout::b_fs_zyx_fsv32);
     k.EnableInputLayout(DataLayout::b_fs_yx_fsv16);
+    k.EnableInputLayout(DataLayout::bs_fs_yx_bsv32_fsv32);
     k.EnableOutputLayout(DataLayout::bfyx);
     k.EnableOutputLayout(DataLayout::bfzyx);
     k.EnableOutputLayout(DataLayout::yxfb);
@@ -30,6 +31,7 @@ ParamsKey PoolingKernelGPUInt8Ref::GetSupportedKey() const {
     k.EnableOutputLayout(DataLayout::b_fs_yx_fsv32);
     k.EnableOutputLayout(DataLayout::b_fs_zyx_fsv32);
     k.EnableOutputLayout(DataLayout::b_fs_yx_fsv16);
+    k.EnableOutputLayout(DataLayout::bs_fs_yx_bsv32_fsv32);
     k.EnableTensorOffset();
     k.EnableTensorPitches();
     k.EnableBatching();

inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/quantize/quantize_kernel_scale_shift_opt.cpp

@@ -8,6 +8,7 @@
 #include <string>
 
 static const size_t sub_group_size = 32;
+static const size_t feature_size = 32;
 
 namespace kernel_selector {
 ParamsKey QuantizeKernelScaleShift::GetSupportedKey() const {
@@ -43,6 +44,14 @@ CommonDispatchData QuantizeKernelScaleShift::SetDefault(const quantize_params& p
         dispatchData.lws[0] = 1;
         dispatchData.lws[1] = sub_group_size;
         dispatchData.lws[2] = 1;
+    } else if (output.GetLayout() == DataLayout::bs_fs_yx_bsv32_fsv32) {
+        dispatchData.gws[0] = output.Y().v * output.X().v;
+        dispatchData.gws[1] = Align(output.Feature().v, feature_size);
+        dispatchData.gws[2] = Align(output.Batch().v, feature_size);
+
+        dispatchData.lws[0] = 1;
+        dispatchData.lws[1] = feature_size;
+        dispatchData.lws[2] = params.engineInfo.maxWorkGroupSize / feature_size;
     } else {
         dispatchData.gws = GetTensorFriendlyWorkGroups(output);
         dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo);
@@ -54,7 +63,8 @@ CommonDispatchData QuantizeKernelScaleShift::SetDefault(const quantize_params& p
 JitConstants QuantizeKernelScaleShift::GetJitConstants(const quantize_params& params, const CommonDispatchData& dispatchData) const {
     JitConstants jit = Parent::GetJitConstants(params, dispatchData);
 
-    if (params.output.GetLayout() == DataLayout::b_fs_yx_fsv16) {
+    if (params.output.GetLayout() == DataLayout::b_fs_yx_fsv16 ||
+        params.output.GetLayout() == DataLayout::bs_fs_yx_bsv32_fsv32) {
         jit.AddConstant(MakeJitConstant("GWS_BATCH", 2));
         jit.AddConstant(MakeJitConstant("GWS_FEATURE", 1));
         jit.AddConstant(MakeJitConstant("GWS_YX", 0));

inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/reorder/reorder_kernel_base.cpp

@@ -177,7 +177,11 @@ ReorderKernelBase::DispatchData ReorderKernelBase::SetDefault(const reorder_para
         }
     }
 
-    if (params.output.GetLayout() == DataLayout::bs_fs_yx_bsv16_fsv16 && params.inputs[0].Feature().v % 16 == 0) {
+    if ((params.output.GetLayout() == DataLayout::bs_fs_yx_bsv16_fsv16 ||
+         params.output.GetLayout() == DataLayout::bs_fs_yx_bsv32_fsv32 ||
+         params.output.GetLayout() == DataLayout::b_fs_yx_fsv16 ||
+         params.output.GetLayout() == DataLayout::bs_fs_yx_bsv32_fsv16) &&
+        params.inputs[0].Feature().v % 16 == 0) {
         dispatchData.lws[0] = 1;
         dispatchData.lws[1] = 16;
         dispatchData.lws[2] = 1;

inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/resample/resample_kernel_opt.cpp

@@ -18,6 +18,15 @@ size_t ResampleKernelOpt::GetOptimalBlockSize(const resample_params& params) con
     return 1;
 }
 
+static size_t GetOptimalDivisor(const size_t input_size, size_t max_val = 16) {
+    for (size_t s = max_val; s > 0; --s) {
+        if (input_size % s == 0) {
+            return s;
+        }
+    }
+    return 1;
+}
+
 ParamsKey ResampleKernelOpt::GetSupportedKey() const {
     ParamsKey k;
     k.EnableInputDataType(Datatype::F16);
@@ -29,8 +38,14 @@ ParamsKey ResampleKernelOpt::GetSupportedKey() const {
     k.EnableOutputDataType(Datatype::UINT8);
     k.EnableOutputDataType(Datatype::INT8);
     k.EnableInputLayout(DataLayout::b_fs_yx_fsv16);
+    k.EnableInputLayout(DataLayout::b_fs_yx_fsv32);
+    k.EnableInputLayout(DataLayout::bs_fs_yx_bsv32_fsv16);
+    k.EnableInputLayout(DataLayout::bs_fs_yx_bsv32_fsv32);
     k.EnableInputLayout(DataLayout::fs_b_yx_fsv32);
     k.EnableOutputLayout(DataLayout::b_fs_yx_fsv16);
+    k.EnableOutputLayout(DataLayout::b_fs_yx_fsv32);
+    k.EnableOutputLayout(DataLayout::bs_fs_yx_bsv32_fsv16);
+    k.EnableOutputLayout(DataLayout::bs_fs_yx_bsv32_fsv32);
     k.EnableOutputLayout(DataLayout::fs_b_yx_fsv32);
     k.EnableDifferentTypes();
     k.EnableTensorOffset();
@@ -56,6 +71,11 @@ ResampleKernelBase::DispatchData ResampleKernelOpt::SetDefault(const kernel_sele
 
         dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, arg.engineInfo);
     } else {
+        auto opt_x_block_size = GetOptimalBlockSize(arg);
+        if (out.X().v > 32 && opt_x_block_size == 1) {
+            opt_x_block_size = GetOptimalDivisor(out.X().v, 32);
+        }
+
         dispatchData.gws[0] = CeilDiv(out.X().v, GetOptimalBlockSize(arg)) * out.Y().v;
         dispatchData.gws[1] = Align(out.Feature().v, sub_group_size);
         dispatchData.gws[2] = arg.output.Batch().v;
@@ -63,6 +83,11 @@ ResampleKernelBase::DispatchData ResampleKernelOpt::SetDefault(const kernel_sele
         dispatchData.lws[0] = 1;
         dispatchData.lws[1] = sub_group_size;
         dispatchData.lws[2] = 1;
+
+        if (arg.output.GetLayout() == DataLayout::bs_fs_yx_bsv32_fsv16
+            || arg.output.GetLayout() == DataLayout::bs_fs_yx_bsv32_fsv32) {
+            dispatchData.lws[2] = GetOptimalDivisor(dispatchData.gws[2]);
+        }
     }
 
     return dispatchData;
@@ -91,7 +116,11 @@ bool ResampleKernelOpt::Validate(const Params& p, const optional_params& o) cons
         params.resampleType != ResampleType::BILINEAR_INTERP)
         return false;
 
-    if (input.GetLayout() != DataLayout::fs_b_yx_fsv32 && input.GetLayout() != DataLayout::b_fs_yx_fsv16)
+    if (input.GetLayout() != DataLayout::fs_b_yx_fsv32 &&
+        input.GetLayout() != DataLayout::b_fs_yx_fsv16 &&
+        input.GetLayout() != DataLayout::b_fs_yx_fsv32 &&
+        input.GetLayout() != DataLayout::bs_fs_yx_bsv32_fsv16 &&
+        input.GetLayout() != DataLayout::bs_fs_yx_bsv32_fsv32)
         return false;
 
     return true;
@@ -100,9 +129,15 @@ bool ResampleKernelOpt::Validate(const Params& p, const optional_params& o) cons
 JitConstants ResampleKernelOpt::GetJitConstants(const resample_params &params) const {
     auto jit = Parent::GetJitConstants(params);
 
-    jit.AddConstant(MakeJitConstant("OUTPUT_X_BLOCK_SIZE", GetOptimalBlockSize(params)));
+    auto opt_x_block_size = GetOptimalBlockSize(params);
+    if (params.output.X().v > 32 && opt_x_block_size == 1) {
+        opt_x_block_size = GetOptimalDivisor(params.output.X().v, 32);
+    }
+
+    jit.AddConstant(MakeJitConstant("OUTPUT_X_BLOCK_SIZE", opt_x_block_size));
+    jit.AddConstant(MakeJitConstant("X_BLOCKS", CeilDiv(params.output.X().v, opt_x_block_size)));
     jit.AddConstant(MakeJitConstant("SUB_GROUP_SIZE", sub_group_size));
-    jit.AddConstant(MakeJitConstant("X_BLOCKS", CeilDiv(params.output.X().v, GetOptimalBlockSize(params))));
+
     size_t vec_size = 0;
     if (params.inputs[0].GetLayout() == DataLayout::fs_b_yx_fsv32) {
         vec_size = 2;

inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/resample/resample_kernel_ref.cpp

@@ -50,6 +50,7 @@ static size_t packing_factor(const resample_params& params) {
     auto get_layout_packing_factor = [](const DataLayout& layout) -> size_t {
         switch (layout) {
         case DataLayout::b_fs_yx_fsv16:
+        case DataLayout::bs_fs_yx_bsv32_fsv16:
             return 16;
         case DataLayout::b_fs_yx_fsv4:
             return 4;

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/activation_ref.cl

@@ -15,11 +15,11 @@
 KERNEL(activation)(
     __global INPUT0_TYPE* input,
     __global OUTPUT_TYPE* output
-#ifdef PARAMETERIZED
-    , __global ADDITIONAL_PARAMS_TYPE* params
-#endif
 #if HAS_FUSED_OPS_DECLS
     , FUSED_OPS_DECLS
+#endif
+#ifdef PARAMETERIZED
+    , __global ADDITIONAL_PARAMS_TYPE* params
 #endif
     )
 {

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_simple_ref.cl

@@ -20,6 +20,8 @@ inline uint FUNC(get_input_index)(uint b, uint f, uint w, uint z, uint y, uint x
     return GET_DATA_BS_FS_ZYX_BSV16_FSV16_INDEX(INPUT0, b, f, z, y, x);
 #elif INPUT0_LAYOUT_BS_FS_YX_BSV16_FSV16
     return GET_DATA_BS_FS_YX_BSV16_FSV16_INDEX(INPUT0, b, f, y, x);
+#elif INPUT0_LAYOUT_BS_FS_YX_BSV32_FSV32
+    return GET_DATA_BS_FS_YX_BSV32_FSV32_INDEX(INPUT0, b, f, y, x);
 #else
 #error concatenation_gpu_simple_ref.cl: input format - not supported
 #endif
@@ -40,6 +42,8 @@ inline uint FUNC(get_output_index)(uint b, uint f, uint w, uint z, uint y, uint
     return GET_DATA_BS_FS_ZYX_BSV16_FSV16_INDEX(OUTPUT, b, f, z, y, x);
 #elif OUTPUT_LAYOUT_BS_FS_YX_BSV16_FSV16
     return GET_DATA_BS_FS_YX_BSV16_FSV16_INDEX(OUTPUT, b, f, y, x);
+#elif OUTPUT_LAYOUT_BS_FS_YX_BSV32_FSV32
+    return GET_DATA_BS_FS_YX_BSV32_FSV32_INDEX(OUTPUT, b, f, y, x);
 #else
 #error concatenation_gpu_simple_ref.cl: output format - not supported
 #endif

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_tiled.cl

@@ -4,6 +4,7 @@
 
 #include "include/common.cl"
 #include "include/data_types.cl"
+#include "include/fetch_data.cl"
 
 // JIT Parameters:
 // SIMD         - sub-group size/simd width, one of {8, 16};

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/fetch_data.cl

@@ -402,7 +402,7 @@ inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index_safe)(uint b, uint f, uint z, uint
                                                   uint f_pad_before, uint f_pad_after,
                                                   uint z_pad_before, uint z_pad_after,
                                                   uint y_pad_before, uint y_pad_after,
-                                                  uint x_pad_before, uint x_pad_after, uint alignmentF, uint alignmentB) {
+                                                  uint x_pad_before, uint x_pad_after, uint alignmentB, uint alignmentF) {
     const uint b_mod = b % b_size;
     const uint f_mod = f_pad_before + (f % f_size);
     const uint fs = f_mod / alignmentF;
@@ -426,26 +426,26 @@ inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index_safe)(uint b, uint f, uint z, uint
     return output_offset;
 }
 
-inline uint FUNC(get_bs_fs_zyx_bsv16_fsv16_index)(uint b, uint f,  uint z, uint y, uint x,
-                                                  uint x_size, uint y_size, uint z_size, uint f_size,
-                                                  uint f_pad_before, uint f_pad_after,
-                                                  uint z_pad_before, uint z_pad_after,
-                                                  uint y_pad_before, uint y_pad_after,
-                                                  uint x_pad_before, uint x_pad_after) {
-    const uint alignment = 16;
+inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index)(uint b, uint f,  uint z, uint y, uint x,
+                                              uint x_size, uint y_size, uint z_size, uint f_size,
+                                              uint f_pad_before, uint f_pad_after,
+                                              uint z_pad_before, uint z_pad_after,
+                                              uint y_pad_before, uint y_pad_after,
+                                              uint x_pad_before, uint x_pad_after,
+                                              uint b_alignment, uint f_alignment) {
     const uint feature = f + f_pad_before;
-    const uint fs = feature / alignment;
-    const uint fsv = feature % alignment;
-    const uint bs = b / alignment;
-    const uint bsv = b % alignment;
+    const uint fs = feature / f_alignment;
+    const uint fsv = feature % f_alignment;
+    const uint bs = b / b_alignment;
+    const uint bsv = b % b_alignment;
 
-    const uint bsv_pitch = alignment;
-    const uint x_pitch = bsv_pitch * alignment;
+    const uint bsv_pitch = f_alignment;
+    const uint x_pitch = bsv_pitch * b_alignment;
     const uint y_pitch = x_pitch * (x_pad_before + x_size + x_pad_after);
     const uint z_pitch = y_pitch * (y_pad_before + y_size + y_pad_after);
     const uint fs_pitch = z_pitch * (z_pad_before + z_size + z_pad_after);
     const uint total_f_size = f_pad_before + f_size + f_pad_after;
-    const uint bs_pitch = fs_pitch * ((total_f_size + alignment - 1) / alignment);
+    const uint bs_pitch = fs_pitch * ((total_f_size + f_alignment - 1) / f_alignment);
 
     const uint output_offset = bs * bs_pitch +
                                fs * fs_pitch +
@@ -458,8 +458,8 @@ inline uint FUNC(get_bs_fs_zyx_bsv16_fsv16_index)(uint b, uint f,  uint z, uint
     return output_offset;
 }
 
-#define GET_DATA_BS_FS_YX_BSV16_FSV16_INDEX(prefix, b, f, y, x) \
-    FUNC_CALL(get_bs_fs_zyx_bsv16_fsv16_index)(                     \
+#define GET_DATA_BS_FS_YX_BSV16_FSV16_INDEX(prefix, b, f, y, x)     \
+    FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index)(                         \
         b, f, 0, y, x,                                              \
         CAT(prefix, _SIZE_X),                                       \
         CAT(prefix, _SIZE_Y),                                       \
@@ -472,10 +472,74 @@ inline uint FUNC(get_bs_fs_zyx_bsv16_fsv16_index)(uint b, uint f,  uint z, uint
         CAT(prefix, _PAD_BEFORE_SIZE_Y),                            \
         CAT(prefix, _PAD_AFTER_SIZE_Y),                             \
         CAT(prefix, _PAD_BEFORE_SIZE_X),                            \
-        CAT(prefix, _PAD_AFTER_SIZE_X))
+        CAT(prefix, _PAD_AFTER_SIZE_X), 16, 16)
+
+#define GET_DATA_BS_FS_YX_BSV32_FSV32_INDEX(prefix, b, f, y, x)     \
+    FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index)(                         \
+        b, f, 0, y, x,                                              \
+        CAT(prefix, _SIZE_X),                                       \
+        CAT(prefix, _SIZE_Y),                                       \
+        CAT(prefix, _SIZE_Z),                                       \
+        CAT(prefix, _FEATURE_NUM),                                  \
+        CAT(prefix, _PAD_BEFORE_FEATURE_NUM),                       \
+        CAT(prefix, _PAD_AFTER_FEATURE_NUM),                        \
+        CAT(prefix, _PAD_BEFORE_SIZE_Z),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_Z),                             \
+        CAT(prefix, _PAD_BEFORE_SIZE_Y),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_Y),                             \
+        CAT(prefix, _PAD_BEFORE_SIZE_X),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_X), 32, 32)
+
+#define GET_DATA_BS_FS_YX_BSV4_FSV4_INDEX(prefix, b, f, y, x)       \
+    FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index)(                         \
+        b, f, 0, y, x,                                              \
+        CAT(prefix, _SIZE_X),                                       \
+        CAT(prefix, _SIZE_Y),                                       \
+        CAT(prefix, _SIZE_Z),                                       \
+        CAT(prefix, _FEATURE_NUM),                                  \
+        CAT(prefix, _PAD_BEFORE_FEATURE_NUM),                       \
+        CAT(prefix, _PAD_AFTER_FEATURE_NUM),                        \
+        CAT(prefix, _PAD_BEFORE_SIZE_Z),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_Z),                             \
+        CAT(prefix, _PAD_BEFORE_SIZE_Y),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_Y),                             \
+        CAT(prefix, _PAD_BEFORE_SIZE_X),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_X), 4, 4)
+
+#define GET_DATA_BS_FS_YX_BSV4_FSV2_INDEX(prefix, b, f, y, x)       \
+    FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index)(                         \
+        b, f, 0, y, x,                                              \
+        CAT(prefix, _SIZE_X),                                       \
+        CAT(prefix, _SIZE_Y),                                       \
+        CAT(prefix, _SIZE_Z),                                       \
+        CAT(prefix, _FEATURE_NUM),                                  \
+        CAT(prefix, _PAD_BEFORE_FEATURE_NUM),                       \
+        CAT(prefix, _PAD_AFTER_FEATURE_NUM),                        \
+        CAT(prefix, _PAD_BEFORE_SIZE_Z),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_Z),                             \
+        CAT(prefix, _PAD_BEFORE_SIZE_Y),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_Y),                             \
+        CAT(prefix, _PAD_BEFORE_SIZE_X),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_X), 4, 2)
+
+#define GET_DATA_BS_FS_YX_BSV32_FSV16_INDEX(prefix, b, f, y, x)     \
+    FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index)(                         \
+        b, f, 0, y, x,                                              \
+        CAT(prefix, _SIZE_X),                                       \
+        CAT(prefix, _SIZE_Y),                                       \
+        CAT(prefix, _SIZE_Z),                                       \
+        CAT(prefix, _FEATURE_NUM),                                  \
+        CAT(prefix, _PAD_BEFORE_FEATURE_NUM),                       \
+        CAT(prefix, _PAD_AFTER_FEATURE_NUM),                        \
+        CAT(prefix, _PAD_BEFORE_SIZE_Z),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_Z),                             \
+        CAT(prefix, _PAD_BEFORE_SIZE_Y),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_Y),                             \
+        CAT(prefix, _PAD_BEFORE_SIZE_X),                            \
+        CAT(prefix, _PAD_AFTER_SIZE_X), 32, 16)
 
 #define GET_DATA_BS_FS_ZYX_BSV16_FSV16_INDEX(prefix, b, f, z, y, x) \
-    FUNC_CALL(get_bs_fs_zyx_bsv16_fsv16_index)(                     \
+    FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index)(                         \
         b, f, z, y, x,                                              \
         CAT(prefix, _SIZE_X),                                       \
         CAT(prefix, _SIZE_Y),                                       \
@@ -488,7 +552,7 @@ inline uint FUNC(get_bs_fs_zyx_bsv16_fsv16_index)(uint b, uint f,  uint z, uint
         CAT(prefix, _PAD_BEFORE_SIZE_Y),                            \
         CAT(prefix, _PAD_AFTER_SIZE_Y),                             \
         CAT(prefix, _PAD_BEFORE_SIZE_X),                            \
-        CAT(prefix, _PAD_AFTER_SIZE_X))
+        CAT(prefix, _PAD_AFTER_SIZE_X), 16, 16)
 
 #define GET_DATA_BS_FS_YX_BSV16_FSV16_INDEX_SAFE(prefix, b, f, y, x) \
     FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index_safe)(                     \
@@ -507,6 +571,74 @@ inline uint FUNC(get_bs_fs_zyx_bsv16_fsv16_index)(uint b, uint f,  uint z, uint
         CAT(prefix, _PAD_BEFORE_SIZE_X),                             \
         CAT(prefix, _PAD_AFTER_SIZE_X), 16, 16)
 
+#define GET_DATA_BS_FS_YX_BSV32_FSV32_INDEX_SAFE(prefix, b, f, y, x) \
+    FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index_safe)(                     \
+        b, f, 0, y, x,                                               \
+        CAT(prefix, _SIZE_X),                                        \
+        CAT(prefix, _SIZE_Y),                                        \
+        CAT(prefix, _SIZE_Z),                                        \
+        CAT(prefix, _FEATURE_NUM),                                   \
+        CAT(prefix, _BATCH_NUM),                                     \
+        CAT(prefix, _PAD_BEFORE_FEATURE_NUM),                        \
+        CAT(prefix, _PAD_AFTER_FEATURE_NUM),                         \
+        CAT(prefix, _PAD_BEFORE_SIZE_Z),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_Z),                              \
+        CAT(prefix, _PAD_BEFORE_SIZE_Y),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_Y),                              \
+        CAT(prefix, _PAD_BEFORE_SIZE_X),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_X), 32, 32)
+
+#define GET_DATA_BS_FS_YX_BSV4_FSV4_INDEX_SAFE(prefix, b, f, y, x)   \
+    FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index_safe)(                     \
+        b, f, 0, y, x,                                               \
+        CAT(prefix, _SIZE_X),                                        \
+        CAT(prefix, _SIZE_Y),                                        \
+        CAT(prefix, _SIZE_Z),                                        \
+        CAT(prefix, _FEATURE_NUM),                                   \
+        CAT(prefix, _BATCH_NUM),                                     \
+        CAT(prefix, _PAD_BEFORE_FEATURE_NUM),                        \
+        CAT(prefix, _PAD_AFTER_FEATURE_NUM),                         \
+        CAT(prefix, _PAD_BEFORE_SIZE_Z),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_Z),                              \
+        CAT(prefix, _PAD_BEFORE_SIZE_Y),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_Y),                              \
+        CAT(prefix, _PAD_BEFORE_SIZE_X),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_X), 4, 4)
+
+#define GET_DATA_BS_FS_YX_BSV4_FSV2_INDEX_SAFE(prefix, b, f, y, x)   \
+    FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index_safe)(                     \
+        b, f, 0, y, x,                                               \
+        CAT(prefix, _SIZE_X),                                        \
+        CAT(prefix, _SIZE_Y),                                        \
+        CAT(prefix, _SIZE_Z),                                        \
+        CAT(prefix, _FEATURE_NUM),                                   \
+        CAT(prefix, _BATCH_NUM),                                     \
+        CAT(prefix, _PAD_BEFORE_FEATURE_NUM),                        \
+        CAT(prefix, _PAD_AFTER_FEATURE_NUM),                         \
+        CAT(prefix, _PAD_BEFORE_SIZE_Z),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_Z),                              \
+        CAT(prefix, _PAD_BEFORE_SIZE_Y),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_Y),                              \
+        CAT(prefix, _PAD_BEFORE_SIZE_X),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_X), 4, 2)
+
+#define GET_DATA_BS_FS_YX_BSV32_FSV16_INDEX_SAFE(prefix, b, f, y, x) \
+    FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index_safe)(                     \
+        b, f, 0, y, x,                                               \
+        CAT(prefix, _SIZE_X),                                        \
+        CAT(prefix, _SIZE_Y),                                        \
+        CAT(prefix, _SIZE_Z),                                        \
+        CAT(prefix, _FEATURE_NUM),                                   \
+        CAT(prefix, _BATCH_NUM),                                     \
+        CAT(prefix, _PAD_BEFORE_FEATURE_NUM),                        \
+        CAT(prefix, _PAD_AFTER_FEATURE_NUM),                         \
+        CAT(prefix, _PAD_BEFORE_SIZE_Z),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_Z),                              \
+        CAT(prefix, _PAD_BEFORE_SIZE_Y),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_Y),                              \
+        CAT(prefix, _PAD_BEFORE_SIZE_X),                             \
+        CAT(prefix, _PAD_AFTER_SIZE_X), 32, 16)
+
 #define GET_DATA_BS_FS_ZYX_BSV16_FSV16_INDEX_SAFE(prefix, b, f, z, y, x) \
     FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index_safe)(                         \
         b, f, z, y, x,                                                   \

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/fetch_weights.cl

@@ -606,6 +606,95 @@ inline uint FUNC(get_os_is_zyx_osa4_isa8_osv8_isv4_swizzled_by_4_index)(uint o,
     return idx;
 }
 
+inline uint FUNC(get_g_os_is_yx_osa4_isa8_osv8_isv4)(uint g, uint o, uint i, uint y, uint x,
+                                                     uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset)
+{
+    const uint isv_idx = i % 4;
+    const uint isa_idx = (i / 4) % 8;
+    const uint is_idx = (i / 32);
+    const uint osv_idx = o % 8;
+    const uint osa_idx = (o / 8) % 4;
+    const uint os_idx = (o / 32);
+
+    const uint ifm_32_aligned = ((size_ifm + 31)/32);
+    const uint ofm_32_aligned = ((size_ofm + 31)/32);
+
+    size_t idx = offset +
+                 isv_idx +
+                 osv_idx * 4 +
+                 isa_idx * 8 * 4 +
+                 osa_idx * 8 * 32 +
+                 x * 32 * 32 +
+                 y * size_x * 32 * 32 +
+                 is_idx * 32 * 32 * size_x * size_y +
+                 os_idx * 32 * 32 * ifm_32_aligned * size_x * size_y +
+                 g * 32 * 32 * ifm_32_aligned * ofm_32_aligned * size_x * size_y;
+
+    return idx;
+}
+
+inline uint FUNC(get_g_os_is_yx_osa4_isa8_osv8_isv2)(uint g, uint o, uint i, uint y, uint x, uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset)
+{
+    const uint isv_idx = i % 2;
+    const uint isa_idx = (i / 2) % 8;
+    const uint is_idx = (i / 16);
+    const uint osv_idx = o % 8;
+    const uint osa_idx = (o / 8) % 4;
+    const uint os_idx = (o / 32);
+
+    const uint ifm_16_aligned = ((size_ifm + 15)/16);
+    const uint ofm_32_aligned = ((size_ifm + 31)/32);
+
+    size_t idx = offset +
+                 isv_idx +
+                 osv_idx * 2 +
+                 isa_idx * 8 * 2 +
+                 osa_idx * 8 * 16 +
+                 x * 32 * 16 +
+                 y * size_x * 32 * 16 +
+                 is_idx * 32 * 16 * size_x * size_y +
+                 os_idx * 32 * 16 * ifm_16_aligned * size_x * size_y +
+                 g * 32 * 16 * ifm_16_aligned * ofm_32_aligned * size_x * size_y;
+
+    return idx;
+}
+
+#define GET_FILTER_OS_IS_YX_OSA4_ISA8_OSV8_ISV4_INDEX(prefix, o, i, y, x) \
+    FUNC_CALL(get_g_os_is_yx_osa4_isa8_osv8_isv4)(                        \
+        0, o, i, y, x,                                                    \
+        CAT(prefix, _SIZE_X),                                             \
+        CAT(prefix, _SIZE_Y),                                             \
+        CAT(prefix, _IFM_NUM),                                            \
+        CAT(prefix, _OFM_NUM),                                            \
+        CAT(prefix, _OFFSET))
+
+#define GET_FILTER_G_OS_IS_YX_OSA4_ISA8_OSV8_ISV4_INDEX(prefix, g, o, i, y, x) \
+    FUNC_CALL(get_g_os_is_yx_osa4_isa8_osv8_isv4)(                             \
+        g, o, i, y, x,                                                         \
+        CAT(prefix, _SIZE_X),                                                  \
+        CAT(prefix, _SIZE_Y),                                                  \
+        CAT(prefix, _IFM_NUM),                                                 \
+        CAT(prefix, _OFM_NUM),                                                 \
+        CAT(prefix, _OFFSET))
+
+#define GET_FILTER_OS_IS_YX_OSA4_ISA8_OSV8_ISV2_INDEX(prefix, o, i, y, x) \
+    FUNC_CALL(get_g_os_is_yx_osa4_isa8_osv8_isv2)(                        \
+        0, o, i, y, x,                                                    \
+        CAT(prefix, _SIZE_X),                                             \
+        CAT(prefix, _SIZE_Y),                                             \
+        CAT(prefix, _IFM_NUM),                                            \
+        CAT(prefix, _OFM_NUM),                                            \
+        CAT(prefix, _OFFSET))
+
+#define GET_FILTER_G_OS_IS_YX_OSA4_ISA8_OSV8_ISV2_INDEX(prefix, g, o, i, y, x) \
+    FUNC_CALL(get_g_os_is_yx_osa4_isa8_osv8_isv2)(                             \
+        g, o, i, y, x,                                                         \
+        CAT(prefix, _SIZE_X),                                                  \
+        CAT(prefix, _SIZE_Y),                                                  \
+        CAT(prefix, _IFM_NUM),                                                 \
+        CAT(prefix, _OFM_NUM),                                                 \
+        CAT(prefix, _OFFSET))
+
 #define GET_FILTER_OS_IS_YX_ISA8_OSV8_ISV4_SWIZZLED_BY_4_INDEX(prefix, o, i, y, x) \
 	FUNC_CALL(get_os_is_yx_isa8_osv8_isv4_swizzled_by_4_index)(                    \
         o, i, y, x, CAT(prefix, _SIZE_X ),                                         \
@@ -722,7 +811,7 @@ inline uint FUNC(get_os_is_y_x8_osv8_isv4_swizzled_by_4_index)(uint o, uint i, u
         CAT(prefix, _SIZE_X),                                    \
         CAT(prefix, _SIZE_Y),                                    \
         CAT(prefix, _OFM_NUM),                                   \
-        CAT(prefix, _IFM_NUM))
+        CAT(prefix, _IFM_NUM), 16, 4)
 
 inline uint FUNC(get_g_os_is_yx_osv16_isv4)(uint g, uint o, uint i, uint y, uint x,
                                           uint i_size,
@@ -730,14 +819,14 @@ inline uint FUNC(get_g_os_is_yx_osv16_isv4)(uint g, uint o, uint i, uint y, uint
                                           uint x_size,
                                           uint y_size,
                                           uint o_num,
-                                          uint i_num)
+                                          uint i_num,
+                                          uint otd,
+                                          uint tile)
 {
-    const uint otd = 16;
     uint out_depth_tile = o / otd;
     uint od             = o - out_depth_tile * otd;
     uint output_slice_size = (o_num + otd - 1) / otd;
 
-    const uint tile = 4;
     uint id_tile = i / tile;
     uint id      = i - id_tile * tile;
     uint input_slice_size = (i_num + tile - 1) / tile;
@@ -753,6 +842,26 @@ inline uint FUNC(get_g_os_is_yx_osv16_isv4)(uint g, uint o, uint i, uint y, uint
     return idx;
 }
 
+#define GET_FILTER_OS_IS_YX_OSV8_ISV2_INDEX(prefix, o, i, y, x) \
+    FUNC_CALL(get_g_os_is_yx_osv16_isv4)(                       \
+        0, o, i, y, x,                                          \
+        CAT(prefix, _IFM_PITCH),                                \
+        CAT(prefix, _OFM_PITCH),                                \
+        CAT(prefix, _SIZE_X),                                   \
+        CAT(prefix, _SIZE_Y),                                   \
+        CAT(prefix, _OFM_NUM),                                  \
+        CAT(prefix, _IFM_NUM), 8, 2)
+
+#define GET_FILTER_OS_IS_YX_OSV8_ISV4_INDEX(prefix, o, i, y, x) \
+    FUNC_CALL(get_g_os_is_yx_osv16_isv4)(                       \
+        0, o, i, y, x,                                          \
+        CAT(prefix, _IFM_PITCH),                                \
+        CAT(prefix, _OFM_PITCH),                                \
+        CAT(prefix, _SIZE_X),                                   \
+        CAT(prefix, _SIZE_Y),                                   \
+        CAT(prefix, _OFM_NUM),                                  \
+        CAT(prefix, _IFM_NUM), 8, 4)
+
 #define GET_FILTER_OS_IS_YX_OSV16_ISV4_INDEX(prefix, o, i, y, x) \
     FUNC_CALL(get_os_is_zyx_osv_isv4)(                           \
         o, i, 0, y, x,                                           \
@@ -1116,3 +1225,23 @@ inline uint FUNC(get_g_os_zyx_is_osv_isv_index)(uint g, uint o, uint i, uint z,
 #define GET_FILTER_G_OS_ZYX_IS_OSV32_ISV4_INDEX(tensor, g, o, i, z, y, x)   GET_FILTER_G_OS_ZYX_IS_OSV_ISV_INDEX(tensor, g, o, i, z, y, x, 32, 4)
 #define GET_FILTER_G_OS_ZYX_IS_OSV32_ISV16_INDEX(tensor, g, o, i, z, y, x)  GET_FILTER_G_OS_ZYX_IS_OSV_ISV_INDEX(tensor, g, o, i, z, y, x, 32, 16)
 #define GET_FILTER_G_OS_ZYX_IS_OSV32_ISV32_INDEX(tensor, g, o, i, z, y, x)  GET_FILTER_G_OS_ZYX_IS_OSV_ISV_INDEX(tensor, g, o, i, z, y, x, 32, 32)
+
+#define GET_FILTER_O_IS_YX_ISV16_INDEX(prefix, o, i, y, x, isv) \
+    CAT(prefix, _OFFSET) +                                      \
+    ((i) % (isv)) +                                             \
+    (o)*CAT(prefix, _OFM_PITCH) +                               \
+    (isv)*(                                                     \
+        (x)*CAT(prefix, _X_PITCH) +                             \
+        (y)*CAT(prefix, _Y_PITCH) +                             \
+        ((i) / (isv))*CAT(prefix, _IFM_PITCH)                   \
+    )
+
+#define GET_FILTER_OS_YXI_OSV16(prefix, o, i, y, x) \
+    CAT(prefix, _OFFSET) +                          \
+    ((o) % (16)) +                                  \
+    (16)*(                                          \
+        (x)*CAT(prefix, _X_PITCH) +                 \
+        (y)*CAT(prefix, _Y_PITCH) +                 \
+        (i)*CAT(prefix, _IFM_PITCH) +               \
+        ((o) / (16))*CAT(prefix, _OFM_PITCH)        \
+    )

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_int8_ref.cl

@@ -75,7 +75,7 @@ KERNEL(pooling_gpu_int8_ref)(
     const uint f    = bf / INPUT0_BATCH_NUM;
     const uint b    = bf % INPUT0_BATCH_NUM;
     const uint z    = 0;
-#elif OUTPUT_LAYOUT_B_FS_YX_FSV16
+#elif OUTPUT_LAYOUT_B_FS_YX_FSV16 || OUTPUT_LAYOUT_BS_FS_YX_BSV32_FSV32
     const uint x = get_global_id(1);
     const uint y = get_global_id(2);
     const uint bf = (uint)get_global_id(0);

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_to_yxfb_batched.cl

@@ -10,21 +10,8 @@
 ///////////////////////// Input Index /////////////////////////
 inline uint FUNC(get_input_index)(uint b, uint f, uint y, uint x)
 {
-#if   INPUT0_SIMPLE
-    return GET_DATA_INDEX(INPUT0, b, f, y, x);
-#elif defined INPUT0_LAYOUT_BS_F_BSV8__AF8  || \
-      defined INPUT0_LAYOUT_BS_F_BSV16__AF8
-    return GET_DATA_BS_FYX_BSV8_INDEX(INPUT0, b, f, y, x, SUB_GROUP_SIZE);
-#elif defined INPUT0_LAYOUT_B_FS_YX_FSV16
-    return GET_DATA_B_FS_YX_FSV16_INDEX(INPUT0, b, f, y, x);
-#elif defined INPUT0_LAYOUT_FS_B_YX_FSV32
-    return GET_DATA_FS_B_YX_FSV32_INDEX(INPUT0, b, f, y, x);
-#elif defined INPUT0_LAYOUT_B_FS_ZYX_FSV16
-    return GET_DATA_B_FS_ZYX_FSV16_INDEX(INPUT0, b, f, 0, y, x);
-#elif defined INPUT0_LAYOUT_BS_FS_ZYX_BSV16_FSV16
-    return GET_DATA_BS_FS_ZYX_BSV16_FSV16_INDEX(INPUT0, b, f, 0, y, x);
-#elif defined INPUT0_LAYOUT_BS_FS_YX_BSV16_FSV16
-    return GET_DATA_BS_FS_YX_BSV16_FSV16_INDEX(INPUT0, b, f, y, x);
+#if INPUT0_DIMS == 4
+    return INPUT0_GET_INDEX(b, f, y, x);
 #else
 #error reorder_data_to_yxfb_batched.cl: input format - not supported
 #endif

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights.cl

@@ -32,6 +32,8 @@ inline uint FUNC(get_input_index)(uint g, uint o, uint i, uint z, uint y, uint x
     return GET_FILTER_OS_IYX_OSV8_INDEX(INPUT0, o, i, y, x, 32);
 #elif defined INPUT0_LAYOUT_OS_IYX_OSV32__AI32
     return GET_FILTER_OS_IYX_OSV32__AI32_INDEX(OUTPUT, o, i, y, x, 32);
+#elif defined INPUT0_LAYOUT_O_IS_YX_ISV16
+    return GET_FILTER_O_IS_YX_ISV16_INDEX(INPUT0, o, i, y, x, 16);
 #elif defined INPUT0_LAYOUT_IYX_OSV64
     return GET_FILTER_OS_IYX_OSV8_INDEX(INPUT0, o, i, y, x, 64);
 #elif defined INPUT0_LAYOUT_OS_IYX_OSV16_ROTATE_180
@@ -100,6 +102,10 @@ inline uint FUNC(get_input_index)(uint g, uint o, uint i, uint z, uint y, uint x
     return GET_FILTER_OS_IS_YX_OSV16_ISV16_INDEX(INPUT0, o, i, y, x);
 #elif defined INPUT0_LAYOUT_OS_IS_ZYX_OSV16_ISV16
     return GET_FILTER_OS_IS_ZYX_OSV16_ISV16_INDEX(INPUT0, o, i, z, y, x);
+#elif defined INPUT0_LAYOUT_OS_IS_YX_OSV8_ISV2
+    return GET_FILTER_OS_IS_YX_OSV8_ISV2_INDEX(INPUT0, o, i, y, x);
+#elif defined INPUT0_LAYOUT_OS_IS_YX_OSV8_ISV4
+    return GET_FILTER_OS_IS_YX_OSV8_ISV4_INDEX(INPUT0, o, i, y, x);
 #elif defined INPUT0_LAYOUT_G_OS_IS_ZYX_OSV16_ISV16
     return GET_FILTER_G_OS_IS_ZYX_OSV16_ISV16_INDEX(INPUT0, g, o, i, z, y, x);
 #elif defined INPUT0_LAYOUT_OS_IS_ZYX_OSV32_ISV16
@@ -146,6 +152,8 @@ inline uint FUNC(get_output_index)(uint g, uint o, uint i, uint z, uint y, uint
     return GET_FILTER_OS_IYX_OSV32__AI32_INDEX(OUTPUT, o, i, y, x, 32);
 #elif defined OUTPUT_LAYOUT_OS_IYX_OSV64
     return GET_FILTER_OS_IYX_OSV8_INDEX(OUTPUT, o, i, y, x, 64);
+#elif defined OUTPUT_LAYOUT_O_IS_YX_ISV16
+    return GET_FILTER_O_IS_YX_ISV16_INDEX(OUTPUT, o, i, y, x, 16);
 #elif defined OUTPUT_LAYOUT_OS_IYX_OSV16_ROTATE_180
     return GET_FILTER_OS_IYX_OSV8_ROTATE_180_INDEX(OUTPUT, o, i, y, x, SUB_GROUP_SIZE);
 #elif defined OUTPUT_LAYOUT_I_YXS_OS_YXSV2_OSV16
@@ -172,6 +180,10 @@ inline uint FUNC(get_output_index)(uint g, uint o, uint i, uint z, uint y, uint
     return GET_FILTER_OS_IS_Y_X8_OSV8_ISV4_SWIZZLED_BY_4(OUTPUT, o, i, y, x);
 #elif defined OUTPUT_LAYOUT_OS_IS_YX_OSV16_ISV4
     return GET_FILTER_OS_IS_YX_OSV16_ISV4_INDEX(OUTPUT, o, i, y, x);
+#elif defined OUTPUT_LAYOUT_OS_IS_YX_OSV8_ISV2
+    return GET_FILTER_OS_IS_YX_OSV8_ISV2_INDEX(OUTPUT, o, i, y, x);
+#elif defined OUTPUT_LAYOUT_OS_IS_YX_OSV8_ISV4
+    return GET_FILTER_OS_IS_YX_OSV8_ISV4_INDEX(OUTPUT, o, i, y, x);
 #elif defined OUTPUT_LAYOUT_OS_IS_YX_OSV32_ISV4_SWIZZLED_BY_2
     return GET_FILTER_OS_IS_YX_OSV32_ISV4_SWIZZLED_BY_2_INDEX(OUTPUT, o, i, y, x);
 #elif defined OUTPUT_LAYOUT_OS_IS_YX_OSV32_ISV4
@@ -179,15 +191,23 @@ inline uint FUNC(get_output_index)(uint g, uint o, uint i, uint z, uint y, uint
 #elif defined OUTPUT_LAYOUT_OS_IS_ZYX_OSV32_ISV4
     return GET_FILTER_OS_IS_ZYX_OSV32_ISV4_INDEX(OUTPUT, o, i, z, y, x);
 #elif defined OUTPUT_LAYOUT_OS_IS_YX_ISA8_OSV8_ISV4_SWIZZLED_BY_4
-    return GET_FILTER_OS_IS_YX_ISA8_OSV8_ISV4_SWIZZLED_BY_4_INDEX(OUTPUT, o, i, y, x);
+    return GET_FILTER_OS_IS_YX_ISA8_OSV8_ISV4_SWIZZLED_BY_4_INDEX(OUTPUT, g, o, i, y, x);
+#elif defined OUTPUT_LAYOUT_G_OS_IS_YX_OSA4_ISA8_OSV8_ISV2
+    return GET_FILTER_G_OS_IS_YX_OSA4_ISA8_OSV8_ISV2_INDEX(OUTPUT, g, o, i, y, x);
+#elif defined OUTPUT_LAYOUT_G_OS_IS_YX_OSA4_ISA8_OSV8_ISV4
+    return GET_FILTER_G_OS_IS_YX_OSA4_ISA8_OSV8_ISV4_INDEX(OUTPUT, g, o, i, y, x);
+#elif defined OUTPUT_LAYOUT_OS_IS_YX_OSA4_ISA8_OSV8_ISV4
+    return GET_FILTER_OS_IS_YX_OSA4_ISA8_OSV8_ISV4_INDEX(OUTPUT, o, i, y, x);
+#elif defined OUTPUT_LAYOUT_OS_IS_YX_OSA4_ISA8_OSV8_ISV2
+    return GET_FILTER_OS_IS_YX_OSA4_ISA8_OSV8_ISV2_INDEX(OUTPUT, o, i, y, x);
 #elif defined OUTPUT_LAYOUT_OS_IS_YX_OSA4_ISA8_OSV8_ISV4_SWIZZLED_BY_4
     return GET_FILTER_OS_IS_YX_OSA4_ISA8_OSV8_ISV4_SWIZZLED_BY_4_INDEX(OUTPUT, o, i, y, x);
 #elif defined OUTPUT_LAYOUT_OS_IS_ZYX_OSA4_ISA8_OSV8_ISV4_SWIZZLED_BY_4
     return GET_FILTER_OS_IS_ZYX_OSA4_ISA8_OSV8_ISV4_SWIZZLED_BY_4_INDEX(OUTPUT, o, i, z, y, x);
 #elif defined OUTPUT_LAYOUT_OS_IS_YX_ISV16_OSV16
     return GET_FILTER_OS_IS_YX_ISV16_OSV16_INDEX(OUTPUT, o, i, y, x, SUB_GROUP_SIZE);
-#elif defined OUTPUT_LAYOUT_OIYX_O16
-    return GET_FILTER_OIYX_O16(OUTPUT, o, i, y, x);
+#elif defined OUTPUT_LAYOUT_OS_YXI_OSV16
+    return GET_FILTER_OS_YXI_OSV16(OUTPUT, o, i, y, x);
 #elif defined OUTPUT_LAYOUT_OS_IS_ZYX_ISV16_OSV16
     return GET_FILTER_OS_IS_ZYX_ISV16_OSV16_INDEX(OUTPUT, o, i, z, y, x, SUB_GROUP_SIZE);
 #elif defined OUTPUT_LAYOUT_IS_OS_ZYX_ISV16_OSV16

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/resample_opt.cl

@@ -246,13 +246,14 @@ KERNEL (resample_opt)(__global INPUT0_TYPE* input,
     if (feature_num >= OUTPUT_FEATURE_NUM)
         return;
 
-    unroll_for (uint out_x = 0; out_x < OUTPUT_X_BLOCK_SIZE; out_x++) {
 #ifdef SAMPLE_TYPE_NEAREST
+    unroll_for (uint out_x = 0; out_x < OUTPUT_X_BLOCK_SIZE; out_x++) {
         const int ix = floor((x + out_x) * SCALES[4]);
         const int iy = floor(y * SCALES[3]);
 
         in_vec_t res = READ_FUNC(input, INPUT0_GET_INDEX(b, feature_block, iy, ix));
 #elif defined(SAMPLE_TYPE_INTERP)
+    unroll_for (uint out_x = 0; out_x < OUTPUT_X_BLOCK_SIZE; out_x++) {
         const ACCUMULATOR_TYPE ix = TO_ACCUMULATOR_TYPE(SCALES[4]) * (x + out_x);
         const ACCUMULATOR_TYPE iy = TO_ACCUMULATOR_TYPE(SCALES[3]) * y;
 
@@ -273,22 +274,23 @@ KERNEL (resample_opt)(__global INPUT0_TYPE* input,
         const acc_vec_t bottom = TO_ACC_VEC_TYPE(bottom_left) + (TO_ACC_VEC_TYPE(bottom_right) - TO_ACC_VEC_TYPE(bottom_left)) * dx;
         acc_vec_t res = top + (bottom - top) * dy;
 #else // defined(SAMPLE_TYPE_LINEAR_ONNX)
-        const int PADDED_Y = INPUT0_SIZE_Y + PADS_BEGIN[3] + PADS_END[3];
-        const int PADDED_X = INPUT0_SIZE_X + PADS_BEGIN[4] + PADS_END[4];
+    const int PADDED_Y = INPUT0_SIZE_Y + PADS_BEGIN[3] + PADS_END[3];
+    const int PADDED_X = INPUT0_SIZE_X + PADS_BEGIN[4] + PADS_END[4];
 
+    const ACCUMULATOR_TYPE iy = FUNC_CALL(get_original_coordinate)(y, SCALES[3], OUTPUT_SIZE_Y, PADDED_Y);
+    float in_y = fmax(0, fmin(iy, PADDED_Y - 1));
+    int in_y1 = min((int)in_y, PADDED_Y - 1);
+    int in_y2 = min(in_y1 + 1, PADDED_Y - 1);
+    const ACCUMULATOR_TYPE dy1 = (in_y1 != in_y2) ? TO_ACCUMULATOR_TYPE(fabs(in_y - in_y1)) : 0.5f;
+    const ACCUMULATOR_TYPE dy2 = (in_y1 != in_y2) ? TO_ACCUMULATOR_TYPE(fabs(in_y - in_y2)) : 0.5f;
+
+    unroll_for (uint out_x = 0; out_x < OUTPUT_X_BLOCK_SIZE; out_x++) {
         const ACCUMULATOR_TYPE ix = FUNC_CALL(get_original_coordinate)(x + out_x, SCALES[4], OUTPUT_SIZE_X, PADDED_X);
-        const ACCUMULATOR_TYPE iy = FUNC_CALL(get_original_coordinate)(y, SCALES[3], OUTPUT_SIZE_Y, PADDED_Y);
-
-        float in_y = fmax(0, fmin(iy, PADDED_Y - 1));
         float in_x = fmax(0, fmin(ix, PADDED_X - 1));
-        int in_y1 = min((int)in_y, PADDED_Y - 1);
-        int in_y2 = min(in_y1 + 1, PADDED_Y - 1);
         int in_x1 = min((int)in_x, PADDED_X - 1);
         int in_x2 = min(in_x1 + 1, PADDED_X - 1);
         const ACCUMULATOR_TYPE dx1 = (in_x1 != in_x2) ? TO_ACCUMULATOR_TYPE(fabs(in_x - in_x1)) : 0.5f;
         const ACCUMULATOR_TYPE dx2 = (in_x1 != in_x2) ? TO_ACCUMULATOR_TYPE(fabs(in_x - in_x2)) : 0.5f;
-        const ACCUMULATOR_TYPE dy1 = (in_y1 != in_y2) ? TO_ACCUMULATOR_TYPE(fabs(in_y - in_y1)) : 0.5f;
-        const ACCUMULATOR_TYPE dy2 = (in_y1 != in_y2) ? TO_ACCUMULATOR_TYPE(fabs(in_y - in_y2)) : 0.5f;
 #if PADDING_USED == 1
         in_y1 -= PADS_BEGIN[3];
         in_y2 -= PADS_BEGIN[3];

inference-engine/thirdparty/clDNN/kernel_selector/core/common/jitter.cpp

@@ -332,13 +332,21 @@ JitDefinitions DataTensorJitConstant::GetDefinitions() const {
                        layout == DataLayout::b_fs_yx_fsv32 ||
                        layout == DataLayout::b_fs_yx_fsv4 ||
                        layout == DataLayout::fs_b_yx_fsv32 ||
-                       layout == DataLayout::bs_fs_yx_bsv16_fsv16) {
+                       layout == DataLayout::bs_fs_yx_bsv16_fsv16 ||
+                       layout == DataLayout::bs_fs_yx_bsv4_fsv4 ||
+                       layout == DataLayout::bs_fs_yx_bsv4_fsv2 ||
+                       layout == DataLayout::bs_fs_yx_bsv32_fsv16 ||
+                       layout == DataLayout::bs_fs_yx_bsv32_fsv32) {
                 auto layout_str = toString(layout);
                 index_func_val = "GET_DATA_" + layout_str + "_INDEX(" + _name + ", b, f, y, x)";
                 raw_index_func_val = "GET_DATA_" + layout_str + "_INDEX(" + _name + ", b, f, y, x)";
                 if (layout == DataLayout::b_fs_yx_fsv16 ||
                     layout == DataLayout::b_fs_yx_fsv32 ||
                     layout == DataLayout::b_fs_yx_fsv4  ||
+                    layout == DataLayout::bs_fs_yx_bsv32_fsv32  ||
+                    layout == DataLayout::bs_fs_yx_bsv32_fsv16  ||
+                    layout == DataLayout::bs_fs_yx_bsv4_fsv4  ||
+                    layout == DataLayout::bs_fs_yx_bsv4_fsv2  ||
                     layout == DataLayout::bs_fs_yx_bsv16_fsv16)
                     safe_index_func_val = "GET_DATA_" + layout_str + "_INDEX_SAFE(" + _name + ", b, f, y, x)";
                 else

inference-engine/thirdparty/clDNN/kernel_selector/core/kernel_selector_common.cpp

@@ -104,6 +104,10 @@ std::string toString(DataLayout l) {
         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_zyx_bsv16_fsv16: return "BS_FS_ZYX_BSV16_FSV16";
+        case kernel_selector::DataLayout::bs_fs_yx_bsv4_fsv4:    return "BS_FS_YX_BSV4_FSV4";
+        case kernel_selector::DataLayout::bs_fs_yx_bsv4_fsv2:    return "BS_FS_YX_BSV4_FSV2";
+        case kernel_selector::DataLayout::bs_fs_yx_bsv32_fsv32:  return "BS_FS_YX_BSV32_FSV32";
+        case kernel_selector::DataLayout::bs_fs_yx_bsv32_fsv16:  return "BS_FS_YX_BSV32_FSV16";
         case kernel_selector::DataLayout::nv12:                  return "NV12";
         case kernel_selector::DataLayout::image_2d_rgba:         return "IMAGE_2D_RGBA";
         default:
@@ -308,6 +312,8 @@ std::string toString(WeightsLayout layout) {
         case WeightsLayout::os_is_zyx_osv16_isv16:                       return "OS_IS_ZYX_OSV16_ISV16";
         case WeightsLayout::os_is_zyx_osv32_isv16:                       return "OS_IS_ZYX_OSV32_ISV16";
         case WeightsLayout::os_is_zyx_osv64_isv16:                       return "OS_IS_ZYX_OSV64_ISV16";
+        case WeightsLayout::o_is_yx_isv16:                               return "O_IS_YX_ISV16";
+        case WeightsLayout::os_yxi_osv16:                                return "OS_YXI_OSV16";
         case WeightsLayout::os_iyx_osv16:                                return "OS_IYX_OSV16";
         case WeightsLayout::os_iyx_osv32:                                return "OS_IYX_OSV32";
         case WeightsLayout::os_iyx_osv32__ai32:                          return "OS_IYX_OSV32__AI32";
@@ -354,6 +360,9 @@ std::string toString(WeightsLayout layout) {
         case WeightsLayout::os_is_yx_isv8_osv16_isv2:                    return "OS_IS_YX_ISV8_OSV16_ISV2";
         case WeightsLayout::os_is_osv32_isv32_swizzled_by_4:             return "OS_IS_OSV32_ISV32_SWIZZLED_BY_4";
         case WeightsLayout::os_i_yxs_osv4_yxsv4:                         return "OS_I_YXS_OSV4_YXSV4";
+        case WeightsLayout::os_is_yx_osa4_isa8_osv8_isv4:                return "OS_IS_YX_OSA4_ISA8_OSV8_ISV4";
+        case WeightsLayout::os_is_yx_osv8_isv4:                          return "OS_IS_YX_OSV8_ISV4";
+        case WeightsLayout::os_is_yx_osv8_isv2:                          return "OS_IS_YX_OSV8_ISV2";
         case WeightsLayout::goiyx:                                       return "GOIYX";
         case WeightsLayout::gioyx:                                       return "GIOYX";
         case WeightsLayout::gyxio:                                       return "GYXIO";
@@ -375,6 +384,9 @@ std::string toString(WeightsLayout layout) {
         case WeightsLayout::gs_oi_yxs_gsv4_yxsv4:                        return "GS_OI_YXS_GSV4_YXSV4";
         case WeightsLayout::gs_oi_yxs_gsv16_yxsv4:                       return "GS_OI_YXS_GSV16_YXSV4";
         case WeightsLayout::gs_oi_yxs_gsv32_yxsv4:                       return "GS_OI_YXS_GSV32_YXSV4";
+        case WeightsLayout::g_os_is_yx_osa4_isa8_osv8_isv4:              return "G_OS_IS_YX_OSA4_ISA8_OSV8_ISV4";
+        case WeightsLayout::g_os_is_yx_osa4_isa8_osv8_isv2:              return "G_OS_IS_YX_OSA4_ISA8_OSV8_ISV2";
+        case WeightsLayout::os_is_yx_osa4_isa8_osv8_isv2:                return "OS_IS_YX_OSA4_ISA8_OSV8_ISV2";
         case WeightsLayout::g_os_is_yx_isv16_osv16:                      return "G_OS_IS_YX_ISV16_OSV16";
         case WeightsLayout::g_os_is_yx_osv16_isv4:                       return "G_OS_IS_YX_OSV16_ISV4";
         case WeightsLayout::g_os_is_zyx_osv16_isv16:                     return "G_OS_IS_ZYX_OSV16_ISV16";

inference-engine/thirdparty/clDNN/src/graph_optimizer/prepare_primitive_fusing.cpp

@@ -399,6 +399,12 @@ void prepare_primitive_fusing::fuse_simple_primitives(program &p) {
                   _lo.get_optimization_attributes().bs_fs_yx_bsv16_fsv16_network)) && node.get_primitive()->groups == 1)
                 return true;
 
+            if (node.get_output_layout().format == format::bs_fs_yx_bsv32_fsv32 || _lo.is_format_optimized(node, format::bs_fs_yx_bsv32_fsv32))
+                return true;
+
+            if (node.get_output_layout().format == format::bs_fs_yx_bsv32_fsv16 || _lo.is_format_optimized(node, format::bs_fs_yx_bsv32_fsv16))
+                return true;
+
             auto in_dt = node.get_dependency(0).get_output_layout().data_type;
 
             // TODO: check if that's enough for correct work
@@ -476,7 +482,7 @@ void prepare_primitive_fusing::fuse_simple_primitives(program &p) {
         };
 
         auto pooling_supports_fusings = [](pooling_node& node) -> bool {
-            auto pooling_mode = node.as<pooling>().get_primitive()->mode;
+            auto pooling_mode = node.get_primitive()->mode;
 
             if (pooling_mode != cldnn::pooling_mode::max_with_argmax)
                 return true;

inference-engine/thirdparty/clDNN/src/impls/ocl/activation.cpp

@@ -97,16 +97,32 @@ attach_activation_impl::attach_activation_impl() {
         std::make_tuple(data_types::f16, format::bfzyx),
         std::make_tuple(data_types::i8, format::bfzyx),
         std::make_tuple(data_types::i32, format::bfzyx),
+
         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),
         std::make_tuple(data_types::u8, format::b_fs_zyx_fsv16),
+
         std::make_tuple(data_types::f32, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::f16, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::i8, format::bs_fs_zyx_bsv16_fsv16),
+        std::make_tuple(data_types::u8, format::bs_fs_zyx_bsv16_fsv16),
+
         std::make_tuple(data_types::f32, format::bs_fs_yx_bsv16_fsv16),
         std::make_tuple(data_types::f16, format::bs_fs_yx_bsv16_fsv16),
         std::make_tuple(data_types::i8, format::bs_fs_yx_bsv16_fsv16),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv16_fsv16),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv16),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv32),
+
         // bfwzyx
         std::make_tuple(data_types::f32, format::bfwzyx),
         std::make_tuple(data_types::f16, format::bfwzyx),

inference-engine/thirdparty/clDNN/src/impls/ocl/concatenation.cpp

@@ -134,6 +134,21 @@ attach_concatenation_impl::attach_concatenation_impl() {
         std::make_tuple(data_types::u8, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::i32, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::i64, format::bs_fs_zyx_bsv16_fsv16),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i32, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i64, format::bs_fs_yx_bsv32_fsv32),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i32, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i64, format::bs_fs_yx_bsv32_fsv16),
+
         std::make_tuple(data_types::f32, format::bs_fs_yx_bsv16_fsv16),
         std::make_tuple(data_types::f16, format::bs_fs_yx_bsv16_fsv16),
         std::make_tuple(data_types::f16, format::b_fs_yx_fsv16),

inference-engine/thirdparty/clDNN/src/impls/ocl/convolution.cpp

@@ -195,10 +195,31 @@ attach_convolution_impl::attach_convolution_impl() {
         std::make_tuple(data_types::f16, format::b_fs_zyx_fsv16),
         std::make_tuple(data_types::f32, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::f16, format::bs_fs_zyx_bsv16_fsv16),
+
         std::make_tuple(data_types::f32, format::bs_fs_yx_bsv16_fsv16),
         std::make_tuple(data_types::f16, format::bs_fs_yx_bsv16_fsv16),
         std::make_tuple(data_types::u8, format::bs_fs_yx_bsv16_fsv16),
         std::make_tuple(data_types::i8, format::bs_fs_yx_bsv16_fsv16),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv32),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv16),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv4_fsv4),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv4_fsv4),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv4_fsv4),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv4_fsv4),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv4_fsv2),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv4_fsv2),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv4_fsv2),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv4_fsv2),
     });
 }
 

inference-engine/thirdparty/clDNN/src/impls/ocl/eltwise.cpp

@@ -124,62 +124,102 @@ namespace detail {
 attach_eltwise_impl::attach_eltwise_impl() {
     implementation_map<eltwise>::add(impl_types::ocl, eltwise_impl::create, {
         std::make_tuple(data_types::f32, format::yxfb),
-        std::make_tuple(data_types::u8, format::bfyx),
         std::make_tuple(data_types::f16, format::yxfb),
         std::make_tuple(data_types::i8, format::yxfb),
         std::make_tuple(data_types::i32, format::yxfb),
         std::make_tuple(data_types::i64, format::yxfb),
+
         std::make_tuple(data_types::f32, format::bfyx),
         std::make_tuple(data_types::f16, format::bfyx),
+        std::make_tuple(data_types::u8, format::bfyx),
         std::make_tuple(data_types::i8, format::bfyx),
         std::make_tuple(data_types::i32, format::bfyx),
         std::make_tuple(data_types::i64, format::bfyx),
+
         std::make_tuple(data_types::f32, format::byxf),
         std::make_tuple(data_types::f16, format::byxf),
         std::make_tuple(data_types::i8, format::byxf),
         std::make_tuple(data_types::i32, format::byxf),
         std::make_tuple(data_types::i64, format::byxf),
+
         std::make_tuple(data_types::f16, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::f32, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::i8, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::u8, format::b_fs_yx_fsv16),
+
         std::make_tuple(data_types::f32, format::bfzyx),
         std::make_tuple(data_types::f16, format::bfzyx),
         std::make_tuple(data_types::i8, format::bfzyx),
         std::make_tuple(data_types::u8, format::bfzyx),
         std::make_tuple(data_types::i32, format::bfzyx),
         std::make_tuple(data_types::i64, format::bfzyx),
+
         std::make_tuple(data_types::f32, format::bfwzyx),
         std::make_tuple(data_types::f16, format::bfwzyx),
         std::make_tuple(data_types::i8, format::bfwzyx),
         std::make_tuple(data_types::u8, format::bfwzyx),
         std::make_tuple(data_types::i32, format::bfwzyx),
         std::make_tuple(data_types::i64, format::bfwzyx),
+
         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),
         std::make_tuple(data_types::u8, format::b_fs_zyx_fsv16),
         std::make_tuple(data_types::i32, format::b_fs_zyx_fsv16),
         std::make_tuple(data_types::i64, format::b_fs_zyx_fsv16),
+
         std::make_tuple(data_types::f32, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::f16, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::i8, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::i32, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::i64, format::bs_fs_zyx_bsv16_fsv16),
+
         std::make_tuple(data_types::f32, format::bs_fs_yx_bsv16_fsv16),
         std::make_tuple(data_types::f16, format::bs_fs_yx_bsv16_fsv16),
+
         std::make_tuple(data_types::i8, format::b_fs_yx_fsv4),
         std::make_tuple(data_types::u8, format::b_fs_yx_fsv4),
         std::make_tuple(data_types::f32, format::b_fs_yx_fsv4),
+
         std::make_tuple(data_types::i8, format::b_fs_yx_fsv32),
         std::make_tuple(data_types::u8, format::b_fs_yx_fsv32),
         std::make_tuple(data_types::f32, format::b_fs_yx_fsv32),
         std::make_tuple(data_types::f16, format::b_fs_yx_fsv32),
+
         std::make_tuple(data_types::i8, format::b_fs_zyx_fsv32),
         std::make_tuple(data_types::u8, format::b_fs_zyx_fsv32),
         std::make_tuple(data_types::f32, format::b_fs_zyx_fsv32),
         std::make_tuple(data_types::f16, format::b_fs_zyx_fsv32),
+
         std::make_tuple(data_types::f16, format::fs_b_yx_fsv32),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i32, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i64, format::bs_fs_yx_bsv32_fsv32),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i32, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i64, format::bs_fs_yx_bsv32_fsv16),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv4_fsv4),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv4_fsv4),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv4_fsv4),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv4_fsv4),
+        std::make_tuple(data_types::i32, format::bs_fs_yx_bsv4_fsv4),
+        std::make_tuple(data_types::i64, format::bs_fs_yx_bsv4_fsv4),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv4_fsv2),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv4_fsv2),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv4_fsv2),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv4_fsv2),
+        std::make_tuple(data_types::i32, format::bs_fs_yx_bsv4_fsv2),
+        std::make_tuple(data_types::i64, format::bs_fs_yx_bsv4_fsv2),
     });
 }
 

inference-engine/thirdparty/clDNN/src/impls/ocl/lrn.cpp

@@ -58,22 +58,41 @@ attach_lrn_impl::attach_lrn_impl() {
         std::make_tuple(data_types::f16, format::yxfb),
         std::make_tuple(data_types::u8, format::yxfb),
         std::make_tuple(data_types::i8, format::yxfb),
+
         std::make_tuple(data_types::f32, format::bfyx),
         std::make_tuple(data_types::f16, format::bfyx),
         std::make_tuple(data_types::u8, format::bfyx),
         std::make_tuple(data_types::i8, format::bfyx),
+
         std::make_tuple(data_types::f32, format::byxf),
         std::make_tuple(data_types::f16, format::byxf),
         std::make_tuple(data_types::u8, format::byxf),
         std::make_tuple(data_types::i8, format::byxf),
+
         std::make_tuple(data_types::f32, format::b_fs_yx_fsv4),
         std::make_tuple(data_types::f16, format::b_fs_yx_fsv4),
         std::make_tuple(data_types::u8, format::b_fs_yx_fsv4),
         std::make_tuple(data_types::i8, format::b_fs_yx_fsv4),
+
         std::make_tuple(data_types::f32, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::f16, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::u8, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::i8, format::b_fs_yx_fsv16),
+
+        std::make_tuple(data_types::f32, format::b_fs_yx_fsv32),
+        std::make_tuple(data_types::f16, format::b_fs_yx_fsv32),
+        std::make_tuple(data_types::i8, format::b_fs_yx_fsv32),
+        std::make_tuple(data_types::u8, format::b_fs_yx_fsv32),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv32),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv16),
     });
 }
 

inference-engine/thirdparty/clDNN/src/impls/ocl/mvn.cpp

@@ -60,24 +60,35 @@ attach_mvn_impl::attach_mvn_impl() {
         std::make_tuple(data_types::f16, format::bfyx),
         std::make_tuple(data_types::u8, format::bfyx),
         std::make_tuple(data_types::i8, format::bfyx),
+
         std::make_tuple(data_types::f32, format::yxfb),
         std::make_tuple(data_types::f16, format::yxfb),
+
         std::make_tuple(data_types::f32, format::byxf),
         std::make_tuple(data_types::f16, format::byxf),
+
         std::make_tuple(data_types::f32, format::bfzyx),
         std::make_tuple(data_types::f16, format::bfzyx),
         std::make_tuple(data_types::u8, format::bfzyx),
         std::make_tuple(data_types::i8, format::bfzyx),
+
         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::u8, format::b_fs_zyx_fsv16),
         std::make_tuple(data_types::i8, format::b_fs_zyx_fsv16),
+
         std::make_tuple(data_types::f32, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::f16, format::bs_fs_zyx_bsv16_fsv16),
+
         std::make_tuple(data_types::f32, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::f16, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::u8, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::i8, format::b_fs_yx_fsv16),
+
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv32),
+
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv16),
     });
 }
 

inference-engine/thirdparty/clDNN/src/impls/ocl/pooling.cpp

@@ -162,48 +162,69 @@ attach_pooling_impl::attach_pooling_impl() {
         std::make_tuple(data_types::f16, format::bfyx),
         std::make_tuple(data_types::i8, format::bfyx),
         std::make_tuple(data_types::u8, format::bfyx),
+
         std::make_tuple(data_types::f32, format::yxfb),
         std::make_tuple(data_types::f16, format::yxfb),
         std::make_tuple(data_types::i8, format::yxfb),
         std::make_tuple(data_types::u8, format::yxfb),
+
         std::make_tuple(data_types::f32, format::byxf),
         std::make_tuple(data_types::f16, format::byxf),
         std::make_tuple(data_types::i8, format::byxf),
         std::make_tuple(data_types::u8, format::byxf),
+
         std::make_tuple(data_types::f16, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::f32, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::i8, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::u8, format::b_fs_yx_fsv16),
+
         std::make_tuple(data_types::f32, format::bs_fs_yx_bsv16_fsv16),
         std::make_tuple(data_types::f16, format::bs_fs_yx_bsv16_fsv16),
         std::make_tuple(data_types::i8, format::bs_fs_yx_bsv16_fsv16),
         std::make_tuple(data_types::u8, format::bs_fs_yx_bsv16_fsv16),
+
         std::make_tuple(data_types::f32, format::bfzyx),
         std::make_tuple(data_types::f16, format::bfzyx),
         std::make_tuple(data_types::i8, format::bfzyx),
         std::make_tuple(data_types::u8, format::bfzyx),
+
         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),
         std::make_tuple(data_types::u8, format::b_fs_zyx_fsv16),
+
         std::make_tuple(data_types::f32, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::f16, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::i8, format::bs_fs_zyx_bsv16_fsv16),
         std::make_tuple(data_types::u8, format::bs_fs_zyx_bsv16_fsv16),
+
         std::make_tuple(data_types::f32, format::b_fs_yx_fsv4),
         std::make_tuple(data_types::f16, format::b_fs_yx_fsv4),
         std::make_tuple(data_types::i8, format::b_fs_yx_fsv4),
         std::make_tuple(data_types::u8, format::b_fs_yx_fsv4),
+
         std::make_tuple(data_types::i8, format::b_fs_yx_fsv32),
         std::make_tuple(data_types::u8, format::b_fs_yx_fsv32),
         std::make_tuple(data_types::f32, format::b_fs_yx_fsv32),
         std::make_tuple(data_types::f16, format::b_fs_yx_fsv32),
+
         std::make_tuple(data_types::i8, format::b_fs_zyx_fsv32),
         std::make_tuple(data_types::u8, format::b_fs_zyx_fsv32),
         std::make_tuple(data_types::f32, format::b_fs_zyx_fsv32),
         std::make_tuple(data_types::f16, format::b_fs_zyx_fsv32),
+
         std::make_tuple(data_types::f16, format::fs_b_yx_fsv32),
         std::make_tuple(data_types::f32, format::fs_b_yx_fsv32),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv32),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv16),
     });
 }
 

inference-engine/thirdparty/clDNN/src/impls/ocl/quantize.cpp

@@ -149,9 +149,16 @@ attach_quantize_impl::attach_quantize_impl() {
         std::make_tuple(data_types::i8, format::bfwzyx),
         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::f32, format::bs_fs_yx_bsv16_fsv16),
-        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv16_fsv16),
-        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv16_fsv16),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv32),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv16),
     });
 }
 

inference-engine/thirdparty/clDNN/src/impls/ocl/resample.cpp

@@ -147,25 +147,46 @@ attach_resample_impl::attach_resample_impl() {
     implementation_map<resample>::add(impl_types::ocl, resample_impl::create, {
         std::make_tuple(data_types::f32, format::yxfb),
         std::make_tuple(data_types::f16, format::yxfb),
-        std::make_tuple(data_types::f32, format::bfyx),
-        std::make_tuple(data_types::f16, format::bfyx),
-        std::make_tuple(data_types::f16, format::fs_b_yx_fsv32),
-        std::make_tuple(data_types::f16, format::b_fs_yx_fsv16),
-        std::make_tuple(data_types::f32, format::b_fs_yx_fsv16),
+
         std::make_tuple(data_types::f32, format::byxf),
         std::make_tuple(data_types::f16, format::byxf),
-        std::make_tuple(data_types::f32, format::bfzyx),
-        std::make_tuple(data_types::f16, format::bfzyx),
+
+        std::make_tuple(data_types::f32, format::bfyx),
+        std::make_tuple(data_types::f16, format::bfyx),
         std::make_tuple(data_types::u8, format::bfyx),
         std::make_tuple(data_types::i8, format::bfyx),
+
+        std::make_tuple(data_types::f32, format::bfzyx),
+        std::make_tuple(data_types::f16, format::bfzyx),
         std::make_tuple(data_types::u8, format::bfzyx),
         std::make_tuple(data_types::i8, format::bfzyx),
+
+        std::make_tuple(data_types::f16, format::fs_b_yx_fsv32),
+
+        std::make_tuple(data_types::f32, format::b_fs_yx_fsv16),
+        std::make_tuple(data_types::f16, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::u8, format::b_fs_yx_fsv16),
         std::make_tuple(data_types::i8, format::b_fs_yx_fsv16),
+
         std::make_tuple(data_types::f32, format::b_fs_yx_fsv4),
         std::make_tuple(data_types::f16, format::b_fs_yx_fsv4),
         std::make_tuple(data_types::u8, format::b_fs_yx_fsv4),
         std::make_tuple(data_types::i8, format::b_fs_yx_fsv4),
+
+        std::make_tuple(data_types::f32, format::b_fs_yx_fsv32),
+        std::make_tuple(data_types::f16, format::b_fs_yx_fsv32),
+        std::make_tuple(data_types::u8, format::b_fs_yx_fsv32),
+        std::make_tuple(data_types::i8, format::b_fs_yx_fsv32),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv32),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv32),
+
+        std::make_tuple(data_types::f32, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::f16, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::u8, format::bs_fs_yx_bsv32_fsv16),
+        std::make_tuple(data_types::i8, format::bs_fs_yx_bsv32_fsv16),
     });
 }
 

inference-engine/thirdparty/clDNN/src/include/to_string_utils.h

@@ -91,6 +91,14 @@ inline std::string fmt_to_str(format fmt) {
             return "fs_b_yx_fsv32";
         case format::bs_fs_yx_bsv16_fsv16:
             return "bs_fs_yx_bsv16_fsv16";
+        case format::bs_fs_yx_bsv32_fsv16:
+            return "bs_fs_yx_bsv32_fsv16";
+        case format::bs_fs_yx_bsv4_fsv2:
+            return "bs_fs_yx_bsv4_fsv2";
+        case format::bs_fs_yx_bsv4_fsv4:
+            return "bs_fs_yx_bsv4_fsv4";
+        case format::bs_fs_yx_bsv32_fsv32:
+            return "bs_fs_yx_bsv32_fsv32";
         case format::b_fs_zyx_fsv16:
             return "b_fs_zyx_fsv16";
         case format::bs_fs_zyx_bsv16_fsv16:
@@ -122,6 +130,10 @@ inline std::string fmt_to_str(format fmt) {
             return "image_2d_weights_winograd_6x3_s1_fbxyb";
         case format::image_2d_weights_winograd_6x3_s1_xfbyb:
             return "image_2d_weights_winograd_6x3_s1_xfbyb";
+        case format::os_yxi_osv16:
+            return "os_yxi_osv16";
+        case format::o_is_yx_isv16:
+            return "o_is_yx_isv16";
         case format::os_iyx_osv16:
             return "os_iyx_osv16";
         case format::os_is_yx_osv16_isv16:
@@ -142,6 +154,14 @@ inline std::string fmt_to_str(format fmt) {
             return "os_is_zyx_isa8_osv8_isv4";
         case format::os_is_zyx_isa8_osv16_isv4:
             return "os_is_zyx_isa8_osv16_isv4";
+        case format::os_is_yx_osa4_isa8_osv8_isv2:
+            return "os_is_yx_osa4_isa8_osv8_isv2";
+        case format::g_os_is_yx_osa4_isa8_osv8_isv2:
+            return "g_os_is_yx_osa4_isa8_osv8_isv2";
+        case format::g_os_is_yx_osa4_isa8_osv8_isv4:
+            return "g_os_is_yx_osa4_isa8_osv8_isv4";
+        case format::os_is_yx_osa4_isa8_osv8_isv4:
+            return "os_is_yx_osa4_isa8_osv8_isv4";
         case format::os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4:
             return "os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4";
         case format::os_is_zyx_osa4_isa8_osv8_isv4_swizzled_by_4:
@@ -150,6 +170,10 @@ inline std::string fmt_to_str(format fmt) {
             return "os_is_yx_isa8_osv8_isv4_swizzled_by_4";
         case format::is_o32_yx_isv32_swizzled_by_4:
             return "is_o32_yx_isv32_swizzled_by_4";
+        case format::os_is_yx_osv8_isv2:
+            return "os_is_yx_osv8_isv2";
+        case format::os_is_yx_osv8_isv4:
+            return "os_is_yx_osv8_isv4";
         case format::os_is_yx_osv16_isv4:
             return "os_is_yx_osv16_isv4";
         case format::os_is_yx_osv32_isv4_swizzled_by_2:

inference-engine/thirdparty/clDNN/src/kernel_selector_helper.cpp

@@ -132,6 +132,14 @@ kernel_selector::data_layout to_data_layout(format f) {
             return kernel_selector::data_layout::bs_fs_zyx_bsv16_fsv16;
         case format::bs_fs_yx_bsv16_fsv16:
             return kernel_selector::data_layout::bs_fs_yx_bsv16_fsv16;
+        case format::bs_fs_yx_bsv32_fsv16:
+            return kernel_selector::data_layout::bs_fs_yx_bsv32_fsv16;
+        case format::bs_fs_yx_bsv4_fsv4:
+            return kernel_selector::data_layout::bs_fs_yx_bsv4_fsv4;
+        case format::bs_fs_yx_bsv4_fsv2:
+            return kernel_selector::data_layout::bs_fs_yx_bsv4_fsv2;
+        case format::bs_fs_yx_bsv32_fsv32:
+            return kernel_selector::data_layout::bs_fs_yx_bsv32_fsv32;
         case format::nv12:
             return kernel_selector::data_layout::nv12;
         case format::image_2d_rgba:
@@ -179,6 +187,14 @@ cldnn::format from_data_layout(kernel_selector::data_layout l) {
             return cldnn::format::b_fs_yx_fsv4;
         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_yx_bsv32_fsv16:
+            return cldnn::format::bs_fs_yx_bsv32_fsv16;
+        case kernel_selector::data_layout::bs_fs_yx_bsv4_fsv2:
+            return cldnn::format::bs_fs_yx_bsv4_fsv2;
+        case kernel_selector::data_layout::bs_fs_yx_bsv4_fsv4:
+            return cldnn::format::bs_fs_yx_bsv4_fsv4;
+        case kernel_selector::data_layout::bs_fs_yx_bsv32_fsv32:
+            return cldnn::format::bs_fs_yx_bsv32_fsv32;
         case kernel_selector::data_layout::nv12:
             return cldnn::format::nv12;
         case kernel_selector::data_layout::image_2d_rgba:
@@ -202,6 +218,10 @@ kernel_selector::weights_layout to_weights_layout(format f, bool is_grouped) {
         case format::yxfb:
         case format::yxio:
             return kernel_selector::weights_layout::yxio;
+        case format::os_yxi_osv16:
+            return kernel_selector::weights_layout::os_yxi_osv16;
+        case format::o_is_yx_isv16:
+            return kernel_selector::weights_layout::o_is_yx_isv16;
         case format::os_iyx_osv16:
             return kernel_selector::weights_layout::os_iyx_osv16;
         case format::os_is_yx_osv16_isv16:
@@ -224,6 +244,14 @@ kernel_selector::weights_layout to_weights_layout(format f, bool is_grouped) {
             return kernel_selector::weights_layout::image_2d_weights_winograd_6x3_s1_fbxyb;
         case format::image_2d_weights_winograd_6x3_s1_xfbyb:
             return kernel_selector::weights_layout::image_2d_weights_winograd_6x3_s1_xfbyb;
+        case format::os_is_yx_osa4_isa8_osv8_isv2:
+            return kernel_selector::weights_layout::os_is_yx_osa4_isa8_osv8_isv2;
+        case format::g_os_is_yx_osa4_isa8_osv8_isv2:
+            return kernel_selector::weights_layout::g_os_is_yx_osa4_isa8_osv8_isv2;
+        case format::g_os_is_yx_osa4_isa8_osv8_isv4:
+            return kernel_selector::weights_layout::g_os_is_yx_osa4_isa8_osv8_isv4;
+        case format::os_is_yx_osa4_isa8_osv8_isv4:
+            return kernel_selector::weights_layout::os_is_yx_osa4_isa8_osv8_isv4;
         case format::os_is_yx_isa8_osv8_isv4:
             return kernel_selector::weights_layout::os_is_yx_isa8_osv8_isv4;
         case format::os_is_yx_isa8_osv16_isv4:
@@ -232,6 +260,10 @@ kernel_selector::weights_layout to_weights_layout(format f, bool is_grouped) {
             return kernel_selector::weights_layout::os_is_zyx_isa8_osv8_isv4;
         case format::os_is_zyx_isa8_osv16_isv4:
             return kernel_selector::weights_layout::os_is_zyx_isa8_osv16_isv4;
+        case format::os_is_yx_osv8_isv2:
+            return kernel_selector::weights_layout::os_is_yx_osv8_isv2;
+        case format::os_is_yx_osv8_isv4:
+            return kernel_selector::weights_layout::os_is_yx_osv8_isv4;
         case format::os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4:
             return kernel_selector::weights_layout::os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4;
         case format::os_is_zyx_osa4_isa8_osv8_isv4_swizzled_by_4:
@@ -354,6 +386,10 @@ cldnn::format::type from_weights_layout(kernel_selector::weights_layout l) {
             return cldnn::format::iyxo;
         case kernel_selector::weights_layout::yxio:
             return cldnn::format::yxfb;
+        case kernel_selector::weights_layout::os_yxi_osv16:
+            return cldnn::format::os_yxi_osv16;
+        case kernel_selector::weights_layout::o_is_yx_isv16:
+            return cldnn::format::o_is_yx_isv16;
         case kernel_selector::weights_layout::os_iyx_osv16:
             return cldnn::format::os_iyx_osv16;
         case kernel_selector::weights_layout::os_is_yx_isv16_osv16:
@@ -384,6 +420,14 @@ cldnn::format::type from_weights_layout(kernel_selector::weights_layout l) {
             return cldnn::format::image_2d_weights_winograd_6x3_s1_fbxyb;
         case kernel_selector::weights_layout::image_2d_weights_winograd_6x3_s1_xfbyb:
             return cldnn::format::image_2d_weights_winograd_6x3_s1_xfbyb;
+        case kernel_selector::weights_layout::os_is_yx_osa4_isa8_osv8_isv2:
+            return cldnn::format::os_is_yx_osa4_isa8_osv8_isv2;
+        case kernel_selector::weights_layout::g_os_is_yx_osa4_isa8_osv8_isv2:
+            return cldnn::format::g_os_is_yx_osa4_isa8_osv8_isv2;
+        case kernel_selector::weights_layout::g_os_is_yx_osa4_isa8_osv8_isv4:
+            return cldnn::format::g_os_is_yx_osa4_isa8_osv8_isv4;
+        case kernel_selector::weights_layout::os_is_yx_osa4_isa8_osv8_isv4:
+            return cldnn::format::os_is_yx_osa4_isa8_osv8_isv4;
         case kernel_selector::weights_layout::os_is_yx_isa8_osv8_isv4:
             return cldnn::format::os_is_yx_isa8_osv8_isv4;
         case kernel_selector::weights_layout::os_is_zyx_isa8_osv8_isv4:
@@ -422,6 +466,10 @@ cldnn::format::type from_weights_layout(kernel_selector::weights_layout l) {
             return cldnn::format::is_os_zyx_isv16_osv16;
         case kernel_selector::weights_layout::is_os_yx_isv16_osv16:
             return cldnn::format::is_os_yx_isv16_osv16;
+        case kernel_selector::weights_layout::os_is_yx_osv8_isv2:
+            return cldnn::format::os_is_yx_osv8_isv2;
+        case kernel_selector::weights_layout::os_is_yx_osv8_isv4:
+            return cldnn::format::os_is_yx_osv8_isv4;
         case kernel_selector::weights_layout::os_is_zyx_isv8_osv16_isv2:
             return cldnn::format::os_is_zyx_isv8_osv16_isv2;
         case kernel_selector::weights_layout::os_zyxi_osv16:

inference-engine/thirdparty/clDNN/src/layout_optimizer.cpp

@@ -955,6 +955,9 @@ bool layout_optimizer::is_format_optimized(const convolution_node& node, const f
             return convolution_fs_b_yx_fsv32_opt(input_layout, output_layout, weights_layout, prim);
         case format::bs_fs_yx_bsv16_fsv16:
             return convolution_bs_fs_yx_bsv16_fsv16_opt(input_layout, output_layout, weights_layout, prim);
+        case format::bs_fs_yx_bsv32_fsv32:
+        case format::bs_fs_yx_bsv32_fsv16:
+            return false;
         default:
             throw std::invalid_argument(
                 "[Layout optimizer] Other formats in is_format_optimized(...) method are not implemented!");

inference-engine/thirdparty/clDNN/src/program_helpers.cpp

@@ -143,6 +143,14 @@ std::pair<bool, bool> program_helpers::are_layouts_identical(layout const& l1, l
         (l2.format == format::b_fs_zyx_fsv32 && l1.format != format::b_fs_zyx_fsv32) ||
         (l1.format == format::b_fs_zyx_fsv16 && l2.format != format::b_fs_zyx_fsv16) ||
         (l2.format == format::b_fs_zyx_fsv16 && l1.format != format::b_fs_zyx_fsv16) ||
+        (l1.format == format::bs_fs_yx_bsv4_fsv4 && l2.format != format::bs_fs_yx_bsv4_fsv4) ||
+        (l2.format == format::bs_fs_yx_bsv4_fsv4 && l1.format != format::bs_fs_yx_bsv4_fsv4) ||
+        (l1.format == format::bs_fs_yx_bsv4_fsv2 && l2.format != format::bs_fs_yx_bsv4_fsv2) ||
+        (l2.format == format::bs_fs_yx_bsv4_fsv2 && l1.format != format::bs_fs_yx_bsv4_fsv2) ||
+        (l1.format == format::bs_fs_yx_bsv32_fsv16 && l2.format != format::bs_fs_yx_bsv32_fsv16) ||
+        (l2.format == format::bs_fs_yx_bsv32_fsv16 && l1.format != format::bs_fs_yx_bsv32_fsv16) ||
+        (l1.format == format::bs_fs_yx_bsv32_fsv32 && l2.format != format::bs_fs_yx_bsv32_fsv32) ||
+        (l2.format == format::bs_fs_yx_bsv32_fsv32 && l1.format != format::bs_fs_yx_bsv32_fsv32) ||
         (l1.format == format::bs_fs_yx_bsv16_fsv16 && l2.format != format::bs_fs_yx_bsv16_fsv16) ||
         (l2.format == format::bs_fs_yx_bsv16_fsv16 && l1.format != format::bs_fs_yx_bsv16_fsv16) ||
         (l1.format == format::bs_fs_zyx_bsv16_fsv16 && l2.format != format::bs_fs_zyx_bsv16_fsv16) ||

inference-engine/thirdparty/clDNN/tests/test_cases/quantize_gpu_test.cpp

@@ -583,3 +583,242 @@ TEST(quantize_gpu, quantize_levels_256_3d_unsigned) {
         EXPECT_EQ(output_ptr[i], ref_data[i]) << " i=" << i;
     }
 }
+
+struct quantize_random_test_params {
+    data_types  input_type;
+    data_types  output_type;
+
+    tensor      input_size;
+
+    format::type in_format;
+    format::type out_format;
+
+    int32_t inputs_num;  // 5: ref
+};
+
+struct quantize_random_test : testing::TestWithParam<quantize_random_test_params>
+{
+    template <typename T>
+    void fill_typed(memory::ptr src, memory::ptr dst) {
+        auto size = dst->get_layout().size;
+        size_t b = size.batch[0];
+        size_t f = size.feature[0];
+        size_t x = size.spatial[0];
+        size_t y = size.spatial[1];
+
+        mem_lock<T> data{src, get_test_stream()};
+        mem_lock<T> ptr{dst, get_test_stream()};
+        for (size_t bi = 0; bi < b; ++bi) {
+            for (size_t fi = 0; fi < f; ++fi) {
+                for (size_t yi = 0; yi < y; ++yi) {
+                    for (size_t xi = 0; xi < x; ++xi) {
+                        auto coords = tensor(batch(bi), feature(fi), spatial(xi, yi, 0, 0));
+                        auto offset = dst->get_layout().get_linear_offset(coords);
+                        auto src_offset = src->get_layout().get_linear_offset(coords);
+                        ptr[offset] = data[src_offset];
+                    }
+                }
+            }
+        }
+    }
+
+    template <typename T>
+    void fill_random_typed(memory::ptr mem, int min, int max, int k) {
+        auto size = mem->get_layout().size;
+        size_t b = size.batch[0];
+        size_t f = size.feature[0];
+        size_t x = size.spatial[0];
+        size_t y = size.spatial[1];
+
+        auto data = generate_random_4d<T>(b, f, y, x, min, max, k);
+        mem_lock<T> ptr{mem, get_test_stream()};
+        for (size_t bi = 0; bi < b; ++bi) {
+            for (size_t fi = 0; fi < f; ++fi) {
+                for (size_t yi = 0; yi < y; ++yi) {
+                    for (size_t xi = 0; xi < x; ++xi) {
+                        auto coords = tensor(batch(bi), feature(fi), spatial(xi, yi, 0, 0));
+                        auto offset = mem->get_layout().get_linear_offset(coords);
+                        ptr[offset] = data[bi][fi][yi][xi];
+                    }
+                }
+            }
+        }
+    }
+
+    void fill_random(memory::ptr mem) {
+        auto dt = mem->get_layout().data_type;
+        switch (dt) {
+        case data_types::f32:
+            fill_random_typed<float>(mem, -127, 127, 2);
+            break;
+        case data_types::f16:
+            fill_random_typed<FLOAT16>(mem, -127, 127, 2);
+            break;
+        case data_types::i8:
+            fill_random_typed<int8_t>(mem, -127, 127, 1);
+            break;
+        case data_types::u8:
+            fill_random_typed<uint8_t>(mem, 0, 255, 1);
+            break;
+        default:
+            break;
+        }
+    }
+
+    template <typename T>
+    bool compare_outputs(const memory::ptr out_ref, const memory::ptr out_opt) {
+        auto output_lay = out_ref->get_layout();
+        auto opt_output_lay = out_opt->get_layout();
+
+        size_t b = output_lay.size.batch[0];
+        size_t f = output_lay.size.feature[0];
+        size_t x = output_lay.size.spatial[0];
+        size_t y = output_lay.size.spatial[1];
+        mem_lock<T> ref_ptr{out_ref, get_test_stream()};
+        mem_lock<T> opt_ptr{out_opt, get_test_stream()};
+        for (size_t bi = 0; bi < b; ++bi) {
+            for (size_t fi = 0; fi < f; ++fi) {
+                for (size_t yi = 0; yi < y; ++yi) {
+                    for (size_t xi = 0; xi < x; ++xi) {
+                        auto ref_out_coords = tensor(batch(bi), feature(fi), spatial(xi, yi, 0, 0));
+                        auto ref_out_offset = output_lay.get_linear_offset(ref_out_coords);
+                        auto ref_out_val = ref_ptr[ref_out_offset];
+
+                        auto opt_out_offset = opt_output_lay.get_linear_offset(ref_out_coords);
+                        auto opt_out_val = opt_ptr[opt_out_offset];
+
+                        EXPECT_EQ(opt_out_val, ref_out_val);
+                    }
+                }
+            }
+        }
+
+        return true;
+    }
+
+    void execute_compare(const quantize_random_test_params& params, bool check_result) {
+        auto& engine = get_test_engine();
+
+        auto in_layout = layout(params.input_type, params.in_format, params.input_size);
+        auto input = engine.allocate_memory(in_layout);
+        fill_random(input);
+
+        auto input_low   = engine.allocate_memory({ data_types::f32, format::bfyx, { 1, 1, 1, 1 } });
+        auto input_high  = engine.allocate_memory({ data_types::f32, format::bfyx, { 1, 1, 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 } });
+
+        // For quantize_gpu_scale_shift_opt
+        auto input_scale   = engine.allocate_memory({ data_types::f32, format::bfyx, { 1, 1, 1, 1 } });
+        auto input_shift   = engine.allocate_memory({ data_types::f32, format::bfyx, { 1, 1, 1, 1 } });
+        auto output_scale  = engine.allocate_memory({ data_types::f32, format::bfyx, { 1, 1, 1, 1 } });
+        auto output_shift  = engine.allocate_memory({ data_types::f32, format::bfyx, { 1, 1, 1, 1 } });
+
+        set_values(input_low,  { 0.0f });
+        set_values(input_high, { 40.0f });
+        set_values(output_low,  { 0.0f });
+        set_values(output_high, { 255.0f });
+
+        set_values(input_scale, { 2.0f });
+        set_values(input_shift, { 4.0f });
+        set_values(output_scale, { 2.0f });
+        set_values(output_shift, { 4.0f });
+
+        // Execute quantize_gpu_ref
+        cldnn::topology topo;
+        if (params.inputs_num == 5) {
+            topo.add(
+                input_layout("input", input->get_layout()),
+                data("input_low", input_low),
+                data("input_high", input_high),
+                data("output_low", output_low),
+                data("output_high", output_high),
+                quantize("quantize", "input", "input_low", "input_high", "output_low", "output_high", 256, params.output_type)
+            );
+        } else {
+            FAIL() << "Not supported inputs number: " << params.inputs_num;
+        }
+
+        auto build_ops = build_options();
+        build_ops.set_option(build_option::outputs({"quantize"}));
+
+        network net(engine, topo, build_ops);
+        net.set_input_data("input", input);
+
+        auto result = net.execute();
+        auto output = result.at("quantize").get_memory();
+
+        auto input_opt = engine.allocate_memory(in_layout);
+        if (params.input_type == data_types::f32) {
+            fill_typed<float>(input, input_opt);
+        } else if (params.input_type == data_types::f16) {
+            fill_typed<FLOAT16>(input, input_opt);
+        } else if (params.input_type == data_types::i8) {
+            fill_typed<int8_t>(input, input_opt);
+        } else if (params.input_type == data_types::u8) {
+            fill_typed<uint8_t>(input, input_opt);
+        } else {
+            FAIL() << "Not supported input data type: " << static_cast<size_t>(params.input_type);
+        }
+
+        cldnn::topology topo_opt;
+        if (params.inputs_num == 5) {
+            topo_opt.add(
+                input_layout("input_opt", input_opt->get_layout()),
+                reorder("input_re", "input_opt", format::bfyx, params.input_type),
+                data("input_low", input_low),
+                data("input_high", input_high),
+                data("output_low", output_low),
+                data("output_high", output_high),
+                quantize("quantize_opt", "input_re", "input_low", "input_high", "output_low", "output_high", 256, params.output_type),
+                reorder("out", "quantize_opt", params.out_format, params.output_type)
+            );
+        } else {
+            FAIL() << "Not supported inputs number: " << params.inputs_num;
+        }
+
+        auto buildops_opt = build_options();
+
+        network net_opt(engine, topo_opt, buildops_opt);
+        net_opt.set_input_data("input_opt", input_opt);
+
+        auto result_opt = net_opt.execute();
+        auto output_opt = result_opt.at("out").get_memory();
+
+        if (check_result == true) {
+            // Check data_types
+            if (params.output_type == data_types::f32) {
+                compare_outputs<float>(output, output_opt);
+            } else if (params.output_type == data_types::f16) {
+                compare_outputs<FLOAT16>(output, output_opt);
+            } else if (params.output_type == data_types::i8) {
+                compare_outputs<int8_t>(output, output_opt);
+            } else if (params.output_type == data_types::u8) {
+                compare_outputs<uint8_t>(output, output_opt);
+            } else {
+                FAIL() << "Not supported output data type: " << static_cast<size_t>(params.output_type);
+            }
+        }
+    }
+};
+
+struct quantize_random_test_param_generator : std::vector<quantize_random_test_params> {
+    quantize_random_test_param_generator& simple_params(data_types input_type, data_types output_type, format::type input_format, format::type output_format, int32_t inputs_num) {
+        push_back(quantize_random_test_params{ input_type, output_type, {1, 32, 2, 2}, input_format, output_format, inputs_num});
+        push_back(quantize_random_test_params{ input_type, output_type, {1, 16, 10, 10}, input_format, output_format, inputs_num});
+        return *this;
+    }
+};
+
+TEST_P(quantize_random_test, random) {
+    auto param = GetParam();
+    execute_compare(param, true);
+}
+
+INSTANTIATE_TEST_SUITE_P(quantize_smoke,
+                        quantize_random_test,
+                        testing::ValuesIn(
+                            quantize_random_test_param_generator()
+                            .simple_params(data_types::f32, data_types::u8, format::bs_fs_yx_bsv32_fsv32, format::bs_fs_yx_bsv32_fsv32, 5)
+                            .simple_params(data_types::f32, data_types::u8, format::b_fs_yx_fsv16, format::b_fs_yx_fsv16, 5)
+                        ));

inference-engine/thirdparty/clDNN/tests/test_cases/resample_gpu_test.cpp

@@ -2037,3 +2037,195 @@ TEST(resample_gpu, interpolate_in1x1x2x4_linear_scale) {
         EXPECT_TRUE(are_equal(answers[i], output_ptr[i])) << i;
     }
 }
+
+struct resample_opt_random_test_params {
+    data_types input_type;
+    tensor input_size;
+    tensor output_size;
+    uint32_t num_filter;
+    resample_type operation_type;
+    uint32_t align_corners;
+    format::type in_format;
+    format::type out_format;
+    std::vector<int32_t> pads_begin;
+    std::vector<int32_t> pads_end;
+};
+
+struct resample_opt_random_test : testing::TestWithParam<resample_opt_random_test_params>
+{
+    bool enable_profiling = false;
+
+    template <typename T>
+    void fill_random_typed(memory::ptr mem, int min, int max, int k) {
+        auto size = mem->get_layout().size;
+        size_t b = size.batch[0];
+        size_t f = size.feature[0];
+        size_t x = size.spatial[0];
+        size_t y = size.spatial[1];
+
+        auto data = generate_random_4d<T>(b, f, y, x, min, max, k);
+        mem_lock<T> ptr{mem, get_test_stream()};
+        for (size_t bi = 0; bi < b; ++bi) {
+            for (size_t fi = 0; fi < f; ++fi) {
+                for (size_t yi = 0; yi < y; ++yi) {
+                    for (size_t xi = 0; xi < x; ++xi) {
+                        auto coords = tensor(batch(bi), feature(fi), spatial(xi, yi, 0, 0));
+                        auto offset = mem->get_layout().get_linear_offset(coords);
+                        ptr[offset] = data[bi][fi][yi][xi];
+                    }
+                }
+            }
+        }
+    }
+
+    void fill_random(memory::ptr mem) {
+        auto dt = mem->get_layout().data_type;
+        switch (dt) {
+        case data_types::f32:
+            fill_random_typed<float>(mem, -127, 127, 2);
+            break;
+        case data_types::f16:
+            fill_random_typed<FLOAT16>(mem, -127, 127, 2);
+            break;
+        case data_types::i8:
+            fill_random_typed<int8_t>(mem, -127, 127, 1);
+            break;
+        case data_types::u8:
+            fill_random_typed<uint8_t>(mem, 0, 255, 1);
+            break;
+        default:
+            break;
+        }
+    }
+
+    template <typename T>
+    bool compare_outputs(const memory::ptr out_ref, const memory::ptr out_opt) {
+        auto output_lay = out_ref->get_layout();
+        auto opt_output_lay = out_opt->get_layout();
+
+        size_t b = output_lay.size.batch[0];
+        size_t f = output_lay.size.feature[0];
+        size_t x = output_lay.size.spatial[0];
+        size_t y = output_lay.size.spatial[1];
+        mem_lock<T> ref_ptr{out_ref, get_test_stream()};
+        mem_lock<T> opt_ptr{out_opt, get_test_stream()};
+        for (size_t bi = 0; bi < b; ++bi) {
+            for (size_t fi = 0; fi < f; ++fi) {
+                for (size_t yi = 0; yi < y; ++yi) {
+                    for (size_t xi = 0; xi < x; ++xi) {
+                        auto ref_out_coords = tensor(batch(bi), feature(fi), spatial(xi, yi, 0, 0));
+                        auto ref_out_offset = output_lay.get_linear_offset(ref_out_coords);
+                        auto ref_out_val = ref_ptr[ref_out_offset];
+
+                        auto opt_out_offset = opt_output_lay.get_linear_offset(ref_out_coords);
+                        auto opt_out_val = opt_ptr[opt_out_offset];
+
+                        EXPECT_EQ(ref_out_offset, opt_out_offset);
+                        if (std::is_same<T, FLOAT16>::value) {
+                            EXPECT_NEAR(static_cast<float>(opt_out_val), static_cast<float>(ref_out_val), 1.e-1f);
+                        } else {
+                            EXPECT_EQ(opt_out_val, ref_out_val);
+                        }
+                    }
+                }
+            }
+        }
+
+        return true;
+    }
+
+    void execute_compare(const resample_opt_random_test_params& params, bool check_result) {
+        auto& engine = get_test_engine();
+
+        auto in_layout = layout(params.input_type, format::bfyx, params.input_size);
+        auto in_mem = engine.allocate_memory(in_layout);
+        fill_random(in_mem);
+
+        /// bfyx
+        cldnn::topology topo;
+        topo.add(input_layout("in", in_layout));
+        auto prim = resample("resample", "in", params.output_size, params.num_filter, params.operation_type);
+        prim.align_corners = params.align_corners;
+        prim.pads_begin = params.pads_begin;
+        prim.pads_end = params.pads_end;
+        topo.add(prim);
+
+        auto build_opts = build_options();
+        build_opts.set_option(build_option::outputs({"resample"}));
+
+        network net(engine, topo, build_opts);
+        net.set_input_data("in", in_mem);
+
+        // first execution of ref
+        auto result = net.execute();
+        auto output = result.at("resample").get_memory();
+
+        cldnn::topology topo_opt;
+        topo_opt.add(input_layout("in", in_layout));
+        topo_opt.add(reorder("in_to_input_type", "in", params.in_format, params.input_type));
+        auto prim_opt = resample("resample_opt", "in_to_input_type", params.output_size, params.num_filter, params.operation_type);
+        prim_opt.align_corners = params.align_corners;
+        prim_opt.pads_begin = params.pads_begin;
+        prim_opt.pads_end = params.pads_end;
+        topo_opt.add(prim_opt);
+        topo_opt.add(reorder("res_to_bfyx", "resample_opt", format::bfyx, params.input_type));
+
+        auto build_opts_opt = build_options();
+        build_opts_opt.set_option(build_option::outputs({"resample_opt", "res_to_bfyx"}));
+
+        network net_opt(engine, topo_opt, build_opts_opt);
+
+        // Use in_mem from ref network
+        net_opt.set_input_data("in", in_mem);
+
+        // first execution of opt
+        auto result_opt = net_opt.execute();
+        auto output_opt = result_opt.at("res_to_bfyx").get_memory();
+
+        if (check_result == true) {
+            // Check data_types
+            if (params.input_type == data_types::f32) {
+                compare_outputs<float>(output, output_opt);
+            } else if (params.input_type == data_types::f16) {
+                compare_outputs<FLOAT16>(output, output_opt);
+            } else if (params.input_type == data_types::i8) {
+                compare_outputs<int8_t>(output, output_opt);
+            } else if (params.input_type == data_types::u8) {
+                compare_outputs<uint8_t>(output, output_opt);
+            } else {
+                FAIL() << "Not supported data type: " << static_cast<size_t>(params.input_type);
+            }
+        }
+    }
+};
+
+TEST_P(resample_opt_random_test, random) {
+    auto param = GetParam();
+    execute_compare(param, true);
+}
+
+INSTANTIATE_TEST_SUITE_P(resample_opt_smoke_nearest,
+                         resample_opt_random_test,
+                         testing::ValuesIn(
+                            std::vector<resample_opt_random_test_params>{
+                                { data_types::i8,  {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::nearest, 1, format::b_fs_yx_fsv16, format::b_fs_yx_fsv16, {}, {}},
+                                { data_types::i8,  {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::nearest, 1, format::b_fs_yx_fsv32, format::b_fs_yx_fsv32, {}, {}},
+                                { data_types::i8,  {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::nearest, 1, format::bs_fs_yx_bsv32_fsv16, format::bs_fs_yx_bsv32_fsv16, {}, {}},
+                                { data_types::i8,  {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::nearest, 1, format::bs_fs_yx_bsv32_fsv32, format::bs_fs_yx_bsv32_fsv32, {}, {}},
+                                { data_types::f16, {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::nearest, 1, format::b_fs_yx_fsv16, format::b_fs_yx_fsv16, {}, {}},
+                                { data_types::f16, {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::nearest, 1, format::b_fs_yx_fsv32, format::b_fs_yx_fsv32, {}, {}},
+                                { data_types::f16, {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::nearest, 1, format::bs_fs_yx_bsv32_fsv16, format::bs_fs_yx_bsv32_fsv16, {}, {}},
+                                { data_types::f16, {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::nearest, 1, format::bs_fs_yx_bsv32_fsv32, format::bs_fs_yx_bsv32_fsv32, {}, {}},
+                            }
+                        ));
+
+INSTANTIATE_TEST_SUITE_P(resample_opt_smoke_linear_onnx,
+                         resample_opt_random_test,
+                         testing::ValuesIn(
+                            std::vector<resample_opt_random_test_params>{
+                                { data_types::f16, {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::linear_onnx, 1, format::b_fs_yx_fsv16, format::b_fs_yx_fsv16, {}, {}},
+                                { data_types::f16, {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::linear_onnx, 1, format::b_fs_yx_fsv32, format::b_fs_yx_fsv32, {}, {}},
+                                { data_types::f16, {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::linear_onnx, 1, format::bs_fs_yx_bsv32_fsv16, format::bs_fs_yx_bsv32_fsv16, {}, {}},
+                                { data_types::f16, {1, 128, 13, 13},  {1, 128, 26, 26},  1, resample_type::linear_onnx, 1, format::bs_fs_yx_bsv32_fsv32, format::bs_fs_yx_bsv32_fsv32, {}, {}},
+                            }
+                        ));