From 61c97edd40d25e7149d97a702783be371c4381e9 Mon Sep 17 00:00:00 2001 From: Taylor Yeonbok Lee Date: Fri, 1 Oct 2021 15:18:15 +0900 Subject: [PATCH] [GPU] Reduce unused macros to reduce loading time (#7435) * Reduce unused macros, where two strategies are used: (1) Extract batch_headers and let them be included only once in each batch. (2) Static reduction in primitive_db_gen.py, which scans each macro's users and exclude the macro if there is no user. * Removed dependency from runtime to kernel_selector * Resolve too large string error * Fix duplicated definition (GET_FILTER_XXXX is defined in both fetch_weight.cl and by runtime. Removed from runtime because the definition is incorrect * Resolve GRN & deconv & gpu_select issues * Fix cldnn unittest issues * Minor fix * Applied review comments * Fix rebase error --- .../clDNN/kernel_selector/CMakeLists.txt | 7 +- .../core/cl_kernels/activation_opt.cl | 4 +- .../core/cl_kernels/activation_ref.cl | 4 +- .../core/cl_kernels/arg_max_min_axis.cl | 4 +- .../core/cl_kernels/arg_max_min_gpu_ref.cl | 4 +- .../core/cl_kernels/arg_max_min_opt.cl | 4 +- .../cl_kernels/average_unpooling_gpu_ref.cl | 4 +- .../core/cl_kernels/batch_to_space_ref.cl | 4 +- .../cl_kernels/binary_convolution_gpu_1x1.cl | 4 +- ...inary_convolution_gpu_1x1_b_fs_yx_fsv16.cl | 4 +- .../binary_convolution_gpu_generic.cl | 4 +- .../cl_kernels/binary_convolution_gpu_ref.cl | 4 +- .../core/cl_kernels/border_gpu_ref.cl | 4 +- .../core/cl_kernels/broadcast_gpu_ref.cl | 4 +- .../cl_kernels/concatenation_gpu_blocked.cl | 4 +- .../concatenation_gpu_depth_bfyx_no_pitch.cl | 4 +- .../concatenation_gpu_fs_b_yx_fsv32.cl | 4 +- .../core/cl_kernels/concatenation_gpu_ref.cl | 4 +- .../concatenation_gpu_simple_ref.cl | 4 +- .../convolution_gpu_b_fs_yx_fsv16_imad_1x1.cl | 6 +- .../convolution_gpu_b_fs_yx_fsv4_1x1.cl | 6 +- .../convolution_gpu_b_fs_yx_fsv4_dw.cl | 6 +- .../convolution_gpu_b_fs_yx_fsv4_int8.cl | 4 +- ...nvolution_gpu_b_fs_yx_fsv_16_32_imad_dw.cl | 6 +- .../convolution_gpu_b_fs_zyx_fsv16_imad.cl | 6 +- .../cl_kernels/convolution_gpu_bfyx_1x1.cl | 5 +- ...convolution_gpu_bfyx_1x1_hgemm_buf_16x1.cl | 4 +- .../convolution_gpu_bfyx_1x1_opt.cl | 4 +- ...volution_gpu_bfyx_depthwise_weights_lwg.cl | 4 +- .../convolution_gpu_bfyx_direct_10_12_16.cl | 6 +- .../cl_kernels/convolution_gpu_bfyx_f16.cl | 4 +- .../convolution_gpu_bfyx_f16_1x1.cl | 4 +- .../convolution_gpu_bfyx_f16_depthwise.cl | 4 +- .../convolution_gpu_bfyx_gemm_like_fp16.cl | 6 +- .../convolution_gpu_bfyx_gemm_like_fp32.cl | 6 +- .../cl_kernels/convolution_gpu_bfyx_iyxo.cl | 4 +- .../convolution_gpu_bfyx_os_iyx_osv16.cl | 4 +- .../convolution_gpu_bfyx_to_bfyx_f16.cl | 4 +- ...lution_gpu_bfyx_to_bs_fs_yx_bsv16_fsv16.cl | 4 +- .../convolution_gpu_bfyx_to_fs_byx_fsv32.cl | 4 +- .../convolution_gpu_fs_byx_fsv32.cl | 4 +- .../convolution_gpu_fs_byx_fsv32_1x1.cl | 4 +- .../convolution_gpu_fs_byx_fsv32_depthwise.cl | 4 +- ...ution_gpu_imad_bs_fs_yx_bsv16_fsv16_1x1.cl | 5 +- ...ution_gpu_imad_bs_fs_yx_bsv16_fsv16_3x3.cl | 5 +- .../convolution_gpu_mmad_b_fs_yx_fsv32.cl | 4 +- .../convolution_gpu_mmad_b_fs_yx_fsv32_dw.cl | 4 +- ...volution_gpu_mmad_bfyx_to_b_fs_yx_fsv32.cl | 4 +- ...nvolution_gpu_mmad_bfyx_to_b_fs_yx_fsv4.cl | 4 +- .../convolution_gpu_winograd_2x3_s1_fused.cl | 2 +- .../convolution_gpu_winograd_6x3_s1_fused.cl | 4 +- .../cl_kernels/convolution_gpu_yxfb_ref.cl | 4 +- .../convolution_gpu_yxfb_yxio_b16_fp16.cl | 4 +- .../convolution_gpu_yxfb_yxio_b16_fp32.cl | 4 +- ..._gpu_yxfb_yxio_b1_block_multiple_x_fp32.cl | 4 +- .../convolution_gpu_yxfb_yxio_b8_fp32.cl | 4 +- .../core/cl_kernels/ctc_greedy_decoder_ref.cl | 4 +- .../core/cl_kernels/cum_sum_partial_sum.cl | 4 +- .../core/cl_kernels/cum_sum_ref.cl | 4 +- .../deconvolution_gpu_b_fs_zyx_fsv16_dw.cl | 6 +- .../cl_kernels/deconvolution_gpu_bfyx_opt.cl | 4 +- ...deconvolution_gpu_imad_along_f_tile_bfx.cl | 6 +- .../deconvolution_gpu_imad_common.cl | 2 +- .../cl_kernels/deconvolution_gpu_imad_ref.cl | 6 +- .../core/cl_kernels/deconvolution_gpu_ref.cl | 4 +- .../deformable_convolution_gpu_bfyx_conv.cl | 4 +- .../deformable_convolution_gpu_bfyx_interp.cl | 4 +- .../deformable_convolution_gpu_bfyx_ref.cl | 6 +- .../cl_kernels/depth_to_space_block2_opt.cl | 4 +- .../core/cl_kernels/depth_to_space_ref.cl | 4 +- .../cl_kernels/detection_output_gpu_ref.cl | 4 +- .../core/cl_kernels/eltwise_b_fs_yx_fsv16.cl | 4 +- .../core/cl_kernels/eltwise_b_fs_yx_fsv4.cl | 4 +- .../core/cl_kernels/eltwise_fs_b_yx_fsv32.cl | 4 +- .../eltwise_mixed_byxf_and_fs_b_yx_fsv32.cl | 4 +- .../core/cl_kernels/eltwise_simple_vload8.cl | 4 +- .../core/cl_kernels/embedding_bag_ref.cl | 4 +- .../cl_kernels/extract_image_patches_ref.cl | 4 +- .../cl_kernels/fully_connected_gpu_MMAD.cl | 6 +- .../fully_connected_gpu_bf_io_gemm.cl | 4 +- ...fully_connected_gpu_bf_io_input_spatial.cl | 4 +- .../fully_connected_gpu_bf_io_ref.cl | 4 +- .../fully_connected_gpu_bf_tiled.cl | 5 +- .../fully_connected_gpu_bfyx_ref.cl | 6 +- ...ully_connected_gpu_bs_f_bsv16_af8_vload.cl | 4 +- .../fully_connected_gpu_bs_f_bsv16_b1.cl | 5 +- ...fully_connected_gpu_bs_f_bsv8_af8_vload.cl | 4 +- .../fully_connected_gpu_fb_io_b8_f8.cl | 4 +- .../fully_connected_gpu_fb_io_b8_f8_vload.cl | 5 +- .../fully_connected_gpu_fb_io_block_fp16.cl | 5 +- .../fully_connected_gpu_fb_io_ref.cl | 4 +- .../fully_connected_gpu_fb_oi_b8_fp32_ref.cl | 4 +- .../fully_connected_gpu_fb_oi_ref.cl | 4 +- .../fully_connected_gpu_fs_byx_fsv32.cl | 2 +- .../cl_kernels/fully_connected_gpu_imad.cl | 4 +- .../fully_connected_gpu_yxfb_ref.cl | 4 +- ...used_conv_eltwise_gpu_bfyx_1x1_opt_fp32.cl | 4 +- .../fused_conv_eltwise_gpu_bfyx_iyxo.cl | 5 +- ...used_conv_eltwise_gpu_bfyx_os_iyx_osv16.cl | 4 +- .../cl_kernels/fused_conv_eltwise_gpu_imad.cl | 2 +- .../cl_kernels/fused_conv_eltwise_gpu_ref.cl | 4 +- ...sed_conv_eltwise_gpu_yxfb_yxio_b16_fp16.cl | 4 +- .../core/cl_kernels/gather_elements_ref.cl | 5 +- .../core/cl_kernels/gather_nd_ref.cl | 2 +- .../core/cl_kernels/gather_ref.cl | 4 +- .../core/cl_kernels/gather_tree_gpu_ref.cl | 4 +- .../core/cl_kernels/gemm_mmad_int8.cl | 4 +- .../core/cl_kernels/gemm_mmad_int8_slm.cl | 4 +- .../core/cl_kernels/gemm_ref.cl | 2 +- .../core/cl_kernels/gemm_tiled_opt.cl | 2 +- .../cl_kernels/gen9_common_conv_bwd_data.cl | 4 +- .../gen9_common_conv_fwd_data_f16.cl | 4 +- .../gen9_common_conv_fwd_data_f32.cl | 4 +- .../core/cl_kernels/generic_eltwise_ref.cl | 4 +- .../core/cl_kernels/grn_ref.cl | 5 +- .../core/cl_kernels/include/acc_type.cl | 8 + .../include/{ => batch_headers}/common.cl | 0 .../include/{ => batch_headers}/data_types.cl | 59 +++---- .../include/{ => batch_headers}/fetch_data.cl | 61 +++---- .../{ => batch_headers}/fetch_weights.cl | 158 +++++++++--------- .../{ => batch_headers}/vec_typedefs.cl | 0 .../core/cl_kernels/include/image_data.cl | 9 + .../core/cl_kernels/include/unit_type.cl | 1 - ...rn_gpu_across_channel_multiple_features.cl | 4 +- ..._across_channel_multiple_features_fsv16.cl | 4 +- .../cl_kernels/lrn_gpu_across_channel_ref.cl | 4 +- .../lrn_gpu_across_channel_yxfb_b8_opt.cl | 4 +- .../core/cl_kernels/lrn_gpu_within_channel.cl | 4 +- .../cl_kernels/lrn_gpu_within_channel_opt.cl | 4 +- .../core/cl_kernels/lrn_ref.cl | 5 +- .../cl_kernels/lrn_within_channel_byxf_opt.cl | 5 +- .../cl_kernels/lstm_dynamic_input_bfyx_opt.cl | 4 +- .../core/cl_kernels/lstm_dynamic_input_ref.cl | 7 +- .../cl_kernels/lstm_dynamic_timeloop_ref.cl | 5 +- .../core/cl_kernels/lstm_elt_gpu_bfyx_ref.cl | 4 +- .../core/cl_kernels/lstm_gemm_gpu_bfyx_ref.cl | 5 +- ...tm_gemv_gpu_subgroup1x64_bfyx_ff_SIMD16.cl | 5 +- ...tm_gemv_gpu_subgroup1x64_bfyx_hh_SIMD16.cl | 4 +- .../core/cl_kernels/max_unpooling_gpu_ref.cl | 4 +- .../cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad.cl | 4 +- .../mvn_gpu_b_fs_yx_fsv16_imad_accumulate.cl | 2 +- .../mvn_gpu_b_fs_yx_fsv16_imad_reduce.cl | 2 +- .../core/cl_kernels/mvn_gpu_bfyx_opt.cl | 4 +- .../cl_kernels/mvn_gpu_ref_across_channels.cl | 4 +- .../cl_kernels/mvn_gpu_ref_within_channels.cl | 4 +- .../normalize_gpu_across_spatial_ref.cl | 4 +- .../normalize_gpu_within_spatial_ref.cl | 4 +- .../core/cl_kernels/one_hot_ref.cl | 4 +- .../core/cl_kernels/permute_ref.cl | 4 +- .../core/cl_kernels/permute_tile_8x8_4x4.cl | 4 +- .../cl_kernels/permute_tile_8x8_4x4_fsv.cl | 4 +- .../cl_kernels/pooling_gpu_average_opt.cl | 4 +- .../cl_kernels/pooling_gpu_b_fs_yx_fsv4.cl | 8 +- .../pooling_gpu_b_fs_zyx_fsv16_imad.cl | 8 +- .../cl_kernels/pooling_gpu_bfyx_block_opt.cl | 4 +- .../core/cl_kernels/pooling_gpu_blocked.cl | 6 +- .../pooling_gpu_bs_fs_yx_bsv16_fsv16.cl | 6 +- .../cl_kernels/pooling_gpu_bsv16_fsv16.cl | 4 +- .../core/cl_kernels/pooling_gpu_byxf_opt.cl | 4 +- .../pooling_gpu_byxf_padding_opt.cl | 4 +- .../cl_kernels/pooling_gpu_fs_b_yx_fsv32.cl | 6 +- .../core/cl_kernels/pooling_gpu_int8_ref.cl | 4 +- .../core/cl_kernels/pooling_gpu_ref.cl | 4 +- .../cl_kernels/pyramid_roi_align_gpu_ref.cl | 4 +- .../core/cl_kernels/quantize_gpu_ref.cl | 4 +- .../quantize_gpu_scale_shift_opt.cl | 4 +- .../cl_kernels/reduce_gpu_b_fs_yx_fsv16.cl | 4 +- .../core/cl_kernels/reduce_ref.cl | 4 +- .../core/cl_kernels/region_yolo_gpu_ref.cl | 2 +- .../core/cl_kernels/reorder_biplanar_nv12.cl | 4 +- .../core/cl_kernels/reorder_data.cl | 5 +- ...eorder_data_b_fs_yx_fsv16_fsv32_to_bfyx.cl | 4 +- .../reorder_data_bfyx_to_blocked_format.cl | 4 +- .../core/cl_kernels/reorder_data_binary.cl | 4 +- .../core/cl_kernels/reorder_data_fast_b1.cl | 4 +- .../reorder_data_to_yxfb_batched.cl | 4 +- .../reorder_from_winograd_2x3_s1.cl | 4 +- .../reorder_fs_b_yx_fsv32_to_bfyx.cl | 2 +- .../core/cl_kernels/reorder_weights.cl | 5 +- .../core/cl_kernels/reorder_weights_binary.cl | 2 +- .../reorder_weights_image_2d_c4_fyx_b.cl | 5 +- .../reorder_weights_image_winograd_6x3_s1.cl | 4 +- .../core/cl_kernels/reorder_weights_opt.cl | 4 +- .../reorder_weights_winograd_2x3_s1.cl | 4 +- .../reorder_weights_winograd_6x3_s1.cl | 4 +- .../core/cl_kernels/reorg_yolo_gpu_ref.cl | 4 +- .../core/cl_kernels/resample_opt.cl | 4 +- .../core/cl_kernels/resample_ref.cl | 4 +- .../core/cl_kernels/reshape_ref.cl | 4 +- .../core/cl_kernels/reverse_sequence_ref.cl | 4 +- .../core/cl_kernels/roi_pooling_ps_ref.cl | 4 +- .../core/cl_kernels/roi_pooling_ref.cl | 4 +- .../cl_kernels/scatter_elements_update_ref.cl | 4 +- .../core/cl_kernels/scatter_nd_update_ref.cl | 4 +- .../core/cl_kernels/scatter_update_ref.cl | 4 +- .../core/cl_kernels/select_gpu_ref.cl | 4 +- .../core/cl_kernels/shuffle_channels_ref.cl | 4 +- .../core/cl_kernels/softmax_gpu_bf.cl | 4 +- .../core/cl_kernels/softmax_gpu_fb.cl | 4 +- .../softmax_gpu_items_class_optimized.cl | 4 +- .../core/cl_kernels/softmax_gpu_ref.cl | 4 +- .../core/cl_kernels/space_to_batch_ref.cl | 4 +- .../core/cl_kernels/space_to_depth_ref.cl | 4 +- .../core/cl_kernels/strided_slice_ref.cl | 4 +- .../core/cl_kernels/tile_ref.cl | 4 +- .../kernel_selector/core/common/jitter.cpp | 2 + .../core/common/primitive_db.cpp | 4 +- .../core/common/primitive_db.h | 2 + .../core/common/primitive_db_gen.py | 156 +++++++++++++++-- .../clDNN/runtime/kernels_cache.cpp | 16 +- .../clDNN/runtime/kernels_cache.hpp | 25 ++- .../thirdparty/clDNN/src/program.cpp | 2 + 212 files changed, 740 insertions(+), 593 deletions(-) create mode 100644 inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/acc_type.cl rename inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/{ => batch_headers}/common.cl (100%) rename inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/{ => batch_headers}/data_types.cl (86%) rename inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/{ => batch_headers}/fetch_data.cl (93%) rename inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/{ => batch_headers}/fetch_weights.cl (89%) rename inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/{ => batch_headers}/vec_typedefs.cl (100%) create mode 100644 inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/image_data.cl diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/CMakeLists.txt b/inference-engine/thirdparty/clDNN/kernel_selector/CMakeLists.txt index a4583766216..4ff8f0b545a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/CMakeLists.txt +++ b/inference-engine/thirdparty/clDNN/kernel_selector/CMakeLists.txt @@ -74,12 +74,14 @@ file(GLOB_RECURSE __CLDNN_Sources__cl_kernels set(__CLDNN_Directory__cg_cache "${CLDNN__CODEGEN_INCDIR}") set(__CLDNN_CGDirectory__cg_cache "${CLDNN__CODEGEN_DIR}/cache") set(__CLDNN_Label__cg_cache "${__CLDNN_Label__core}\\codegen") +set(__CLDNN_Label__cg_cache_batch_headers "${__CLDNN_Label__core}\\codegen") set(__CLDNN_File__cg_cache__prim_db "ks_primitive_db.inc") +set(__CLDNN_File__cg_cache__prim_db_batch_headers "ks_primitive_db_batch_headers.inc") set(__CLDNN_Sources__cg_cache "${__CLDNN_Directory__cg_cache}/${__CLDNN_File__cg_cache__prim_db}" + "${__CLDNN_Directory__cg_cache}/${__CLDNN_File__cg_cache__prim_db_batch_headers}" ) - set(__CLDNN_AllSources ${__CLDNN_Sources__main} ${__CLDNN_Sources__core} @@ -144,12 +146,13 @@ endif() # =================================== Custom pre- and post-steps ======================================= add_custom_command(OUTPUT "${__CLDNN_CGDirectory__cg_cache}/${__CLDNN_File__cg_cache__prim_db}" COMMAND "${CMAKE_COMMAND}" -E make_directory "${__CLDNN_CGDirectory__cg_cache}" - COMMAND "${PYTHON_EXECUTABLE}" "${__CLDNN_Directory__core_common}/primitive_db_gen.py" -out_path "${__CLDNN_CGDirectory__cg_cache}" -out_file_name "${__CLDNN_File__cg_cache__prim_db}" -kernels "${__CLDNN_Directory__cl_kernels}" + COMMAND "${PYTHON_EXECUTABLE}" "${__CLDNN_Directory__core_common}/primitive_db_gen.py" -out_path "${__CLDNN_CGDirectory__cg_cache}" -out_file_name_prim_db "${__CLDNN_File__cg_cache__prim_db}" -out_file_name_batch_headers "${__CLDNN_File__cg_cache__prim_db_batch_headers}" -kernels "${__CLDNN_Directory__cl_kernels}" DEPENDS ${__CLDNN_Sources__cl_kernels} "${__CLDNN_Directory__core_common}/primitive_db_gen.py" COMMENT "Generating ${__CLDNN_File__cg_cache__prim_db} ..." ) add_custom_command(OUTPUT "${__CLDNN_Directory__cg_cache}/${__CLDNN_File__cg_cache__prim_db}" COMMAND "${CMAKE_COMMAND}" -E copy_if_different "${__CLDNN_CGDirectory__cg_cache}/${__CLDNN_File__cg_cache__prim_db}" "${__CLDNN_Directory__cg_cache}/${__CLDNN_File__cg_cache__prim_db}" + COMMAND "${CMAKE_COMMAND}" -E copy_if_different "${__CLDNN_CGDirectory__cg_cache}/${__CLDNN_File__cg_cache__prim_db_batch_headers}" "${__CLDNN_Directory__cg_cache}/${__CLDNN_File__cg_cache__prim_db_batch_headers}" DEPENDS "${__CLDNN_CGDirectory__cg_cache}/${__CLDNN_File__cg_cache__prim_db}" ${__CLDNN_Sources__cl_kernels} "${__CLDNN_Directory__core_common}/primitive_db_gen.py" COMMENT "Updating file if the file changed (${__CLDNN_File__cg_cache__prim_db}) ..." ) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/activation_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/activation_opt.cl index b80f5044e5c..35b7936a42b 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/activation_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/activation_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" KERNEL(activation)( __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/activation_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/activation_ref.cl index e4ef903a9bb..0aacaaa07e1 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/activation_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/activation_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #ifdef PARAMETERIZED #define GET_INDEX(prefix, num, idx_order) CAT(CAT(prefix, num), _GET_INDEX_SAFE)(idx_order) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_axis.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_axis.cl index 8dbf7a029b5..68d47ff8989 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_axis.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_axis.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #ifdef BATCH_AXIS #define VALUES_NUM INPUT0_BATCH_NUM diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_gpu_ref.cl index 551be7f23c9..b2d09d502fe 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" #define GLOBAL_SIZE 128 #define LOCAL_SIZE GLOBAL_SIZE diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_opt.cl index 428bd11ebdd..c61c3ac2fed 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" #ifndef SG_SIZE #define SG_SIZE 16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/average_unpooling_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/average_unpooling_gpu_ref.cl index eddb12fba4e..4c4dee16944 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/average_unpooling_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/average_unpooling_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(average_unpooling_gpu)(const __global UNIT_TYPE* input, __global UNIT_TYPE* output) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/batch_to_space_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/batch_to_space_ref.cl index 4a64c093813..10b373a320a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/batch_to_space_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/batch_to_space_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(batch_to_space_ref)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_1x1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_1x1.cl index 499424fa0d6..6f46f6f83e5 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_1x1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_1x1.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define OC_BLOCK_SIZE 32 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_1x1_b_fs_yx_fsv16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_1x1_b_fs_yx_fsv16.cl index 8fc1e8af12b..7b3744b9735 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_1x1_b_fs_yx_fsv16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_1x1_b_fs_yx_fsv16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/unit_type.cl" #define OC_BLOCK_SIZE 16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_generic.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_generic.cl index 757f2c2d55e..093466e5f1b 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_generic.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_generic.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define OC_BLOCK_SIZE 32 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_ref.cl index 479687e9384..2e41db6e878 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(binary_convolution_ref)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/border_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/border_gpu_ref.cl index 3ceb04000ab..38234b49b5f 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/border_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/border_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(border_gpu_ref)( diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/broadcast_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/broadcast_gpu_ref.cl index 173d264ba90..b222c5bece2 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/broadcast_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/broadcast_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(broadcast_gpu_ref)( diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_blocked.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_blocked.cl index e440a4933d5..54375aaa5f2 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_blocked.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_blocked.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #define WORK_GROUP_SIZE 16 #define IC_BLOCK 16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_depth_bfyx_no_pitch.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_depth_bfyx_no_pitch.cl index 313894dab67..2be9c5e16a0 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_depth_bfyx_no_pitch.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_depth_bfyx_no_pitch.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" // // In this kernel we are processing "fyx" as flatten 1D "elements". diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_fs_b_yx_fsv32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_fs_b_yx_fsv32.cl index 8c721e5c849..c6651f6f0d0 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_fs_b_yx_fsv32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_fs_b_yx_fsv32.cl @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" #include "include/unit_type.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_ref.cl index 3da2fab0e29..a7ec0edbdb4 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define GET_INDEX(prefix, ORDER) CAT(prefix, _GET_INDEX)(ORDER) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_simple_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_simple_ref.cl index e492b19fe06..6689d28c4ee 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_simple_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_simple_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" ///////////////////////// Input Index ///////////////////////// inline uint FUNC(get_input_index)(uint b, uint f, uint w, uint z, uint y, uint x) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv16_imad_1x1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv16_imad_1x1.cl index 4b3ad7e51b9..250c98cdf01 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv16_imad_1x1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv16_imad_1x1.cl @@ -2,10 +2,10 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" #include "include/imad.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" #define TYPE_N_(type, n) type##n #define TYPE_N(type, n) TYPE_N_(type, n) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_1x1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_1x1.cl index e6b08df0ac2..541c24c084a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_1x1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_1x1.cl @@ -2,10 +2,10 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" #include "include/imad.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" // ====================================================================================== // Host side jit-constants: diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_dw.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_dw.cl index 239e28ac4c6..9a0a762fd54 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_dw.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_dw.cl @@ -3,9 +3,9 @@ // #include "include/imad.cl" -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" // ====================================================================================== // Host side jit-constants: diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_int8.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_int8.cl index 2a825d8ee8d..75505fe234f 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_int8.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_int8.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" #include "include/imad.cl" #define INPUT0_PACKED_TYPE uint diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv_16_32_imad_dw.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv_16_32_imad_dw.cl index bcf1fc4379f..42eb8c191dd 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv_16_32_imad_dw.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv_16_32_imad_dw.cl @@ -3,9 +3,9 @@ // #include "include/imad.cl" -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" // ====================================================================================== // Host side jit-constants: diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_zyx_fsv16_imad.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_zyx_fsv16_imad.cl index a25db2b63ba..fb272dfcfbe 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_zyx_fsv16_imad.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_zyx_fsv16_imad.cl @@ -2,10 +2,10 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" #include "include/imad.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" #define TYPE_N_(type, n) type##n #define TYPE_N(type, n) TYPE_N_(type, n) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1.cl index b704ae24bba..49b8ef4fce0 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/sub_group.cl" #if FP16_UNIT_USED @@ -128,4 +128,3 @@ KERNEL(convolution_bfyx_1x1)( #undef CONCAT_TOKEN #undef CONCAT_TOKEN_HANDLER1 #undef MULTIPLY_BLOCKS_16x16 -#undef MAKE_VECTOR_TYPE diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1_hgemm_buf_16x1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1_hgemm_buf_16x1.cl index cb5eab11507..8cbe88d9d8c 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1_hgemm_buf_16x1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1_hgemm_buf_16x1.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/gemm_common.cl" #define MULT(C_, A_, i_) \ diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1_opt.cl index f633ecbafee..7843daffff5 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define SIMD_SIZE 8 __attribute__((intel_reqd_sub_group_size(SIMD_SIZE))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_depthwise_weights_lwg.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_depthwise_weights_lwg.cl index 5b99c0b6274..cfaa1677f20 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_depthwise_weights_lwg.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_depthwise_weights_lwg.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #if FP16_UNIT_USED #define ALIGNED_BLOCK_READ(ptr, byte_offset) as_half(intel_sub_group_block_read_us8((const __global ushort*)(ptr) + (byte_offset))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_direct_10_12_16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_direct_10_12_16.cl index e045323e745..7505b615437 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_direct_10_12_16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_direct_10_12_16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" ////////////////////////////////////////////////////////////////////////////// // Direct Convolution @@ -26,7 +26,7 @@ KERNEL(convolution_f16_10x12x16)( #endif uint split_idx) { -#include "include/vec_typedefs.cl" +#include "include/batch_headers/vec_typedefs.cl" const unsigned global_x = (uint)get_global_id(0); const unsigned global_y = (uint)get_global_id(1); diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16.cl index 212f4bc04dc..223fa7f5a0a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define INPUT_TYPE INPUT0_TYPE #define INPUT_TYPE2 MAKE_VECTOR_TYPE(INPUT0_TYPE, 2) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16_1x1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16_1x1.cl index cc4ba6fa81b..5d47b2d63a0 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16_1x1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16_1x1.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/unit_type.cl" #if X_BLOCK_SIZE > 1 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16_depthwise.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16_depthwise.cl index ad17b16d3ff..c98ce3be591 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16_depthwise.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16_depthwise.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_gemm_like_fp16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_gemm_like_fp16.cl index 7fd07490ae1..a166b7f2c64 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_gemm_like_fp16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_gemm_like_fp16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #if defined(cl_intel_subgroups_short) #define TILE_M 1 @@ -20,7 +20,7 @@ KERNEL(convolution_f16)( #endif uint split_idx) { -#include "include/vec_typedefs.cl" +#include "include/batch_headers/vec_typedefs.cl" const unsigned group_x = get_group_id(0); const unsigned group_y = get_group_id(1); diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_gemm_like_fp32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_gemm_like_fp32.cl index 23c956d06ad..558cee90ac2 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_gemm_like_fp32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_gemm_like_fp32.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/sub_group.cl" #define TILE_M 2 @@ -20,7 +20,7 @@ KERNEL(convolution_f32)( #endif uint split_idx) { -#include "include/vec_typedefs.cl" +#include "include/batch_headers/vec_typedefs.cl" const unsigned group_x = get_group_id(0); const unsigned group_y = get_group_id(1); diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_iyxo.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_iyxo.cl index b5489893da7..20102117278 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_iyxo.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_iyxo.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" __attribute__((intel_reqd_sub_group_size(SUB_GROUP_SIZE))) __attribute__((reqd_work_group_size(1, 1, SUB_GROUP_SIZE))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_os_iyx_osv16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_os_iyx_osv16.cl index b622f1c1698..6a3bfa15c4f 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_os_iyx_osv16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_os_iyx_osv16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" // --------------------------------------------------------------------------------------------------------------------- diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_bfyx_f16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_bfyx_f16.cl index d3c8b435118..e13894f8af9 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_bfyx_f16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_bfyx_f16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define FEATURE_SLICE_SIZE 16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_bs_fs_yx_bsv16_fsv16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_bs_fs_yx_bsv16_fsv16.cl index 250660513b3..4268b89933c 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_bs_fs_yx_bsv16_fsv16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_bs_fs_yx_bsv16_fsv16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/unit_type.cl" #define BATCH_SLICE_SIZE 16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_fs_byx_fsv32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_fs_byx_fsv32.cl index 4eb31daf594..890e1b34472 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_fs_byx_fsv32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_fs_byx_fsv32.cl @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" #include "include/unit_type.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32.cl index 6c72dfb3c15..092049d15df 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32_1x1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32_1x1.cl index 725441511c7..252a3295a85 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32_1x1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32_1x1.cl @@ -3,8 +3,8 @@ // #include "include/unit_type.cl" -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32_depthwise.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32_depthwise.cl index 15476c0fc50..bf0f5b4f917 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32_depthwise.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32_depthwise.cl @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" #include "include/unit_type.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_imad_bs_fs_yx_bsv16_fsv16_1x1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_imad_bs_fs_yx_bsv16_fsv16_1x1.cl index acdfa0983e7..1efb1d188b2 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_imad_bs_fs_yx_bsv16_fsv16_1x1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_imad_bs_fs_yx_bsv16_fsv16_1x1.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" #include "include/imad.cl" #if QUANTIZATION_TERM #define ACCUMULATOR_TYPE int @@ -20,7 +20,6 @@ #define AS_TYPE_N_(type, n, x) as_##type##n(x) #define AS_TYPE_N(type, n, x) AS_TYPE_N_(type, n, x) #define AS_INPUT0_TYPE_4(x) AS_TYPE_N(INPUT0_TYPE, 4, x) -#define MAKE_VECTOR_TYPE(elem_type, size) CAT(elem_type, size) #define OUTPUT_TYPE16 MAKE_VECTOR_TYPE(OUTPUT_TYPE, 16) #define BATCH_SLICE_SIZE 16 #define FEATURE_SLICE_SIZE 16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_imad_bs_fs_yx_bsv16_fsv16_3x3.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_imad_bs_fs_yx_bsv16_fsv16_3x3.cl index e4af8e1799d..caf0cf5516b 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_imad_bs_fs_yx_bsv16_fsv16_3x3.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_imad_bs_fs_yx_bsv16_fsv16_3x3.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" #include "include/imad.cl" #if QUANTIZATION_TERM #define ACCUMULATOR_TYPE int @@ -16,7 +16,6 @@ #define ACTIVATION_TYPE INPUT0_TYPE #define TO_ACTIVATION_TYPE(x) TO_INPUT0_TYPE(x) #endif -#define MAKE_VECTOR_TYPE(elem_type, size) CAT(elem_type, size) #define OUTPUT_TYPE16 MAKE_VECTOR_TYPE(OUTPUT_TYPE, 16) #define BATCH_SLICE_SIZE 16 #define FEATURE_SLICE_SIZE 16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_b_fs_yx_fsv32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_b_fs_yx_fsv32.cl index bb24fbb79f6..b052183b5f8 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_b_fs_yx_fsv32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_b_fs_yx_fsv32.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/mmad.cl" #define AS_TYPE(type, val) CAT(as_, type)(val) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_b_fs_yx_fsv32_dw.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_b_fs_yx_fsv32_dw.cl index 01045cac5ab..6969d2f2d8a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_b_fs_yx_fsv32_dw.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_b_fs_yx_fsv32_dw.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(convolution_mmad_b_fs_yx_fsv32_dw)( __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_bfyx_to_b_fs_yx_fsv32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_bfyx_to_b_fs_yx_fsv32.cl index b03dd28a5f7..6a3f7de816e 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_bfyx_to_b_fs_yx_fsv32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_bfyx_to_b_fs_yx_fsv32.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/imad.cl" #define CEIL_DIV(x, y) (1 + ((x) - 1) / (y)) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_bfyx_to_b_fs_yx_fsv4.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_bfyx_to_b_fs_yx_fsv4.cl index 1115cd60cc8..9aca517a91f 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_bfyx_to_b_fs_yx_fsv4.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_bfyx_to_b_fs_yx_fsv4.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/imad.cl" #define CEIL_DIV(x, y) (1 + ((x) - 1) / (y)) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_winograd_2x3_s1_fused.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_winograd_2x3_s1_fused.cl index cbcc1d61cf6..d2826bb6361 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_winograd_2x3_s1_fused.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_winograd_2x3_s1_fused.cl @@ -8,7 +8,7 @@ // Output matrix dimensions: M x N // -------------------------------------------------------------------------------------------------------------------------------- -#include "include/common.cl" +#include "include/batch_headers/common.cl" #define DOT4i0( _result, _A, _B, i) \ diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_winograd_6x3_s1_fused.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_winograd_6x3_s1_fused.cl index bfabed08d03..4450e9ce2ea 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_winograd_6x3_s1_fused.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_winograd_6x3_s1_fused.cl @@ -8,8 +8,8 @@ // Output matrix dimensions: M x N // -------------------------------------------------------------------------------------------------------------------------------- -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" #define DOT8i_0( _result, _A, _B, i) \ diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_ref.cl index 826865ec92f..3407ab36ebd 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(convolution_gpu_yxfb_ref)( const __global UNIT_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b16_fp16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b16_fp16.cl index 2d0cc481bcc..3e14f21a910 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b16_fp16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b16_fp16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/sub_group.cl" __attribute__((intel_reqd_sub_group_size(16))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b16_fp32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b16_fp32.cl index 9da54043c61..e335fd3246e 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b16_fp32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b16_fp32.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/sub_group.cl" KERNEL(convolution_gpu_yxfb_yxio_b16)( diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b1_block_multiple_x_fp32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b1_block_multiple_x_fp32.cl index 610813b9629..3e1c7c11b81 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b1_block_multiple_x_fp32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b1_block_multiple_x_fp32.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/sub_group.cl" __attribute__((reqd_work_group_size(LOCAL_WORK_GROUP_SIZE, 1, 1))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b8_fp32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b8_fp32.cl index 1e54c1c3f9f..0b5945c4044 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b8_fp32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b8_fp32.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/sub_group.cl" __attribute__((reqd_work_group_size(LOCAL_WORK_GROUP_SIZE, 1, 1))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/ctc_greedy_decoder_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/ctc_greedy_decoder_ref.cl index dab9a3496ec..7ff50a3160a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/ctc_greedy_decoder_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/ctc_greedy_decoder_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(ctc_greedy_decoder_ref)(const __global INPUT0_TYPE* probabilities ,const __global INPUT1_TYPE* sequence_indicators diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/cum_sum_partial_sum.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/cum_sum_partial_sum.cl index 0184c7b0ea4..0222444f841 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/cum_sum_partial_sum.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/cum_sum_partial_sum.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" ///////////////////////// Input Index ///////////////////////// inline uint FUNC(get_input_index)(uint b, uint f, uint w, uint z, uint y, uint x) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/cum_sum_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/cum_sum_ref.cl index 87e5a44a691..f82d07ddae5 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/cum_sum_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/cum_sum_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" ///////////////////////// Input Index ///////////////////////// inline uint FUNC(get_input_index)(uint b, uint f, uint w, uint z, uint y, uint x) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_b_fs_zyx_fsv16_dw.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_b_fs_zyx_fsv16_dw.cl index 2ef0bbca9e6..417e7f5d909 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_b_fs_zyx_fsv16_dw.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_b_fs_zyx_fsv16_dw.cl @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #include "deconvolution_gpu_imad_common.cl" diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_bfyx_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_bfyx_opt.cl index 64a8ae7a8d1..d4252c8a1c3 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_bfyx_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_bfyx_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define WORK_GROUP_GROUP_SIZE 16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_along_f_tile_bfx.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_along_f_tile_bfx.cl index 69d2453722f..7c5d21e2090 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_along_f_tile_bfx.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_along_f_tile_bfx.cl @@ -2,10 +2,10 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" #include "include/imad.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" #include "deconvolution_gpu_imad_common.cl" diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_common.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_common.cl index 52826aca3f6..651b35d7a63 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_common.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_common.cl @@ -2,7 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" #define CEIL_DIV(a, b) (((a) + ((b) - 1)) / (b)) #define ALIGN(a, b) (CEIL_DIV(a, b) * (b)) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_ref.cl index 78a94855c22..e276700e73e 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_ref.cl @@ -2,10 +2,10 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" #include "include/imad.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" #include "deconvolution_gpu_imad_common.cl" diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_ref.cl index eff0c831507..0bcf450041a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(deconvolution_gpu_yxfb_ref)( const __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_conv.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_conv.cl index 90f3e140ed8..e77400341f7 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_conv.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_conv.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/unit_type.cl" #define FEATURE_SLICE_SIZE 16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_interp.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_interp.cl index 7469dbbe296..b3c57029e16 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_interp.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_interp.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" __attribute__((intel_reqd_sub_group_size(16))) KERNEL(deformable_convolution_gpu_bfyx_interp)( diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_ref.cl index 55f7e0c7d95..9edea25b7eb 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_ref.cl @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" KERNEL(deformable_convolution_gpu_bfyx_ref)( const __global INPUT0_TYPE* data, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/depth_to_space_block2_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/depth_to_space_block2_opt.cl index 985711b4aab..1748fe2eac8 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/depth_to_space_block2_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/depth_to_space_block2_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(depth_to_space_block2_opt)(const __global half* input, __global half* output) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/depth_to_space_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/depth_to_space_ref.cl index c65f66ac1b7..df803cd33cb 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/depth_to_space_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/depth_to_space_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(depth_to_space_ref)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/detection_output_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/detection_output_gpu_ref.cl index 5d9ae37def8..1341f975b80 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/detection_output_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/detection_output_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/common.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/common.cl" #include "include/detection_output_common.cl" // DetectionOuput - performs non-maximuim suppression to generate the detection output diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_b_fs_yx_fsv16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_b_fs_yx_fsv16.cl index a165f48424b..cb937024438 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_b_fs_yx_fsv16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_b_fs_yx_fsv16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define FEATURE_SLICE_SIZE 16 #define unroll_for __attribute__((opencl_unroll_hint())) for diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_b_fs_yx_fsv4.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_b_fs_yx_fsv4.cl index 6c7f755f972..9c0fb1e353d 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_b_fs_yx_fsv4.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_b_fs_yx_fsv4.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define OUTPUT_TYPE_BLOCK MAKE_VECTOR_TYPE(OUTPUT_TYPE, VEC_SIZE) #define TO_TYPE(type, val) CAT(convert_, type)(val) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_fs_b_yx_fsv32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_fs_b_yx_fsv32.cl index eb35b427265..bc467525963 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_fs_b_yx_fsv32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_fs_b_yx_fsv32.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(eltwise_fs_b_yx_fsv32)( INPUTS_DECLS diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_mixed_byxf_and_fs_b_yx_fsv32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_mixed_byxf_and_fs_b_yx_fsv32.cl index 7bd4f470b0f..2a81fc37b1e 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_mixed_byxf_and_fs_b_yx_fsv32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_mixed_byxf_and_fs_b_yx_fsv32.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/unit_type.cl" // Kernel works only for sub_group size of 16 with 32 features slice size and process 2 features per WI diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_simple_vload8.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_simple_vload8.cl index 135997e61a7..148aa7ce8b2 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_simple_vload8.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_simple_vload8.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(eltwise_gpu_vload8)(INPUTS_DECLS __global OUTPUT_TYPE* output) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/embedding_bag_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/embedding_bag_ref.cl index e4e9c298f32..7fc22c741c8 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/embedding_bag_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/embedding_bag_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #ifdef PACKED_SUM KERNEL(embedding_bag_ref)( diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/extract_image_patches_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/extract_image_patches_ref.cl index 6f7317702ab..478183a2526 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/extract_image_patches_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/extract_image_patches_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(extract_image_patches_ref)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_MMAD.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_MMAD.cl index f47c2de1f0e..b2f17aa4594 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_MMAD.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_MMAD.cl @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" #include "include/mmad.cl" #define INPUT_PACKED_TYPE_8 CAT(INPUT_PACKED_TYPE, 8) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_gemm.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_gemm.cl index 6d97e4fe835..3b22c1b57fa 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_gemm.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_gemm.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #if defined(__fc_f16) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_input_spatial.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_input_spatial.cl index ba3f9b587ef..601d384a71a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_input_spatial.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_input_spatial.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" // Required JIT constants: // - FP16_SUPPORTED - [0/1] Value indicating whether device supports FP16 OpenCL extension (cl_khr_fp16). diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_ref.cl index 78742821f1e..8b299a97e1b 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" // Required JIT constants: // - FP16_SUPPORTED - [0/1] Value indicating whether device supports FP16 OpenCL extension (cl_khr_fp16). diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_tiled.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_tiled.cl index 5c0387115c4..4aceb2fd74b 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_tiled.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_tiled.cl @@ -2,9 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" // JIT Parameters: // SIMD - sub-group size/simd width, one of {8, 16}; diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bfyx_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bfyx_ref.cl index 9aac3fcc9a7..05ab582cde8 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bfyx_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bfyx_ref.cl @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" KERNEL(fc)( const __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv16_af8_vload.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv16_af8_vload.cl index 28326ed3b3c..ace92cd3545 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv16_af8_vload.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv16_af8_vload.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/sub_group.cl" // Block read - currently block is 4 bytes aligned. diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv16_b1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv16_b1.cl index d413761f9be..4da22a4de19 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv16_b1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv16_b1.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" // --------------------------------------------------------------------------------------------------------------------- // Just-in-time macro definitions: @@ -486,7 +486,6 @@ KERNEL (fully_connected_gpu_bx_bs_x_bsv16_b1)( #undef CONCAT_TOKEN_HANDLER1 #undef CONCAT_TOKEN -#undef MAKE_VECTOR_TYPE #undef CVT_UNIT #undef CHUNK_UNITS_TYPE #undef AS_CHUNK diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv8_af8_vload.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv8_af8_vload.cl index 5ec70585cee..1375935454a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv8_af8_vload.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv8_af8_vload.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/sub_group.cl" #if FP16_UNIT_USED diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_b8_f8.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_b8_f8.cl index 7cf440ad42e..3e9094759da 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_b8_f8.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_b8_f8.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/sub_group.cl" __attribute__((reqd_work_group_size(8, 1, 1))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_b8_f8_vload.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_b8_f8_vload.cl index ad37c8c6a7c..aac7c1268af 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_b8_f8_vload.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_b8_f8_vload.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/sub_group.cl" #if FP16_UNIT_USED @@ -224,7 +224,6 @@ KERNEL (fully_connected_gpu_xb_xb_b8_x8_vload)( #undef SUB_GROUP_SIZE #undef ALIGNED_BLOCK_READ8 -#undef MAKE_VECTOR_TYPE #undef CONCAT_TOKEN #undef CONCAT_TOKEN_HANDLER1 #undef MULTIPLY_BLOCKS_8x8 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_block_fp16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_block_fp16.cl index eff4b3111d3..376a257c783 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_block_fp16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_block_fp16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" // --------------------------------------------------------------------------------------------------------------------- // Just-in-time macro definitions: @@ -275,7 +275,6 @@ KERNEL (fully_connected_gpu_xb_xb_block_fp16)( #undef CONCAT_TOKEN_HANDLER1 #undef CONCAT_TOKEN -#undef MAKE_VECTOR_TYPE #undef CVT_UNIT #undef CHUNK_UNITS_TYPE #undef AS_CHUNK diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_ref.cl index 1fa78d7e54a..2707cd3bfa7 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" // Required JIT constants: // - FP16_SUPPORTED - [0/1] Value indicating whether device supports FP16 OpenCL extension (cl_khr_fp16). diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_oi_b8_fp32_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_oi_b8_fp32_ref.cl index a4f8bbb0235..193e770a9ad 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_oi_b8_fp32_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_oi_b8_fp32_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" __attribute__((reqd_work_group_size(8, 1, 1))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_oi_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_oi_ref.cl index 568ace580d9..4be0e33a503 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_oi_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_oi_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" // Required JIT constants: // - FP16_SUPPORTED - [0/1] Value indicating whether device supports FP16 OpenCL extension (cl_khr_fp16). diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fs_byx_fsv32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fs_byx_fsv32.cl index 9e1afb4e4c3..1ed279d7c7a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fs_byx_fsv32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fs_byx_fsv32.cl @@ -2,7 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" +#include "include/batch_headers/common.cl" #include "include/unit_type.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_imad.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_imad.cl index 4117d5a63e6..2b5402136ec 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_imad.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_imad.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/imad.cl" #define SIMD_SIZE 16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_yxfb_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_yxfb_ref.cl index 1fc91ba0a73..9f86457a8e4 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_yxfb_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_yxfb_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/reshape_dims.cl" // Required JIT constants: diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_1x1_opt_fp32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_1x1_opt_fp32.cl index 22dce26d0fb..205e16301f6 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_1x1_opt_fp32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_1x1_opt_fp32.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define SIMD_SIZE 8 __attribute__((intel_reqd_sub_group_size(SIMD_SIZE))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_iyxo.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_iyxo.cl index a038a0fda6b..bcb7fc7a40f 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_iyxo.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_iyxo.cl @@ -2,8 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/image_data.cl" __attribute__((intel_reqd_sub_group_size(SUB_GROUP_SIZE))) __attribute__((reqd_work_group_size(1, 1, SUB_GROUP_SIZE))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_os_iyx_osv16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_os_iyx_osv16.cl index 77d979128bb..57674e4cfa9 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_os_iyx_osv16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_os_iyx_osv16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" // --------------------------------------------------------------------------------------------------------------------- diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_imad.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_imad.cl index e618369d22e..4aa3cb861aa 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_imad.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_imad.cl @@ -2,7 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/imad.cl" #if QUANTIZATION_TERM # define ACCUMULATOR_TYPE int diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_ref.cl index 732885177f7..4b2918e22b7 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" KERNEL(kernel_name)( const __global INPUT0_TYPE *conv_input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_yxfb_yxio_b16_fp16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_yxfb_yxio_b16_fp16.cl index dcfe8f60a8e..842a98714e7 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_yxfb_yxio_b16_fp16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_yxfb_yxio_b16_fp16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/sub_group.cl" __attribute__((intel_reqd_sub_group_size(16))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_elements_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_elements_ref.cl index d03c1c85b13..2c4101f5db2 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_elements_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_elements_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define GET_OUTPUT_INDEX(prefix, idx_order) CAT(prefix, _GET_INDEX)(idx_order) @@ -18,7 +18,6 @@ KERNEL(gather_elements_ref)(const __global INPUT0_TYPE* data, const uint dim0 = get_global_id(0); const uint dim1 = get_global_id(1); const uint dim2 = get_global_id(2); - // Calculate indice index #if INPUT1_DIMS == 4 #define ORDER b,f,y,x diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl index 5ba38da8432..1c024128796 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl @@ -2,7 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" #define GET_UPDATES_INDEX(prefix, idx_order) CAT(prefix, _GET_INDEX)(idx_order) #define GET_OUTPUT_INDEX(out_order) OUTPUT_GET_INDEX(out_order) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_ref.cl index 978edce3356..8b114350f2f 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #ifdef INDEX_DIM inline uint FUNC(get_positive_index)(int in) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_tree_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_tree_gpu_ref.cl index c8ae6f39f11..fa71712dc30 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_tree_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_tree_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(gather_tree_gpu_ref)( const __global INPUT0_TYPE* step_input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_mmad_int8.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_mmad_int8.cl index 5ba378d7fd0..ad49b6ae800 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_mmad_int8.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_mmad_int8.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/mmad.cl" #define PACK_SIZE 4 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_mmad_int8_slm.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_mmad_int8_slm.cl index 6aac3b7672f..845d842bb91 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_mmad_int8_slm.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_mmad_int8_slm.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/mmad.cl" #define AS_TYPE(type, val) CAT(as_, type)(val) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_ref.cl index e476027e140..cfb403bb4b7 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_ref.cl @@ -2,7 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" // Required JIT definitions: // TRANSPOSE_INPUT0 [1/0] - whether to tranpose first input. diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_tiled_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_tiled_opt.cl index f87fbe7c587..ae79242b369 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_tiled_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gemm_tiled_opt.cl @@ -2,7 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/unit_type.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_bwd_data.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_bwd_data.cl index 0c49cc51e9b..f732acc922a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_bwd_data.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_bwd_data.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #define INPUT_TYPE8 MAKE_VECTOR_TYPE(INPUT0_TYPE, 8) #define OUTPUT_TYPE8 MAKE_VECTOR_TYPE(OUTPUT_TYPE, 8) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_fwd_data_f16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_fwd_data_f16.cl index 85e31e12908..96198ba83d9 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_fwd_data_f16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_fwd_data_f16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/unit_type.cl" #define WITH_ELTWISE 1 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_fwd_data_f32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_fwd_data_f32.cl index 4b7cefd5fc9..32a1d6f8275 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_fwd_data_f32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gen9_common_conv_fwd_data_f32.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/unit_type.cl" #define WITH_ELTWISE 1 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/generic_eltwise_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/generic_eltwise_ref.cl index 30b1e798fee..ebc302a7a5f 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/generic_eltwise_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/generic_eltwise_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #if ELTWISE_LAYOUT_BASED || QUANTIZATION_TERM || ELTWISE_BROADCAST #define GET_INDEX(prefix, num, idx_order) CAT(CAT(prefix, num), _GET_INDEX_SAFE)(idx_order) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/grn_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/grn_ref.cl index d52c9dc4ee0..f02f0972892 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/grn_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/grn_ref.cl @@ -2,8 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/acc_type.cl" KERNEL(grn_ref)( const __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/acc_type.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/acc_type.cl new file mode 100644 index 00000000000..19023bb03c8 --- /dev/null +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/acc_type.cl @@ -0,0 +1,8 @@ +// Defines default accumulator type. +// TODO: currently we calculate on float32 because it's lot of "add" operation and it stuck on the value "8192.0f" +// TODO: Get rid of this include and generate proper accumulator type on host (when needed) +#if !defined(ACCUMULATOR_TYPE) + #define ACCUMULATOR_TYPE float + #define TO_ACCUMULATOR_TYPE(v) (float)(v) + #define ACCUMULATOR_TYPE_ZERO 0.0f +#endif diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/common.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/common.cl similarity index 100% rename from inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/common.cl rename to inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/common.cl diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/data_types.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/data_types.cl similarity index 86% rename from inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/data_types.cl rename to inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/data_types.cl index 7643eeb94fe..9dd0bec0f0f 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/data_types.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/data_types.cl @@ -2,17 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -// Defines default accumulator type. -// TODO: currently we calculate on float32 because it's lot of "add" operation and it stuck on the value "8192.0f" -// TODO: Get rid of this include and generate proper accumulator type on host (when needed) -#if !defined(ACCUMULATOR_TYPE) - #define ACCUMULATOR_TYPE float - #define TO_ACCUMULATOR_TYPE(v) (float)(v) - #define ACCUMULATOR_TYPE_ZERO 0.0f -#endif - -inline void FUNC(sub_group_block_write_uchar16)(__global uchar* outPtr, uchar16 v) { +inline void sub_group_block_write_uchar16(__global uchar* outPtr, uchar16 v) { #ifdef cl_intel_subgroups_char intel_sub_group_block_write_uc16(outPtr, v); #else @@ -37,7 +28,7 @@ inline void FUNC(sub_group_block_write_uchar16)(__global uchar* outPtr, uchar16 #endif } -inline uchar16 FUNC(sub_group_block_read_uchar16)(const __global uchar* ptr) __attribute__((overloadable)) { +inline uchar16 sub_group_block_read_uchar16(const __global uchar* ptr) __attribute__((overloadable)) { #ifdef cl_intel_subgroups_char // WA for compiler support // return intel_sub_group_block_read_uc16(ptr); @@ -68,7 +59,7 @@ inline uchar16 FUNC(sub_group_block_read_uchar16)(const __global uchar* ptr) __a #endif } -inline uchar16 FUNC(sub_group_block_read_uchar16)(const __local uchar* ptr) __attribute__((overloadable)) { +inline uchar16 sub_group_block_read_uchar16(const __local uchar* ptr) __attribute__((overloadable)) { #if LOCAL_BLOCK_IO_SUPPORTED && defined(cl_intel_subgroup_local_block_io) && defined(cl_intel_subgroups_char) // WA for compiler support // return intel_sub_group_block_read_uc16(ptr); @@ -99,7 +90,7 @@ inline uchar16 FUNC(sub_group_block_read_uchar16)(const __local uchar* ptr) __at #endif } -inline void FUNC(sub_group_block_write_uchar8)(__global uchar* outPtr, uchar8 v) +inline void sub_group_block_write_uchar8(__global uchar* outPtr, uchar8 v) { #ifdef cl_intel_subgroups_char intel_sub_group_block_write_uc8(outPtr, v); @@ -117,7 +108,7 @@ inline void FUNC(sub_group_block_write_uchar8)(__global uchar* outPtr, uchar8 v) #endif } -inline uchar8 FUNC(sub_group_block_read_uchar8)(const __global uchar* ptr) __attribute__((overloadable)) { +inline uchar8 sub_group_block_read_uchar8(const __global uchar* ptr) __attribute__((overloadable)) { #ifdef cl_intel_subgroups_char return intel_sub_group_block_read_uc8(ptr); #else @@ -138,7 +129,7 @@ inline uchar8 FUNC(sub_group_block_read_uchar8)(const __global uchar* ptr) __att #endif } -inline uchar8 FUNC(sub_group_block_read_uchar8)(const __local uchar* ptr) __attribute__((overloadable)) { +inline uchar8 sub_group_block_read_uchar8(const __local uchar* ptr) __attribute__((overloadable)) { #if LOCAL_BLOCK_IO_SUPPORTED && defined(cl_intel_subgroup_local_block_io) && defined(cl_intel_subgroups_char) return intel_sub_group_block_read_uc8(ptr); #else @@ -159,7 +150,7 @@ inline uchar8 FUNC(sub_group_block_read_uchar8)(const __local uchar* ptr) __attr #endif } -inline void FUNC(sub_group_block_write_uchar4)(__global uchar* outPtr, uchar4 v) { +inline void sub_group_block_write_uchar4(__global uchar* outPtr, uchar4 v) { #ifdef cl_intel_subgroups_char intel_sub_group_block_write_uc4(outPtr, v); #else @@ -172,7 +163,7 @@ inline void FUNC(sub_group_block_write_uchar4)(__global uchar* outPtr, uchar4 v) #endif } -inline uchar4 FUNC(sub_group_block_read_uchar4)(const __global uchar* ptr) __attribute__((overloadable)) { +inline uchar4 sub_group_block_read_uchar4(const __global uchar* ptr) __attribute__((overloadable)) { #ifdef cl_intel_subgroups_char return intel_sub_group_block_read_uc4(ptr); #else @@ -189,7 +180,7 @@ inline uchar4 FUNC(sub_group_block_read_uchar4)(const __global uchar* ptr) __att #endif } -inline uchar4 FUNC(sub_group_block_read_uchar4)(const __local uchar* ptr) __attribute__((overloadable)) { +inline uchar4 sub_group_block_read_uchar4(const __local uchar* ptr) __attribute__((overloadable)) { #if LOCAL_BLOCK_IO_SUPPORTED && defined(cl_intel_subgroup_local_block_io) && defined(cl_intel_subgroups_char) return intel_sub_group_block_read_uc4(ptr); #else @@ -206,7 +197,7 @@ inline uchar4 FUNC(sub_group_block_read_uchar4)(const __local uchar* ptr) __attr #endif } -inline void FUNC(sub_group_block_write_uchar2)(__global uchar* outPtr, uchar2 v) { +inline void sub_group_block_write_uchar2(__global uchar* outPtr, uchar2 v) { #ifdef cl_intel_subgroups_char intel_sub_group_block_write_uc2(outPtr, v); #else @@ -217,7 +208,7 @@ inline void FUNC(sub_group_block_write_uchar2)(__global uchar* outPtr, uchar2 v) #endif } -inline uchar2 FUNC(sub_group_block_read_uchar2)(const __global uchar* ptr) __attribute__((overloadable)) { +inline uchar2 sub_group_block_read_uchar2(const __global uchar* ptr) __attribute__((overloadable)) { #ifdef cl_intel_subgroups_char return intel_sub_group_block_read_uc2(ptr); #else @@ -232,7 +223,7 @@ inline uchar2 FUNC(sub_group_block_read_uchar2)(const __global uchar* ptr) __att #endif } -inline uchar2 FUNC(sub_group_block_read_uchar2)(const __local uchar* ptr) __attribute__((overloadable)) { +inline uchar2 sub_group_block_read_uchar2(const __local uchar* ptr) __attribute__((overloadable)) { #if LOCAL_BLOCK_IO_SUPPORTED && defined(cl_intel_subgroup_local_block_io) && defined(cl_intel_subgroups_char) return intel_sub_group_block_read_uc2(ptr); #else @@ -247,7 +238,7 @@ inline uchar2 FUNC(sub_group_block_read_uchar2)(const __local uchar* ptr) __attr #endif } -inline void FUNC(sub_group_block_write_uchar)(__global uchar* outPtr, uchar v) { +inline void sub_group_block_write_uchar(__global uchar* outPtr, uchar v) { #ifdef cl_intel_subgroups_char intel_sub_group_block_write_uc(outPtr, v); #else @@ -257,7 +248,7 @@ inline void FUNC(sub_group_block_write_uchar)(__global uchar* outPtr, uchar v) { #endif } -inline uchar FUNC(sub_group_block_read_uchar)(const __global uchar* ptr) __attribute__((overloadable)) { +inline uchar sub_group_block_read_uchar(const __global uchar* ptr) __attribute__((overloadable)) { #ifdef cl_intel_subgroups_char return intel_sub_group_block_read_uc(ptr); #else @@ -271,7 +262,7 @@ inline uchar FUNC(sub_group_block_read_uchar)(const __global uchar* ptr) __attri #endif } -inline uchar FUNC(sub_group_block_read_uchar)(const __local uchar* ptr) __attribute__((overloadable)) { +inline uchar sub_group_block_read_uchar(const __local uchar* ptr) __attribute__((overloadable)) { #if LOCAL_BLOCK_IO_SUPPORTED && defined(cl_intel_subgroup_local_block_io) && defined(cl_intel_subgroups_char) return intel_sub_group_block_read_uc(ptr); #else @@ -356,17 +347,17 @@ inline uchar FUNC(sub_group_block_read_uchar)(const __local uchar* ptr) __attrib #define BLOCK_WRITE_FUNC_size4 intel_sub_group_block_write #define BLOCK_WRITE_FUNC(type_size) CAT(BLOCK_WRITE_FUNC_size, type_size) -#define BLOCK_READ_UC_1(ptr) FUNC_CALL(sub_group_block_read_uchar)(ptr) -#define BLOCK_READ_UC_2(ptr) FUNC_CALL(sub_group_block_read_uchar2)(ptr) -#define BLOCK_READ_UC_4(ptr) FUNC_CALL(sub_group_block_read_uchar4)(ptr) -#define BLOCK_READ_UC_8(ptr) FUNC_CALL(sub_group_block_read_uchar8)(ptr) -#define BLOCK_READ_UC_16(ptr) FUNC_CALL(sub_group_block_read_uchar16)(ptr) +#define BLOCK_READ_UC_1(ptr) sub_group_block_read_uchar(ptr) +#define BLOCK_READ_UC_2(ptr) sub_group_block_read_uchar2(ptr) +#define BLOCK_READ_UC_4(ptr) sub_group_block_read_uchar4(ptr) +#define BLOCK_READ_UC_8(ptr) sub_group_block_read_uchar8(ptr) +#define BLOCK_READ_UC_16(ptr) sub_group_block_read_uchar16(ptr) -#define BLOCK_WRITE_UC_1(ptr, val) FUNC_CALL(sub_group_block_write_uchar)(ptr, val) -#define BLOCK_WRITE_UC_2(ptr, val) FUNC_CALL(sub_group_block_write_uchar2)(ptr, val) -#define BLOCK_WRITE_UC_4(ptr, val) FUNC_CALL(sub_group_block_write_uchar4)(ptr, val) -#define BLOCK_WRITE_UC_8(ptr, val) FUNC_CALL(sub_group_block_write_uchar8)(ptr, val) -#define BLOCK_WRITE_UC_16(ptr, val) FUNC_CALL(sub_group_block_write_uchar16)(ptr, val) +#define BLOCK_WRITE_UC_1(ptr, val) sub_group_block_write_uchar(ptr, val) +#define BLOCK_WRITE_UC_2(ptr, val) sub_group_block_write_uchar2(ptr, val) +#define BLOCK_WRITE_UC_4(ptr, val) sub_group_block_write_uchar4(ptr, val) +#define BLOCK_WRITE_UC_8(ptr, val) sub_group_block_write_uchar8(ptr, val) +#define BLOCK_WRITE_UC_16(ptr, val) sub_group_block_write_uchar16(ptr, val) #define BLOCK_READN_FUNC_size1(vector_size) CAT(BLOCK_READ_UC_, vector_size) #define BLOCK_READN_FUNC_SIZE_DEF(type_size, vector_size) MAKE_VECTOR_TYPE(BLOCK_READ_FUNC(type_size), vector_size) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/fetch_data.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/fetch_data.cl similarity index 93% rename from inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/fetch_data.cl rename to inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/fetch_data.cl index b50850d3f00..5ba559dca1f 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/fetch_data.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/fetch_data.cl @@ -87,7 +87,7 @@ ((b) / (sub_group_size))*CAT(prefix, _BATCH_PITCH) \ ) -inline uint FUNC(get_b_fs_yx_fsv_index)(uint b, uint f, uint y, uint x, +inline uint get_b_fs_yx_fsv_index(uint b, uint f, uint y, uint x, uint x_size, uint y_size, uint f_size, uint b_size, uint b_pad_before, uint b_pad_after, uint f_pad_before, uint f_pad_after, @@ -111,7 +111,7 @@ inline uint FUNC(get_b_fs_yx_fsv_index)(uint b, uint f, uint y, uint x, return output_offset; } -inline uint FUNC(get_b_fs_yx_fsv_index_safe)(uint b, uint f, uint y, uint x, +inline uint get_b_fs_yx_fsv_index_safe(uint b, uint f, uint y, uint x, uint x_size, uint y_size, uint f_size, uint b_size, uint b_pad_before, uint b_pad_after, uint f_pad_before, uint f_pad_after, @@ -136,7 +136,7 @@ inline uint FUNC(get_b_fs_yx_fsv_index_safe)(uint b, uint f, uint y, uint x, } #define GET_DATA_B_FS_YX_FSV16_INDEX(prefix, b, f, y, x) \ - FUNC_CALL(get_b_fs_yx_fsv_index)( \ + get_b_fs_yx_fsv_index( \ b, f, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ @@ -152,7 +152,7 @@ inline uint FUNC(get_b_fs_yx_fsv_index_safe)(uint b, uint f, uint y, uint x, CAT(prefix, _PAD_AFTER_SIZE_X), 16) #define GET_DATA_B_FS_YX_FSV16_INDEX_SAFE(prefix, b, f, y, x) \ - FUNC_CALL(get_b_fs_yx_fsv_index_safe)( \ + get_b_fs_yx_fsv_index_safe( \ b, f, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ @@ -168,7 +168,7 @@ inline uint FUNC(get_b_fs_yx_fsv_index_safe)(uint b, uint f, uint y, uint x, CAT(prefix, _PAD_AFTER_SIZE_X), 16) #define GET_DATA_B_FS_YX_FSV4_INDEX(prefix, b, f, y, x) \ - FUNC_CALL(get_b_fs_yx_fsv_index)( \ + get_b_fs_yx_fsv_index( \ b, f, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ @@ -184,7 +184,7 @@ inline uint FUNC(get_b_fs_yx_fsv_index_safe)(uint b, uint f, uint y, uint x, CAT(prefix, _PAD_AFTER_SIZE_X), 4) #define GET_DATA_B_FS_YX_FSV4_INDEX_SAFE(prefix, b, f, y, x) \ - FUNC_CALL(get_b_fs_yx_fsv_index_safe)( \ + get_b_fs_yx_fsv_index_safe( \ b, f, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ @@ -200,7 +200,7 @@ inline uint FUNC(get_b_fs_yx_fsv_index_safe)(uint b, uint f, uint y, uint x, CAT(prefix, _PAD_AFTER_SIZE_X), 4) #define GET_DATA_B_FS_YX_FSV32_INDEX(prefix, b, f, y, x) \ - FUNC_CALL(get_b_fs_yx_fsv_index)( \ + get_b_fs_yx_fsv_index( \ b, f, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ @@ -216,7 +216,7 @@ inline uint FUNC(get_b_fs_yx_fsv_index_safe)(uint b, uint f, uint y, uint x, CAT(prefix, _PAD_AFTER_SIZE_X), 32) #define GET_DATA_B_FS_YX_FSV32_INDEX_SAFE(prefix, b, f, y, x) \ - FUNC_CALL(get_b_fs_yx_fsv_index_safe)( \ + get_b_fs_yx_fsv_index_safe( \ b, f, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ @@ -232,7 +232,7 @@ inline uint FUNC(get_b_fs_yx_fsv_index_safe)(uint b, uint f, uint y, uint x, CAT(prefix, _PAD_AFTER_SIZE_X), 32) #define GET_DATA_FS_B_YX_FSV32_INDEX(prefix, b, f, y, x) \ - FUNC_CALL(get_fs_b_yx_fsv32_index)( \ + get_fs_b_yx_fsv32_index( \ b, f, y, x, \ CAT(prefix, _PAD_BEFORE_SIZE_X), \ CAT(prefix, _SIZE_X), \ @@ -243,7 +243,7 @@ inline uint FUNC(get_b_fs_yx_fsv_index_safe)(uint b, uint f, uint y, uint x, CAT(prefix, _PAD_BEFORE_FEATURE_NUM), \ CAT(prefix, _BATCH_NUM)) -inline uint FUNC(get_fs_b_yx_fsv32_index)(uint b, uint f, uint y, uint x, +inline uint get_fs_b_yx_fsv32_index(uint b, uint f, uint y, uint x, uint x_pad_before, uint x_size, uint x_pad_after, uint y_pad_before, uint y_size, uint y_pad_after, uint f_pad_before, @@ -278,7 +278,7 @@ inline uint FUNC(get_fs_b_yx_fsv32_index)(uint b, uint f, uint y, uint x, } #define GET_DATA_B_FS_ZYX_FSV16_INDEX(prefix, b, f, z, y, x) \ - FUNC_CALL(get_b_fs_zyx_fsv_index)( \ + get_b_fs_zyx_fsv_index( \ b, f, z, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ @@ -294,7 +294,7 @@ inline uint FUNC(get_fs_b_yx_fsv32_index)(uint b, uint f, uint y, uint x, CAT(prefix, _PAD_AFTER_SIZE_X), 16) #define GET_DATA_B_FS_ZYX_FSV16_INDEX_SAFE(prefix, b, f, z, y, x) \ - FUNC_CALL(get_b_fs_zyx_fsv_index_safe)( \ + get_b_fs_zyx_fsv_index_safe( \ b, f, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -311,7 +311,7 @@ inline uint FUNC(get_fs_b_yx_fsv32_index)(uint b, uint f, uint y, uint x, #define GET_DATA_B_FS_ZYX_FSV32_INDEX(prefix, b, f, z, y, x) \ - FUNC_CALL(get_b_fs_zyx_fsv_index)( \ + get_b_fs_zyx_fsv_index( \ b, f, z, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ @@ -327,7 +327,7 @@ inline uint FUNC(get_fs_b_yx_fsv32_index)(uint b, uint f, uint y, uint x, CAT(prefix, _PAD_AFTER_SIZE_X), 32) #define GET_DATA_B_FS_ZYX_FSV32_INDEX_SAFE(prefix, b, f, z, y, x) \ - FUNC_CALL(get_b_fs_zyx_fsv_index_safe)( \ + get_b_fs_zyx_fsv_index_safe( \ b, f, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -342,7 +342,7 @@ inline uint FUNC(get_fs_b_yx_fsv32_index)(uint b, uint f, uint y, uint x, CAT(prefix, _PAD_BEFORE_SIZE_X), \ CAT(prefix, _PAD_AFTER_SIZE_X), 32) -inline uint FUNC(get_b_fs_zyx_fsv_index)(uint b, uint f, uint z, uint y, uint x, +inline uint get_b_fs_zyx_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, @@ -370,7 +370,7 @@ inline uint FUNC(get_b_fs_zyx_fsv_index)(uint b, uint f, uint z, uint y, uint x return output_offset; } -inline uint FUNC(get_b_fs_zyx_fsv_index_safe)(uint b, uint f, uint z, uint y, uint x, +inline uint get_b_fs_zyx_fsv_index_safe(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, @@ -397,7 +397,7 @@ inline uint FUNC(get_b_fs_zyx_fsv_index_safe)(uint b, uint f, uint z, uint y, u return output_offset; } -inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index_safe)(uint b, uint f, uint z, uint y, uint x, +inline uint get_bs_fs_zyx_bsv_fsv_index_safe(uint b, uint f, uint z, uint y, uint x, uint x_size, uint y_size, uint z_size, uint f_size, uint b_size, uint f_pad_before, uint f_pad_after, uint z_pad_before, uint z_pad_after, @@ -426,7 +426,7 @@ 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_bsv_fsv_index)(uint b, uint f, uint z, uint y, uint x, +inline uint 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, @@ -459,7 +459,7 @@ inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index)(uint b, uint f, uint z, uint y, u } #define GET_DATA_BS_FS_YX_BSV16_FSV16_INDEX(prefix, b, f, y, x) \ - FUNC_CALL(get_bs_fs_zyx_bsv_fsv_index)( \ + get_bs_fs_zyx_bsv_fsv_index( \ b, f, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -475,7 +475,7 @@ inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index)(uint b, uint f, uint z, uint y, u 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)( \ + get_bs_fs_zyx_bsv_fsv_index( \ b, f, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -491,7 +491,7 @@ inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index)(uint b, uint f, uint z, uint y, u 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)( \ + get_bs_fs_zyx_bsv_fsv_index( \ b, f, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -507,7 +507,7 @@ inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index)(uint b, uint f, uint z, uint y, u 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)( \ + get_bs_fs_zyx_bsv_fsv_index( \ b, f, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -523,7 +523,7 @@ inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index)(uint b, uint f, uint z, uint y, u 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)( \ + get_bs_fs_zyx_bsv_fsv_index( \ b, f, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -539,7 +539,7 @@ inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index)(uint b, uint f, uint z, uint y, u 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_bsv_fsv_index)( \ + get_bs_fs_zyx_bsv_fsv_index( \ b, f, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -555,7 +555,7 @@ inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index)(uint b, uint f, uint z, uint y, u 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)( \ + get_bs_fs_zyx_bsv_fsv_index_safe( \ b, f, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -640,7 +640,7 @@ inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index)(uint b, uint f, uint z, uint y, u 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)( \ + get_bs_fs_zyx_bsv_fsv_index_safe( \ b, f, z, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ @@ -656,12 +656,3 @@ inline uint FUNC(get_bs_fs_zyx_bsv_fsv_index)(uint b, uint f, uint z, uint y, u CAT(prefix, _PAD_BEFORE_SIZE_X), \ CAT(prefix, _PAD_AFTER_SIZE_X), 16, 16) -#define DECLARE_SAMPLER const sampler_t imageSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST - -#if FP16_UNIT_USED - #define IMAGE_READ(image, coord) read_imageh((image), imageSampler, (coord)) - #define IMAGE_WRITE(image, coord, val) write_imageh((image), (coord), (val)) -#else - #define IMAGE_READ(image, coord) read_imagef((image), imageSampler, (coord)) - #define IMAGE_WRITE(image, coord, val) write_imagef((image), (coord), (val)) -#endif diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/fetch_weights.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/fetch_weights.cl similarity index 89% rename from inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/fetch_weights.cl rename to inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/fetch_weights.cl index 43f73705237..bcfa110c4c8 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/fetch_weights.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/fetch_weights.cl @@ -51,7 +51,7 @@ ) #define GET_FILTER_OS_IS_YX_ISV8_OSV16_ISV2_INDEX(prefix, o, i, y, x, sub_group_size) \ - FUNC_CALL(get_os_is_zyx_isv8_osv16_isv2_index)( \ + get_os_is_zyx_isv8_osv16_isv2_index( \ 0, o, i, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -63,7 +63,7 @@ ) #define GET_FILTER_OS_IS_ZYX_ISV8_OSV16_ISV2_INDEX(prefix, o, i, z, y, x, sub_group_size) \ - FUNC_CALL(get_os_is_zyx_isv8_osv16_isv2_index)( \ + get_os_is_zyx_isv8_osv16_isv2_index( \ 0, o, i, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -74,7 +74,7 @@ CAT(prefix, _OFFSET) \ ) -inline uint FUNC(get_os_is_zyx_osv_isv_index)(uint o, uint i, uint z, uint y, uint x, +inline uint get_os_is_zyx_osv_isv_index(uint o, uint i, uint z, uint y, uint x, uint x_size, uint y_size, uint z_size, uint i_size, uint o_size, uint osv_size, uint isv_size) { const uint isv = i % isv_size; @@ -100,7 +100,7 @@ inline uint FUNC(get_os_is_zyx_osv_isv_index)(uint o, uint i, uint z, uint y, ui return output_offset; } -inline uint FUNC(get_g_os_is_zyx_osv_isv_index)(uint g, uint o, uint i, uint z, uint y, uint x, +inline uint get_g_os_is_zyx_osv_isv_index(uint g, uint o, uint i, uint z, uint y, uint x, uint x_size, uint y_size, uint z_size, uint i_size, uint o_size, uint osv_size, uint isv_size) { const uint isv = i % isv_size; @@ -129,7 +129,7 @@ inline uint FUNC(get_g_os_is_zyx_osv_isv_index)(uint g, uint o, uint i, uint z, } #define GET_FILTER_G_OS_IS_ZYX_OSV16_ISV16_INDEX(prefix, g, o, i, z, y, x) \ - FUNC_CALL(get_g_os_is_zyx_osv_isv_index)( \ + get_g_os_is_zyx_osv_isv_index( \ g, o, i, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -140,7 +140,7 @@ inline uint FUNC(get_g_os_is_zyx_osv_isv_index)(uint g, uint o, uint i, uint z, 16) #define GET_FILTER_OS_IS_YX_OSV16_ISV16_INDEX(prefix, o, i, y, x) \ - FUNC_CALL(get_os_is_zyx_osv_isv_index)( \ + get_os_is_zyx_osv_isv_index( \ o, i, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -151,7 +151,7 @@ inline uint FUNC(get_g_os_is_zyx_osv_isv_index)(uint g, uint o, uint i, uint z, 16) #define GET_FILTER_OS_IS_ZYX_OSV16_ISV16_INDEX(prefix, o, i, z, y, x) \ - FUNC_CALL(get_os_is_zyx_osv_isv_index)( \ + get_os_is_zyx_osv_isv_index( \ o, i, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -162,7 +162,7 @@ inline uint FUNC(get_g_os_is_zyx_osv_isv_index)(uint g, uint o, uint i, uint z, 16) #define GET_FILTER_OS_IS_ZYX_OSV32_ISV16_INDEX(prefix, o, i, z, y, x) \ - FUNC_CALL(get_os_is_zyx_osv_isv_index)( \ + get_os_is_zyx_osv_isv_index( \ o, i, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -173,7 +173,7 @@ inline uint FUNC(get_g_os_is_zyx_osv_isv_index)(uint g, uint o, uint i, uint z, 16) #define GET_FILTER_OS_IS_ZYX_OSV64_ISV16_INDEX(prefix, o, i, z, y, x) \ - FUNC_CALL(get_os_is_zyx_osv_isv_index)( \ + get_os_is_zyx_osv_isv_index( \ o, i, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -184,7 +184,7 @@ inline uint FUNC(get_g_os_is_zyx_osv_isv_index)(uint g, uint o, uint i, uint z, 16) #define GET_FILTER_G_OS_IS_YX_ISV8_OSV16_ISV2_INDEX(prefix, g, o, i, y, x, sub_group_size) \ - FUNC_CALL(get_os_is_zyx_isv8_osv16_isv2_index)( \ + get_os_is_zyx_isv8_osv16_isv2_index( \ g, o, i, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -196,7 +196,7 @@ inline uint FUNC(get_g_os_is_zyx_osv_isv_index)(uint g, uint o, uint i, uint z, ) #define GET_FILTER_G_OS_IS_ZYX_ISV8_OSV16_ISV2_INDEX(prefix, g, o, i, z, y, x, sub_group_size) \ - FUNC_CALL(get_os_is_zyx_isv8_osv16_isv2_index)( \ + get_os_is_zyx_isv8_osv16_isv2_index( \ g, o, i, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -207,7 +207,7 @@ inline uint FUNC(get_g_os_is_zyx_osv_isv_index)(uint g, uint o, uint i, uint z, CAT(prefix, _OFFSET) \ ) -inline uint FUNC(get_os_is_zyx_isv8_osv16_isv2_index)(uint g, uint o, uint i, uint z, uint y, uint x, uint x_size, uint y_size, uint z_size, +inline uint get_os_is_zyx_isv8_osv16_isv2_index(uint g, uint o, uint i, uint z, uint y, uint x, uint x_size, uint y_size, uint z_size, uint g_size, uint o_size, uint i_size, uint offset) { const uint group_offset = g * o_size * i_size * z_size * y_size * x_size; @@ -227,7 +227,7 @@ inline uint FUNC(get_os_is_zyx_isv8_osv16_isv2_index)(uint g, uint o, uint i, u return idx; } -inline uint FUNC(get_os_zyxi_osv16_index)(uint o, uint i, uint z, uint y, uint x, uint i_size, uint o_size, uint x_size, uint y_size, uint z_size) +inline uint get_os_zyxi_osv16_index(uint o, uint i, uint z, uint y, uint x, uint i_size, uint o_size, uint x_size, uint y_size, uint z_size) { const size_t idx = o%16 + (o / 16)*i_size*x_size*y_size*z_size*16 + 16 *(i+ x*i_size + y*i_size*x_size + z*i_size*x_size*y_size); @@ -235,7 +235,7 @@ inline uint FUNC(get_os_zyxi_osv16_index)(uint o, uint i, uint z, uint y, uint x } #define GET_FILTER_OS_ZYXI_OSV16(prefix, o, i, z, y, x) \ - FUNC_CALL(get_os_zyxi_osv16_index)( \ + get_os_zyxi_osv16_index( \ o, i, z, y, x, CAT(prefix, _IFM_NUM), \ CAT(prefix, _OFM_NUM), \ CAT(prefix, _SIZE_X), \ @@ -338,7 +338,7 @@ inline uint FUNC(get_os_zyxi_osv16_index)(uint o, uint i, uint z, uint y, uint x ((o) / (sub_group_size))*CAT(prefix, _OFM_PITCH) \ ) -inline uint FUNC(get_gi_yxs_os_yxsv2_osv_index)(uint g, uint o, uint i, uint y, uint x, uint x_size, uint g_pitch, uint i_pitch, +inline uint get_gi_yxs_os_yxsv2_osv_index(uint g, uint o, uint i, uint y, uint x, uint x_size, uint g_pitch, uint i_pitch, uint y_pitch, uint x_pitch, uint offset, uint sub_group_size) { const uint aligned_ofm_line = x_pitch; @@ -363,7 +363,7 @@ inline uint FUNC(get_gi_yxs_os_yxsv2_osv_index)(uint g, uint o, uint i, uint y, } #define GET_FILTER_I_YXS_OS_YXSV2_OSV_INDEX(prefix, o, i, y, x, sub_group_size) \ - FUNC_CALL(get_gi_yxs_os_yxsv2_osv_index)( \ + get_gi_yxs_os_yxsv2_osv_index( \ 0, o, i, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _GROUPS_PITCH), \ @@ -373,7 +373,7 @@ inline uint FUNC(get_gi_yxs_os_yxsv2_osv_index)(uint g, uint o, uint i, uint y, CAT(prefix, _OFFSET), \ sub_group_size) -inline uint FUNC(get_giy_xs_os_xsv2_osv_index)(uint g, uint o, uint i, uint y, uint x, uint x_size, uint g_pitch, +inline uint get_giy_xs_os_xsv2_osv_index(uint g, uint o, uint i, uint y, uint x, uint x_size, uint g_pitch, uint i_pitch, uint y_pitch, uint x_pitch, uint offset, uint sub_group_size) { const uint aligned_ofm_line = x_pitch; @@ -409,7 +409,7 @@ inline uint FUNC(get_giy_xs_os_xsv2_osv_index)(uint g, uint o, uint i, uint y, u } #define GET_FILTER_IY_XS_OS_XSV2_OSV_INDEX(prefix, o, i, y, x, sub_group_size) \ - FUNC_CALL(get_giy_xs_os_xsv2_osv_index)( \ + get_giy_xs_os_xsv2_osv_index( \ 0, o, i, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _GROUPS_PITCH), \ @@ -419,7 +419,7 @@ inline uint FUNC(get_giy_xs_os_xsv2_osv_index)(uint g, uint o, uint i, uint y, u CAT(prefix, _OFFSET), \ sub_group_size) -inline uint FUNC(get_os_is_yx_isa8_osv8_isv4_index)(uint o, uint i, uint y, uint x, uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset) +inline uint get_os_is_yx_isa8_osv8_isv4_index(uint o, uint i, uint y, uint x, uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset) { const uint f_32_aligned = ((size_ifm + 31)/32) * 32; const uint isv2_idx = i % 4; @@ -438,14 +438,14 @@ inline uint FUNC(get_os_is_yx_isa8_osv8_isv4_index)(uint o, uint i, uint y, uint } #define GET_FILTER_OS_IS_YX_ISA8_OSV8_ISV4_INDEX(prefix, o, i, y, x) \ - FUNC_CALL(get_os_is_yx_isa8_osv8_isv4_index)( \ + get_os_is_yx_isa8_osv8_isv4_index( \ o, i, y, x, CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ CAT(prefix, _IFM_NUM), \ CAT(prefix, _OFM_NUM), \ CAT(prefix, _OFFSET)) -inline uint FUNC(get_os_is_zyx_isa8_osv8_isv4_index)(uint o, uint i, uint z, uint y, uint x, +inline uint get_os_is_zyx_isa8_osv8_isv4_index(uint o, uint i, uint z, uint y, uint x, uint size_x, uint size_y, uint size_z, uint size_ifm, uint size_ofm, uint offset) { @@ -467,7 +467,7 @@ inline uint FUNC(get_os_is_zyx_isa8_osv8_isv4_index)(uint o, uint i, uint z, uin } #define GET_FILTER_OS_IS_ZYX_ISA8_OSV8_ISV4_INDEX(prefix, o, i, z, y, x) \ - FUNC_CALL(get_os_is_zyx_isa8_osv8_isv4_index)( \ + get_os_is_zyx_isa8_osv8_isv4_index( \ o, i, z, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ @@ -476,7 +476,7 @@ inline uint FUNC(get_os_is_zyx_isa8_osv8_isv4_index)(uint o, uint i, uint z, uin CAT(prefix, _OFM_NUM), \ CAT(prefix, _OFFSET)) -inline uint FUNC(get_os_is_yx_isa8_osv16_isv4_index)(uint o, uint i, uint y, uint x, uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset) +inline uint get_os_is_yx_isa8_osv16_isv4_index(uint o, uint i, uint y, uint x, uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset) { const uint f_32_aligned = ((size_ifm + 31)/32) * 32; const uint isv2_idx = i % 4; @@ -495,14 +495,14 @@ inline uint FUNC(get_os_is_yx_isa8_osv16_isv4_index)(uint o, uint i, uint y, uin } #define GET_FILTER_OS_IS_YX_ISA8_OSV16_ISV4_INDEX(prefix, o, i, y, x) \ - FUNC_CALL(get_os_is_yx_isa8_osv16_isv4_index)( \ + get_os_is_yx_isa8_osv16_isv4_index( \ o, i, y, x, CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ CAT(prefix, _IFM_NUM), \ CAT(prefix, _OFM_NUM), \ CAT(prefix, _OFFSET)) -inline uint FUNC(get_os_is_zyx_isa8_osv16_isv4_index)(uint o, uint i, uint z, uint y, uint x, +inline uint get_os_is_zyx_isa8_osv16_isv4_index(uint o, uint i, uint z, uint y, uint x, uint size_x, uint size_y, uint size_z, uint size_ifm, uint size_ofm, uint offset) { @@ -524,7 +524,7 @@ inline uint FUNC(get_os_is_zyx_isa8_osv16_isv4_index)(uint o, uint i, uint z, ui } #define GET_FILTER_OS_IS_ZYX_ISA8_OSV16_ISV4_INDEX(prefix, o, i, z, y, x) \ - FUNC_CALL(get_os_is_zyx_isa8_osv16_isv4_index)( \ + get_os_is_zyx_isa8_osv16_isv4_index( \ o, i, z, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ @@ -533,7 +533,7 @@ inline uint FUNC(get_os_is_zyx_isa8_osv16_isv4_index)(uint o, uint i, uint z, ui CAT(prefix, _OFM_NUM), \ CAT(prefix, _OFFSET)) -inline uint FUNC(get_os_is_yx_isa8_osv8_isv4_swizzled_by_4_index)(uint o, uint i, uint y, uint x, uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset) +inline uint get_os_is_yx_isa8_osv8_isv4_swizzled_by_4_index(uint o, uint i, uint y, uint x, uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset) { const uint o_swizzled = (o % 4) * 8 + ((o % 32) / 4) + (o / 32) * 32; @@ -553,7 +553,7 @@ inline uint FUNC(get_os_is_yx_isa8_osv8_isv4_swizzled_by_4_index)(uint o, uint i return idx; } -inline uint FUNC(get_os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4_index)(uint o, uint i, uint y, uint x, uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset) +inline uint get_os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4_index(uint o, uint i, uint y, uint x, uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset) { const uint o_swizzled = (o % 4) * 8 + ((o % 32) / 4) + (o / 32) * 32; const uint isv_idx = i % 4; @@ -578,7 +578,7 @@ inline uint FUNC(get_os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4_index)(uint o, u return idx; } -inline uint FUNC(get_os_is_zyx_osa4_isa8_osv8_isv4_swizzled_by_4_index)(uint o, uint i, uint z, uint y, uint x, +inline uint get_os_is_zyx_osa4_isa8_osv8_isv4_swizzled_by_4_index(uint o, uint i, uint z, uint y, uint x, uint size_x, uint size_y, uint size_z, uint size_ifm, uint size_ofm, uint offset) { @@ -606,7 +606,7 @@ 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 z, uint y, uint x, +inline uint get_g_os_is_yx_osa4_isa8_osv8_isv4(uint g, uint o, uint i, uint z, uint y, uint x, uint size_x, uint size_y, uint size_z, uint size_ifm, uint size_ofm, uint offset) { const uint isv_idx = i % 4; @@ -634,8 +634,8 @@ inline uint FUNC(get_g_os_is_yx_osa4_isa8_osv8_isv4)(uint g, uint o, uint i, uin return idx; } -inline uint FUNC(get_g_os_is_yx_osa4_isa8_osv8_isv2)(uint g, uint o, uint i, uint z, uint y, uint x, - uint size_x, uint size_y, uint size_z, uint size_ifm, uint size_ofm, uint offset) +inline uint get_g_os_is_yx_osa4_isa8_osv8_isv2(uint g, uint o, uint i, uint z, uint y, uint x, + uint size_x, uint size_y, uint size_z, uint size_ifm, uint size_ofm, uint offset) { const uint isv_idx = i % 2; const uint isa_idx = (i / 2) % 8; @@ -662,7 +662,7 @@ inline uint FUNC(get_g_os_is_yx_osa4_isa8_osv8_isv2)(uint g, uint o, uint i, uin return idx; } -inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_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) +inline uint get_g_os_is_yx_osa2_isa8_osv16_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; @@ -688,7 +688,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv4)(uint g, uint o, uint i, ui return idx; } -inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_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) +inline uint get_g_os_is_yx_osa2_isa8_osv16_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; @@ -715,7 +715,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui } #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)( \ + get_g_os_is_yx_osa4_isa8_osv8_isv4( \ 0, o, i, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -725,7 +725,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui CAT(prefix, _OFFSET)) #define GET_FILTER_OS_IS_ZYX_OSA4_ISA8_OSV8_ISV4_INDEX(prefix, o, i, z, y, x) \ - FUNC_CALL(get_g_os_is_yx_osa4_isa8_osv8_isv4)( \ + get_g_os_is_yx_osa4_isa8_osv8_isv4( \ 0, o, i, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -735,7 +735,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui 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)( \ + get_g_os_is_yx_osa4_isa8_osv8_isv4( \ g, o, i, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -745,7 +745,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui CAT(prefix, _OFFSET)) #define GET_FILTER_G_OS_IS_ZYX_OSA4_ISA8_OSV8_ISV4_INDEX(prefix, g, o, i, z, y, x) \ - FUNC_CALL(get_g_os_is_yx_osa4_isa8_osv8_isv4)( \ + get_g_os_is_yx_osa4_isa8_osv8_isv4( \ g, o, i, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -755,7 +755,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui 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)( \ + get_g_os_is_yx_osa4_isa8_osv8_isv2( \ 0, o, i, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -765,7 +765,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui CAT(prefix, _OFFSET)) #define GET_FILTER_OS_IS_ZYX_OSA4_ISA8_OSV8_ISV2_INDEX(prefix, o, i, z, y, x) \ - FUNC_CALL(get_g_os_is_yx_osa4_isa8_osv8_isv2)( \ + get_g_os_is_yx_osa4_isa8_osv8_isv2( \ 0, o, i, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -775,7 +775,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui 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)( \ + get_g_os_is_yx_osa4_isa8_osv8_isv2( \ g, o, i, 0, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -785,7 +785,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui CAT(prefix, _OFFSET)) #define GET_FILTER_G_OS_IS_ZYX_OSA4_ISA8_OSV8_ISV2_INDEX(prefix, g, o, i, z, y, x) \ - FUNC_CALL(get_g_os_is_yx_osa4_isa8_osv8_isv2)( \ + get_g_os_is_yx_osa4_isa8_osv8_isv2( \ g, o, i, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -795,7 +795,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui CAT(prefix, _OFFSET)) #define GET_FILTER_OS_IS_YX_OSA2_ISA8_OSV16_ISV4_INDEX(prefix, o, i, y, x) \ - FUNC_CALL(get_g_os_is_yx_osa2_isa8_osv16_isv4)( \ + get_g_os_is_yx_osa2_isa8_osv16_isv4( \ 0, o, i, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -804,7 +804,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui CAT(prefix, _OFFSET)) #define GET_FILTER_G_OS_IS_YX_OSA2_ISA8_OSV16_ISV4_INDEX(prefix, g, o, i, y, x) \ - FUNC_CALL(get_g_os_is_yx_osa2_isa8_osv16_isv4)( \ + get_g_os_is_yx_osa2_isa8_osv16_isv4( \ g, o, i, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -813,7 +813,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui CAT(prefix, _OFFSET)) #define GET_FILTER_OS_IS_YX_OSA2_ISA8_OSV16_ISV2_INDEX(prefix, o, i, y, x) \ - FUNC_CALL(get_g_os_is_yx_osa2_isa8_osv16_isv2)( \ + get_g_os_is_yx_osa2_isa8_osv16_isv2( \ 0, o, i, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -822,7 +822,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui CAT(prefix, _OFFSET)) #define GET_FILTER_G_OS_IS_YX_OSA2_ISA8_OSV16_ISV2_INDEX(prefix, g, o, i, y, x) \ - FUNC_CALL(get_g_os_is_yx_osa2_isa8_osv16_isv2)( \ + get_g_os_is_yx_osa2_isa8_osv16_isv2( \ g, o, i, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -831,7 +831,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui 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)( \ + get_os_is_yx_isa8_osv8_isv4_swizzled_by_4_index( \ o, i, y, x, CAT(prefix, _SIZE_X ), \ CAT(prefix, _SIZE_Y), \ CAT(prefix, _IFM_NUM), \ @@ -839,7 +839,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui CAT(prefix, _OFFSET)) #define GET_FILTER_OS_IS_YX_OSA4_ISA8_OSV8_ISV4_SWIZZLED_BY_4_INDEX(prefix, o, i, y, x) \ - FUNC_CALL(get_os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4_index)( \ + get_os_is_yx_osa4_isa8_osv8_isv4_swizzled_by_4_index( \ o, i, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -848,7 +848,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui CAT(prefix, _OFFSET)) #define GET_FILTER_OS_IS_ZYX_OSA4_ISA8_OSV8_ISV4_SWIZZLED_BY_4_INDEX(prefix, o, i, z, y, x) \ - FUNC_CALL(get_os_is_zyx_osa4_isa8_osv8_isv4_swizzled_by_4_index)( \ + get_os_is_zyx_osa4_isa8_osv8_isv4_swizzled_by_4_index( \ o, i, z, y, x, \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y), \ @@ -858,7 +858,7 @@ inline uint FUNC(get_g_os_is_yx_osa2_isa8_osv16_isv2)(uint g, uint o, uint i, ui CAT(prefix, _OFFSET)) -inline uint FUNC(get_is_o_yx_isv32_index)(uint o, uint i, uint y, uint x, uint i_size, uint o_size, uint x_size, uint y_size) +inline uint get_is_o_yx_isv32_index(uint o, uint i, uint y, uint x, uint i_size, uint o_size, uint x_size, uint y_size) { const uint i_aligned_to_32 = ((i_size + 31) / 32) * 32; const uint i_val = i % 32; @@ -868,14 +868,14 @@ inline uint FUNC(get_is_o_yx_isv32_index)(uint o, uint i, uint y, uint x, uint i } #define GET_FILTER_IS_O_YX_ISV32(prefix, o, i, y, x) \ - FUNC_CALL(get_is_o_yx_isv32_index)( \ + get_is_o_yx_isv32_index( \ o, i, y, x, \ CAT(prefix, _IFM_NUM), \ CAT(prefix, _OFM_NUM), \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y)) -inline uint FUNC(get_is_o32_yx_isv32_swizzled_by_4_index)(uint o, uint i, uint y, uint x, uint i_size, uint o_size, uint x_size, uint y_size) +inline uint get_is_o32_yx_isv32_swizzled_by_4_index(uint o, uint i, uint y, uint x, uint i_size, uint o_size, uint x_size, uint y_size) { const uint o_aligned_to_32 = ((o_size + 31) / 32) * 32; const uint o_swizzled = (o % 4) * 8 + ((o % 32) / 4) + (o / 32) * 32; @@ -887,14 +887,14 @@ inline uint FUNC(get_is_o32_yx_isv32_swizzled_by_4_index)(uint o, uint i, uint y } #define GET_FILTER_IS_O32_YX_ISV32_SWIZZLED_BY_4(prefix, o, i, y, x) \ - FUNC_CALL(get_is_o32_yx_isv32_swizzled_by_4_index)( \ + get_is_o32_yx_isv32_swizzled_by_4_index( \ o, i, y, x, \ CAT(prefix, _IFM_NUM), \ CAT(prefix, _OFM_NUM), \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y)) -inline uint FUNC(get_os_is_y_x8_osv8_isv4_index)(uint o, uint i, uint y, uint x, uint i_size, uint o_size, uint x_size, uint y_size) +inline uint get_os_is_y_x8_osv8_isv4_index(uint o, uint i, uint y, uint x, uint i_size, uint o_size, uint x_size, uint y_size) { const uint i_aligned_to_4 = ((i_size + 3) / 4) * 4; const uint o_aligned_to_8 = ((o_size + 7) / 8) * 8; @@ -908,14 +908,14 @@ inline uint FUNC(get_os_is_y_x8_osv8_isv4_index)(uint o, uint i, uint y, uint x, } #define GET_FILTER_OS_IS_Y_X8_OSV8_ISV4(prefix, o, i, y, x) \ - FUNC_CALL(get_os_is_y_x8_osv8_isv4_index)( \ + get_os_is_y_x8_osv8_isv4_index( \ o, i, y, x, \ CAT(prefix, _IFM_NUM), \ CAT(prefix, _OFM_NUM), \ CAT(prefix, _SIZE_X), \ CAT(prefix, _SIZE_Y)) -inline uint FUNC(get_os_is_y_x8_osv8_isv4_swizzled_by_4_index)(uint o, uint i, uint y, uint x, uint i_size, uint o_size, uint x_size, uint y_size) +inline uint get_os_is_y_x8_osv8_isv4_swizzled_by_4_index(uint o, uint i, uint y, uint x, uint i_size, uint o_size, uint x_size, uint y_size) { const uint i_aligned_to_4 = ((i_size + 3) / 4) * 4; const uint o_aligned_to_8 = ((o_size + 7) / 8) * 8; @@ -930,7 +930,7 @@ inline uint FUNC(get_os_is_y_x8_osv8_isv4_swizzled_by_4_index)(uint o, uint i, u } #define GET_FILTER_OS_IS_Y_X8_OSV8_ISV4_SWIZZLED_BY_4(prefix, o, i, y, x) \ - FUNC_CALL(get_os_is_y_x8_osv8_isv4_swizzled_by_4_index)( \ + get_os_is_y_x8_osv8_isv4_swizzled_by_4_index( \ o, i, y, x, \ CAT(prefix, _IFM_NUM), \ CAT(prefix, _OFM_NUM), \ @@ -939,7 +939,7 @@ inline uint FUNC(get_os_is_y_x8_osv8_isv4_swizzled_by_4_index)(uint o, uint i, u #define GET_FILTER_G_OS_IS_YX_OSV16_ISV4_INDEX(prefix, g, o, i, y, x) \ - FUNC_CALL(get_g_os_is_yx_osv16_isv4)( \ + get_g_os_is_yx_osv16_isv4( \ g, o, i, y, x, \ CAT(prefix, _IFM_PITCH), \ CAT(prefix, _OFM_PITCH), \ @@ -948,7 +948,7 @@ inline uint FUNC(get_os_is_y_x8_osv8_isv4_swizzled_by_4_index)(uint o, uint i, u CAT(prefix, _OFM_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, +inline uint get_g_os_is_yx_osv16_isv4(uint g, uint o, uint i, uint y, uint x, uint i_size, uint o_size, uint x_size, @@ -978,7 +978,7 @@ inline uint FUNC(get_g_os_is_yx_osv16_isv4)(uint g, uint o, uint i, uint y, uint } #define GET_FILTER_OS_IS_YX_OSV8_ISV2_INDEX(prefix, o, i, y, x) \ - FUNC_CALL(get_g_os_is_yx_osv16_isv4)( \ + get_g_os_is_yx_osv16_isv4( \ 0, o, i, y, x, \ CAT(prefix, _IFM_PITCH), \ CAT(prefix, _OFM_PITCH), \ @@ -988,7 +988,7 @@ inline uint FUNC(get_g_os_is_yx_osv16_isv4)(uint g, uint o, uint i, uint y, uint 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)( \ + get_g_os_is_yx_osv16_isv4( \ 0, o, i, y, x, \ CAT(prefix, _IFM_PITCH), \ CAT(prefix, _OFM_PITCH), \ @@ -998,7 +998,7 @@ inline uint FUNC(get_g_os_is_yx_osv16_isv4)(uint g, uint o, uint i, uint y, uint 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)( \ + get_os_is_zyx_osv_isv4( \ o, i, 0, y, x, \ CAT(prefix, _IFM_PITCH), \ CAT(prefix, _OFM_PITCH), \ @@ -1007,7 +1007,7 @@ inline uint FUNC(get_g_os_is_yx_osv16_isv4)(uint g, uint o, uint i, uint y, uint 16) #define GET_FILTER_OS_IS_YX_OSV32_ISV4_INDEX(prefix, o, i, y, x) \ - FUNC_CALL(get_os_is_zyx_osv_isv4)( \ + get_os_is_zyx_osv_isv4( \ o, i, 0, y, x, \ CAT(prefix, _IFM_PITCH), \ CAT(prefix, _OFM_PITCH), \ @@ -1016,7 +1016,7 @@ inline uint FUNC(get_g_os_is_yx_osv16_isv4)(uint g, uint o, uint i, uint y, uint 32) #define GET_FILTER_OS_IS_ZYX_OSV32_ISV4_INDEX(prefix, o, i, z, y, x) \ - FUNC_CALL(get_os_is_zyx_osv_isv4)( \ + get_os_is_zyx_osv_isv4( \ o, i, z, y, x, \ CAT(prefix, _IFM_PITCH), \ CAT(prefix, _OFM_PITCH), \ @@ -1024,7 +1024,7 @@ inline uint FUNC(get_g_os_is_yx_osv16_isv4)(uint g, uint o, uint i, uint y, uint CAT(prefix, _SIZE_Y), \ 32) -inline uint FUNC(get_os_is_zyx_osv_isv4)(uint o, uint i, uint z, uint y, uint x, +inline uint get_os_is_zyx_osv_isv4(uint o, uint i, uint z, uint y, uint x, uint i_size, uint o_size, uint x_size, @@ -1050,14 +1050,14 @@ inline uint FUNC(get_os_is_zyx_osv_isv4)(uint o, uint i, uint z, uint y, uint x, } #define GET_FILTER_OS_IS_YX_OSV32_ISV4_SWIZZLED_BY_2_INDEX(prefix, o, i, y, x) \ - FUNC_CALL(get_os_is_yx_osv32_isv4_swizzled_by_2)( \ + get_os_is_yx_osv32_isv4_swizzled_by_2( \ o, i, y, x, \ CAT(prefix, _OFM_NUM), \ CAT(prefix, _IFM_NUM), \ CAT(prefix, _SIZE_Y), \ CAT(prefix, _SIZE_X)) -inline uint FUNC(get_os_is_yx_osv32_isv4_swizzled_by_2)(uint o, uint i, uint y, uint x, +inline uint get_os_is_yx_osv32_isv4_swizzled_by_2(uint o, uint i, uint y, uint x, uint o_size, uint i_size, uint y_size, @@ -1084,7 +1084,7 @@ inline uint FUNC(get_os_is_yx_osv32_isv4_swizzled_by_2)(uint o, uint i, uint y, return idx; } -inline uint FUNC(get_os_is_osv32_isv32_swizzled_by_4_index)(uint o, uint i, uint y, uint x, uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset) +inline uint get_os_is_osv32_isv32_swizzled_by_4_index(uint o, uint i, uint y, uint x, uint size_x, uint size_y, uint size_ifm, uint size_ofm, uint offset) { const uint size_ifm_a = ((size_ifm + 31)/32) * 32; @@ -1106,14 +1106,14 @@ inline uint FUNC(get_os_is_osv32_isv32_swizzled_by_4_index)(uint o, uint i, uint } #define GET_FILTER_OS_IS_OSV32_ISV32_SWIZZLED_BY_4_INDEX(prefix, o, i, y, x)\ - FUNC_CALL(get_os_is_osv32_isv32_swizzled_by_4_index)(\ + get_os_is_osv32_isv32_swizzled_by_4_index(\ o, i, y, x, CAT(prefix, _SIZE_X ),\ CAT(prefix, _SIZE_Y),\ CAT(prefix, _IFM_NUM),\ CAT(prefix, _OFM_NUM),\ CAT(prefix, _OFFSET)) -inline uint FUNC(get_os_i_yxs_osv_yxsv4_index)(uint o, uint i, uint y, uint x, uint i_size, uint size_x, uint size_y, uint osv) { +inline uint get_os_i_yxs_osv_yxsv4_index(uint o, uint i, uint y, uint x, uint i_size, uint size_x, uint size_y, uint osv) { const uint yxsv = 4; uint yx = y * size_x + x; uint yx_size_aligned = (size_x * size_y + yxsv - 1) / yxsv * yxsv; @@ -1132,7 +1132,7 @@ inline uint FUNC(get_os_i_yxs_osv_yxsv4_index)(uint o, uint i, uint y, uint x, u } #define GET_FILTER_OS_I_YXS_OSV4_YXSV4_INDEX(prefix, o, i, y, x) \ - FUNC_CALL(get_os_i_yxs_osv_yxsv4_index)( \ + get_os_i_yxs_osv_yxsv4_index( \ o, i, y, x, \ CAT(prefix, _IFM_NUM), \ CAT(prefix, _SIZE_X), \ @@ -1233,7 +1233,7 @@ inline uint FUNC(get_os_i_yxs_osv_yxsv4_index)(uint o, uint i, uint y, uint x, u ) #define GET_FILTER_GI_YXS_OS_YXSV2_OSV_INDEX(prefix, g, o, i, y, x, sub_group_size) \ - FUNC_CALL(get_gi_yxs_os_yxsv2_osv_index)( \ + get_gi_yxs_os_yxsv2_osv_index( \ g, o, i, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _GROUPS_PITCH), \ @@ -1244,7 +1244,7 @@ inline uint FUNC(get_os_i_yxs_osv_yxsv4_index)(uint o, uint i, uint y, uint x, u sub_group_size) #define GET_FILTER_GIY_XS_OS_XSV2_OSV_INDEX(prefix, g, o, i, y, x, sub_group_size) \ - FUNC_CALL(get_giy_xs_os_xsv2_osv_index)( \ + get_giy_xs_os_xsv2_osv_index( \ g, o, i, y, x, \ CAT(prefix, _SIZE_X ), \ CAT(prefix, _GROUPS_PITCH), \ @@ -1254,7 +1254,7 @@ inline uint FUNC(get_os_i_yxs_osv_yxsv4_index)(uint o, uint i, uint y, uint x, u CAT(prefix, _OFFSET), \ sub_group_size) -inline uint FUNC(get_gs_oi_yxs_gsv_yxsv4_index)(uint g, uint o, uint i, uint y, uint x, uint o_size, uint i_size, uint size_x, uint size_y, const uint gsv) { +inline uint get_gs_oi_yxs_gsv_yxsv4_index(uint g, uint o, uint i, uint y, uint x, uint o_size, uint i_size, uint size_x, uint size_y, const uint gsv) { const uint yxsv = 4; uint yx = y * size_x + x; uint yx_size_aligned = (size_x * size_y + yxsv - 1) / yxsv * yxsv; @@ -1273,7 +1273,7 @@ inline uint FUNC(get_gs_oi_yxs_gsv_yxsv4_index)(uint g, uint o, uint i, uint y, } #define GET_FILTER_GS_OI_YXS_GSV4_YXSV4_INDEX(prefix, g, o, i, y, x) \ - FUNC_CALL(get_gs_oi_yxs_gsv_yxsv4_index)( \ + get_gs_oi_yxs_gsv_yxsv4_index( \ g, o, i, y, x, \ CAT(prefix, _OFM_NUM), \ CAT(prefix, _IFM_NUM), \ @@ -1282,7 +1282,7 @@ inline uint FUNC(get_gs_oi_yxs_gsv_yxsv4_index)(uint g, uint o, uint i, uint y, 4) #define GET_FILTER_GS_OI_YXS_GSV16_YXSV4_INDEX(prefix, g, o, i, y, x) \ - FUNC_CALL(get_gs_oi_yxs_gsv_yxsv4_index)( \ + get_gs_oi_yxs_gsv_yxsv4_index( \ g, o, i, y, x, \ CAT(prefix, _OFM_NUM), \ CAT(prefix, _IFM_NUM), \ @@ -1291,7 +1291,7 @@ inline uint FUNC(get_gs_oi_yxs_gsv_yxsv4_index)(uint g, uint o, uint i, uint y, 16) #define GET_FILTER_GS_OI_YXS_GSV32_YXSV4_INDEX(prefix, g, o, i, y, x) \ - FUNC_CALL(get_gs_oi_yxs_gsv_yxsv4_index)( \ + get_gs_oi_yxs_gsv_yxsv4_index( \ g, o, i, y, x, \ CAT(prefix, _OFM_NUM), \ CAT(prefix, _IFM_NUM), \ @@ -1311,7 +1311,7 @@ inline uint FUNC(get_gs_oi_yxs_gsv_yxsv4_index)(uint g, uint o, uint i, uint y, ((o) / (sub_group_size))*CAT(prefix, _OFM_PITCH) \ ) -inline uint FUNC(get_g_os_zyx_is_osv_isv_index)(uint g, uint o, uint i, uint z, uint y, uint x, +inline uint get_g_os_zyx_is_osv_isv_index(uint g, uint o, uint i, uint z, uint y, uint x, uint g_size, uint o_size, uint i_size, uint z_size, uint y_size, uint x_size, uint osv, uint isv) { uint is_size = (i_size + isv - 1) / isv; @@ -1344,7 +1344,7 @@ 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_OSV_ISV_INDEX(tensor, g, o, i, z, y, x, osv, isv) \ - FUNC_CALL(get_g_os_zyx_is_osv_isv_index)( \ + get_g_os_zyx_is_osv_isv_index( \ g, o, i, z, y, x, \ CAT(tensor, _GROUPS_NUM), \ CAT(tensor, _OFM_NUM), \ diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/vec_typedefs.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/vec_typedefs.cl similarity index 100% rename from inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/vec_typedefs.cl rename to inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/batch_headers/vec_typedefs.cl diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/image_data.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/image_data.cl new file mode 100644 index 00000000000..3fe5486bc3a --- /dev/null +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/image_data.cl @@ -0,0 +1,9 @@ +#define DECLARE_SAMPLER const sampler_t imageSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST + +#if FP16_UNIT_USED + #define IMAGE_READ(image, coord) read_imageh((image), imageSampler, (coord)) + #define IMAGE_WRITE(image, coord, val) write_imageh((image), (coord), (val)) +#else + #define IMAGE_READ(image, coord) read_imagef((image), imageSampler, (coord)) + #define IMAGE_WRITE(image, coord, val) write_imagef((image), (coord), (val)) +#endif diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/unit_type.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/unit_type.cl index fba62c463e6..17c86158fcf 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/unit_type.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/include/unit_type.cl @@ -31,7 +31,6 @@ #define UNIT_BLOCK_RW_TYPE4 MAKE_VECTOR_TYPE(UNIT_BLOCK_RW_TYPE, 4) #define UNIT_BLOCK_RW_TYPE8 MAKE_VECTOR_TYPE(UNIT_BLOCK_RW_TYPE, 8) -#define AS_TYPE(type, val) CAT(as_, type)(val) #define UNIT_BLOCK_READ_FUNC2 CAT(UNIT_BLOCK_READ_FUNC, 2) #define UNIT_BLOCK_READ_FUNC4 CAT(UNIT_BLOCK_READ_FUNC, 4) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_multiple_features.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_multiple_features.cl index 5b99aa970a1..7b4f6f624bf 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_multiple_features.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_multiple_features.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #ifdef FORCE_SIMD_16 __attribute__((intel_reqd_sub_group_size(16))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_multiple_features_fsv16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_multiple_features_fsv16.cl index afa7bbac5b0..37582f03607 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_multiple_features_fsv16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_multiple_features_fsv16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" __attribute__((intel_reqd_sub_group_size(16))) KERNEL (lrn_gpu_across_channel_multiple_features_fsv16)( diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_ref.cl index 4aa19b9223c..452e1baa18d 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" KERNEL (lrn_gpu_across_channel_ref)( const __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_yxfb_b8_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_yxfb_b8_opt.cl index c987d0ea495..d45f063a8b2 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_yxfb_b8_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_across_channel_yxfb_b8_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define INPUT_VECTOR_TYPE MAKE_VECTOR_TYPE(INPUT0_TYPE, 8) #define OUTPUT_VECTOR_TYPE MAKE_VECTOR_TYPE(OUTPUT_TYPE, 8) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_within_channel.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_within_channel.cl index b82b47b32d2..bccfd745880 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_within_channel.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_within_channel.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" KERNEL (lrn_gpu_within_channel)( const __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_within_channel_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_within_channel_opt.cl index 8c8da428e35..1e91956d032 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_within_channel_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_gpu_within_channel_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" KERNEL (lrn_gpu_within_channel_opt)( const __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_ref.cl index ae984b5ffed..71004b94568 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_ref.cl @@ -2,8 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/acc_type.cl" KERNEL(normalization)( __global const INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_within_channel_byxf_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_within_channel_byxf_opt.cl index 7f43639d675..e2b22d3738f 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_within_channel_byxf_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lrn_within_channel_byxf_opt.cl @@ -2,8 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/acc_type.cl" #define VECTOR_TYPE MAKE_VECTOR_TYPE(INPUT0_TYPE, 8) #define ACCUMULATOR_VECTOR_TYPE MAKE_VECTOR_TYPE(INPUT0_TYPE, 8) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_input_bfyx_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_input_bfyx_opt.cl index d054e349f9b..e6a18a56762 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_input_bfyx_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_input_bfyx_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/unit_type.cl" #include "include/sub_group.cl" diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_input_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_input_ref.cl index 8031c92a97e..712d853755f 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_input_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_input_ref.cl @@ -2,9 +2,10 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/fetch_weights.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/fetch_weights.cl" +#include "include/acc_type.cl" KERNEL(lstm_dynamic_input_ref)( const __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_timeloop_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_timeloop_ref.cl index 2df423de9b9..2abd8b54507 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_timeloop_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_dynamic_timeloop_ref.cl @@ -2,8 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/acc_type.cl" #define ACTIVATION_LOGISTIC(input) (UNIT_VAL_ONE/(UNIT_VAL_ONE + exp(-input))) #define ACTIVATION_HYPERBOLIC_TAN(input) (tanh(input)) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_elt_gpu_bfyx_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_elt_gpu_bfyx_ref.cl index 76dc5fc4803..3bb468ba1b3 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_elt_gpu_bfyx_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_elt_gpu_bfyx_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" // tempGEMM = [ batch, 1, direction, 4 * hidden_size ] // cell = [ batch, 1, direction, hidden_size ] optional diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemm_gpu_bfyx_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemm_gpu_bfyx_ref.cl index 3eb0536b936..9d08927ac04 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemm_gpu_bfyx_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemm_gpu_bfyx_ref.cl @@ -2,8 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/acc_type.cl" #ifndef DIRECTION #define DIRECTION 0 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemv_gpu_subgroup1x64_bfyx_ff_SIMD16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemv_gpu_subgroup1x64_bfyx_ff_SIMD16.cl index b88359c0ac8..c70eb5180ec 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemv_gpu_subgroup1x64_bfyx_ff_SIMD16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemv_gpu_subgroup1x64_bfyx_ff_SIMD16.cl @@ -2,8 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/acc_type.cl" #ifndef DIRECTION #define DIRECTION 0 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemv_gpu_subgroup1x64_bfyx_hh_SIMD16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemv_gpu_subgroup1x64_bfyx_hh_SIMD16.cl index 86f5aba2766..60b926eaf3b 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemv_gpu_subgroup1x64_bfyx_hh_SIMD16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/lstm_gemv_gpu_subgroup1x64_bfyx_hh_SIMD16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #ifndef DIRECTION #define DIRECTION 0 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/max_unpooling_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/max_unpooling_gpu_ref.cl index 46bbe583be1..c0ee49c7066 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/max_unpooling_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/max_unpooling_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(pooling_gpu)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output, const __global float* arg_max) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad.cl index dc77da647fd..d5c1600dfb7 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad.cl @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/imad.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" #include "mvn_gpu_b_fs_yx_fsv16_imad_accumulate.cl" #include "mvn_gpu_b_fs_yx_fsv16_imad_reduce.cl" diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad_accumulate.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad_accumulate.cl index 7ae45dd5f0d..39c463c3f50 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad_accumulate.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad_accumulate.cl @@ -2,7 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" // ============================================================================================================================== // DECLARE_PACKED_ACCUMULATE(Name, AccT, InputT, SliceSize, SlicePitch, Items, Workers, AccOp) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad_reduce.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad_reduce.cl index dafd6e47b37..456403a6112 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad_reduce.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_b_fs_yx_fsv16_imad_reduce.cl @@ -2,7 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" +#include "include/batch_headers/common.cl" // ============================================================================================================================== // DECLARE_SG_PACKED_REDUCE_ADD(Name, Type, VecSize, PostOp) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_bfyx_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_bfyx_opt.cl index 1dc2754def2..f9888834c2d 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_bfyx_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_bfyx_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" __attribute__((reqd_work_group_size(LWS, 1, 1))) KERNEL (mvn_gpu_bfyx_opt)( diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_ref_across_channels.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_ref_across_channels.cl index 4d740781931..1d361e677a4 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_ref_across_channels.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_ref_across_channels.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" KERNEL (mvn_gpu_ref_across_channels)( diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_ref_within_channels.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_ref_within_channels.cl index 77201027a65..b621479e159 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_ref_within_channels.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/mvn_gpu_ref_within_channels.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" KERNEL (mvn_gpu_ref_within_channels)( const __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/normalize_gpu_across_spatial_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/normalize_gpu_across_spatial_ref.cl index 2e8e23a0c77..243d78767c8 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/normalize_gpu_across_spatial_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/normalize_gpu_across_spatial_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" KERNEL (normalize_gpu_across_spatial_bfyx)( const __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/normalize_gpu_within_spatial_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/normalize_gpu_within_spatial_ref.cl index 757b49ac898..2da82bef003 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/normalize_gpu_within_spatial_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/normalize_gpu_within_spatial_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" KERNEL (normalize_gpu_within_spatial_bfyx)( const __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/one_hot_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/one_hot_ref.cl index 748f0276d4f..5afaac0c8d7 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/one_hot_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/one_hot_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #if OUTPUT_LAYOUT_BFZYX #define GET_COORDS_INDEX(prefix, coords) GET_DATA_INDEX_5D(prefix, coords[0], coords[1], coords[2], coords[3], coords[4]) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_ref.cl index e51affa0037..ebeb81d038e 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL (permute_ref)( const __global INPUT0_TYPE* input, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_tile_8x8_4x4.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_tile_8x8_4x4.cl index 9b190a27000..56214ba3395 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_tile_8x8_4x4.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_tile_8x8_4x4.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for #define CEIL_DIV(A, B) (((A) + (B) - 1) / (B)) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_tile_8x8_4x4_fsv.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_tile_8x8_4x4_fsv.cl index 2dc70de53c4..29e29262e22 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_tile_8x8_4x4_fsv.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/permute_tile_8x8_4x4_fsv.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for #define CEIL_DIV(A, B) (((A) + (B) - 1) / (B)) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_average_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_average_opt.cl index 0776269007d..a6f0f98dcd9 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_average_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_average_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" __attribute__((intel_reqd_sub_group_size(SUB_GROUP_SIZE))) __attribute__((reqd_work_group_size(SUB_GROUP_SIZE, 1, 1))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_b_fs_yx_fsv4.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_b_fs_yx_fsv4.cl index e1fe673115f..4580d01b3ea 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_b_fs_yx_fsv4.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_b_fs_yx_fsv4.cl @@ -2,13 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #define ALIGN_TO(val, multiple) (((val) + (multiple) - 1) / (multiple) * (multiple)) -#define AS_TYPE(type, val) CAT(as_, type)(val) #define INPUT_VEC4 MAKE_VECTOR_TYPE(INPUT0_TYPE, 4) #define ACTIVATION_VEC4 MAKE_VECTOR_TYPE(ACTIVATION_TYPE, 4) @@ -183,7 +182,6 @@ KERNEL(pooling_gpu_b_fs_yx_fsv4)( } #undef ALIGN_TO -#undef AS_TYPE #undef INIT_VAL #undef INPUT_VEC4 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_b_fs_zyx_fsv16_imad.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_b_fs_zyx_fsv16_imad.cl index 17af51c8137..ac583a5a954 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_b_fs_zyx_fsv16_imad.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_b_fs_zyx_fsv16_imad.cl @@ -2,13 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #define ALIGN_TO(val, multiple) (((val) + (multiple) - 1) / (multiple) * (multiple)) -#define AS_TYPE(type, val) CAT(as_, type)(val) #define IN_VEC16 MAKE_VECTOR_TYPE(INPUT0_TYPE, 16) #define OUT_VEC16 MAKE_VECTOR_TYPE(OUTPUT_TYPE, 16) @@ -466,7 +465,6 @@ KERNEL(pooling_gpu_b_fs_zyx_fsv16)( #endif // GLOBAL_POOLING #undef ALIGN_TO -#undef AS_TYPE #undef IN_VEC16 #undef OUT_VEC16 #undef ACTIVATION_VEC16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bfyx_block_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bfyx_block_opt.cl index 2dd6a2fa9c4..2c670930e25 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bfyx_block_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bfyx_block_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #if MAX_POOLING || MAX_WITH_ARGMAX_POOLING #define INIT_VAL ACCUMULATOR_VAL_MIN diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_blocked.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_blocked.cl index c2db898dc0c..b6395ece5f8 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_blocked.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_blocked.cl @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #define FEATURE_SLICE_SIZE 16 #if X_BLOCK_SIZE > 1 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bs_fs_yx_bsv16_fsv16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bs_fs_yx_bsv16_fsv16.cl index 3ed786fdd12..71ca0314bca 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bs_fs_yx_bsv16_fsv16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bs_fs_yx_bsv16_fsv16.cl @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #define ALIGN_TO(val, multiple) (((val) + (multiple)-1) / (multiple) * (multiple)) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bsv16_fsv16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bsv16_fsv16.cl index b95b82b7385..58ba7e5bd4a 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bsv16_fsv16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_bsv16_fsv16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define INPUT0_SIZE_X_WITH_PADDING (INPUT0_PAD_BEFORE_SIZE_X + INPUT0_SIZE_X + INPUT0_PAD_AFTER_SIZE_X) #define INPUT0_SIZE_Y_WITH_PADDING (INPUT0_PAD_BEFORE_SIZE_Y + INPUT0_SIZE_Y + INPUT0_PAD_AFTER_SIZE_Y) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_byxf_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_byxf_opt.cl index ac403bc8c45..51685b581f9 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_byxf_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_byxf_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define INPUT_VEC8 MAKE_VECTOR_TYPE(INPUT0_TYPE, 8) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_byxf_padding_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_byxf_padding_opt.cl index cc8aa6929a2..15c32b7d3f7 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_byxf_padding_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_byxf_padding_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define INPUT0_VEC8 MAKE_VECTOR_TYPE(INPUT0_TYPE,8) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_fs_b_yx_fsv32.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_fs_b_yx_fsv32.cl index 8ae4a593f1f..d98c6a49bed 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_fs_b_yx_fsv32.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_fs_b_yx_fsv32.cl @@ -2,9 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #if MAX_POOLING #define INIT_VAL ACCUMULATOR_VAL_MIN diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_int8_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_int8_ref.cl index 4a44d357c83..456d615b191 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_int8_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_int8_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #if MAX_POOLING #define INIT_VAL ACCUMULATOR_VAL_MIN diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_ref.cl index 3994bed38e8..6d2d3a7c123 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pooling_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #if MAX_POOLING || MAX_WITH_ARGMAX_POOLING #define INIT_VAL ACCUMULATOR_VAL_MIN diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pyramid_roi_align_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pyramid_roi_align_gpu_ref.cl index b05429457ff..694d79d64c5 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pyramid_roi_align_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/pyramid_roi_align_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define PYRAMID_LEVELS 4 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/quantize_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/quantize_gpu_ref.cl index 1a1a8103055..1bdcefc634c 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/quantize_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/quantize_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #ifdef SUB_GROUP_SIZE __attribute__((intel_reqd_sub_group_size(SUB_GROUP_SIZE))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/quantize_gpu_scale_shift_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/quantize_gpu_scale_shift_opt.cl index b0239a42b84..5b5517db571 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/quantize_gpu_scale_shift_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/quantize_gpu_scale_shift_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #ifdef SUB_GROUP_SIZE __attribute__((intel_reqd_sub_group_size(SUB_GROUP_SIZE))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reduce_gpu_b_fs_yx_fsv16.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reduce_gpu_b_fs_yx_fsv16.cl index f6714deb175..148e379db66 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reduce_gpu_b_fs_yx_fsv16.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reduce_gpu_b_fs_yx_fsv16.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define SIMD 16 #define FSV 16 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reduce_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reduce_ref.cl index 4e202173f41..8f87158f3ac 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reduce_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reduce_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" inline uint FUNC(calc_linear_offset)(uint b, uint f, uint w, uint z, uint y, uint x) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/region_yolo_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/region_yolo_gpu_ref.cl index 0b927017e88..72b231f3c8b 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/region_yolo_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/region_yolo_gpu_ref.cl @@ -2,7 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" inline INPUT0_TYPE FUNC(logistic_activate)(INPUT0_TYPE x) { return 1. / (1. + exp(-x)); diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_biplanar_nv12.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_biplanar_nv12.cl index 2b6006ed8e4..b4cda3999bc 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_biplanar_nv12.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_biplanar_nv12.cl @@ -3,9 +3,9 @@ // #include "include/reshape_dims.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" ///////////////////////// Output Index ///////////////////////// inline uint FUNC(get_output_index)(uint b, uint f, uint w, uint z, uint y, uint x) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data.cl index 2280a8df61b..a799cf199f5 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data.cl @@ -3,9 +3,10 @@ // #include "include/reshape_dims.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" +#include "include/image_data.cl" #define INPUT_TYPE4 MAKE_VECTOR_TYPE(INPUT_REORDER_TYPE, 4) #define OUTPUT_TYPE4 MAKE_VECTOR_TYPE(OUTPUT_REORDER_TYPE, 4) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_b_fs_yx_fsv16_fsv32_to_bfyx.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_b_fs_yx_fsv16_fsv32_to_bfyx.cl index 9e8238eb1a7..9a43a711058 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_b_fs_yx_fsv16_fsv32_to_bfyx.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_b_fs_yx_fsv16_fsv32_to_bfyx.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for #define CEIL_DIV(A, B) (((A) + (B) - 1) / (B)) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_bfyx_to_blocked_format.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_bfyx_to_blocked_format.cl index 6fb342adbfa..ac738a50aee 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_bfyx_to_blocked_format.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_bfyx_to_blocked_format.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for #define CEIL_DIV(A, B) (((A) + (B) - 1) / (B)) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_binary.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_binary.cl index 62216f86ed7..b7dbc4948dc 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_binary.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_binary.cl @@ -3,9 +3,9 @@ // #include "include/reshape_dims.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" #if !INPUT0_LAYOUT_BFYX && !INPUT0_LAYOUT_B_FS_YX_32FP #error "Data binary reorder: unsupported input layout" diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_fast_b1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_fast_b1.cl index d1a6fe71b0f..e5436b1f036 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_fast_b1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_fast_b1.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #include "include/reshape_dims.cl" diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_to_yxfb_batched.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_to_yxfb_batched.cl index e083ef2726a..c82d7c530a0 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_to_yxfb_batched.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data_to_yxfb_batched.cl @@ -3,9 +3,9 @@ // #include "include/reshape_dims.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" ///////////////////////// Input Index ///////////////////////// inline uint FUNC(get_input_index)(uint b, uint f, uint y, uint x) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_from_winograd_2x3_s1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_from_winograd_2x3_s1.cl index cad4b2f5491..0e9713f0008 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_from_winograd_2x3_s1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_from_winograd_2x3_s1.cl @@ -6,8 +6,8 @@ // Convert the results using the inverse F(2,3) Winograd transform. // -------------------------------------------------------------------------------------------------------------------------------- -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(reorder_from_winograd_2x3_s1)(global const UNIT_TYPE* input_winograd, global float* output) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_fs_b_yx_fsv32_to_bfyx.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_fs_b_yx_fsv32_to_bfyx.cl index 0f29beab02f..afb5d58c315 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_fs_b_yx_fsv32_to_bfyx.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_fs_b_yx_fsv32_to_bfyx.cl @@ -2,7 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" +#include "include/batch_headers/fetch_data.cl" #include "include/unit_type.cl" __attribute__((reqd_work_group_size(1, LWS1, 1))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights.cl index 3975722e2d3..782bfe72499 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights.cl @@ -2,9 +2,10 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_weights.cl" +#include "include/batch_headers/fetch_weights.cl" #include "include/reshape_dims.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" +#include "include/image_data.cl" #define DECLARE_SAMPLER const sampler_t imageSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_binary.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_binary.cl index a8a57293824..4447b65ddef 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_binary.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_binary.cl @@ -3,7 +3,7 @@ // #include "include/reshape_dims.cl" -#include "include/data_types.cl" +#include "include/batch_headers/data_types.cl" #define OFM_BLOCK_SIZE 32 #define IFM_PACK_SIZE 32 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_image_2d_c4_fyx_b.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_image_2d_c4_fyx_b.cl index a52220f4e79..fc395f3da27 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_image_2d_c4_fyx_b.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_image_2d_c4_fyx_b.cl @@ -2,8 +2,9 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/image_data.cl" KERNEL (reorder_weights_image_2d_c4_fyx_b)(const __global INPUT0_TYPE* input, write_only image2d_t output) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_image_winograd_6x3_s1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_image_winograd_6x3_s1.cl index 8df6e3d81d5..18577f276bb 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_image_winograd_6x3_s1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_image_winograd_6x3_s1.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(reorder_weights_image_winograd_6x3_s1)(const __global INPUT0_TYPE* input, write_only image2d_t output) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_opt.cl index a30c9da2124..952438491fe 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" INIT_INPUT0_INDEX_FUNC_HERE INIT_OUTPUT_INDEX_FUNC_HERE diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_winograd_2x3_s1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_winograd_2x3_s1.cl index 39134ffd1c9..f1b680edcc3 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_winograd_2x3_s1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_winograd_2x3_s1.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(reorder_weights_winograd_2x3_s1)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_winograd_6x3_s1.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_winograd_6x3_s1.cl index 6d1907d6c0f..098caacde79 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_winograd_6x3_s1.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights_winograd_6x3_s1.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(reorder_weights_winograd_6x3_s1)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorg_yolo_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorg_yolo_gpu_ref.cl index a72e5527d7a..c2324ca48cf 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorg_yolo_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorg_yolo_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" #if OUTPUT_LAYOUT_BFYX #define IW INPUT0_SIZES[0] diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/resample_opt.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/resample_opt.cl index 2c9382802b4..c96ee525a49 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/resample_opt.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/resample_opt.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" #define unroll_for __attribute__((opencl_unroll_hint)) for diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/resample_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/resample_ref.cl index 4d10116d352..5326a42e9d2 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/resample_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/resample_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/fetch_data.cl" -#include "include/data_types.cl" +#include "include/batch_headers/fetch_data.cl" +#include "include/batch_headers/data_types.cl" inline uint FUNC(get_input_index)(uint b, uint f, uint z, uint y, uint x) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reshape_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reshape_ref.cl index 064867f8c4f..07bacee13a7 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reshape_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reshape_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" KERNEL (reshape_ref)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reverse_sequence_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reverse_sequence_ref.cl index f3bc8738180..1376a2752f5 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reverse_sequence_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reverse_sequence_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(reverse_sequence_ref)(const __global INPUT0_TYPE* input, const __global INPUT1_TYPE* seq_lengths, __global OUTPUT_TYPE* output) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/roi_pooling_ps_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/roi_pooling_ps_ref.cl index df6bea32487..f8bf7d3d516 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/roi_pooling_ps_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/roi_pooling_ps_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" // Each RoI is described by 5 elements [batch_id xmin ymin xmax ymax] #define ROI_NUM_ELEMENTS 5 diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/roi_pooling_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/roi_pooling_ref.cl index 61f08bfe8b5..6b9775c8a4c 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/roi_pooling_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/roi_pooling_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" /**************************************************************************** diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_elements_update_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_elements_update_ref.cl index 96f187e8e07..c45f2d7b775 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_elements_update_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_elements_update_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define GET_UPDATES_INDEX(prefix, idx_order) CAT(prefix, _GET_INDEX)(idx_order) #define GET_OUTPUT_INDEX(idx_order) OUTPUT_GET_INDEX(idx_order) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_nd_update_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_nd_update_ref.cl index cad99cafed6..104a2b4fe09 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_nd_update_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_nd_update_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define GET_UPDATES_INDEX(prefix, idx_order) CAT(prefix, _GET_INDEX)(idx_order) #define GET_OUTPUT_INDEX(idx_order) OUTPUT_GET_INDEX(idx_order) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_update_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_update_ref.cl index 0536e310c5e..f3919a9ef04 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_update_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/scatter_update_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define AXIS_B (0) #define AXIS_F (1) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/select_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/select_gpu_ref.cl index b61720a6a7f..3b1c595ac92 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/select_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/select_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #define GET_INDEX(prefix, num) \ CAT(CAT(prefix, num), _OFFSET) + \ diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/shuffle_channels_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/shuffle_channels_ref.cl index 79a9636716e..ff0b9176fb9 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/shuffle_channels_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/shuffle_channels_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(shuffle_channels_ref)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_bf.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_bf.cl index 901b8bbe702..18df70be0ec 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_bf.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_bf.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" __attribute__((reqd_work_group_size(LWS, 1, 1))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_fb.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_fb.cl index a4e77dce152..1383fa61062 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_fb.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_fb.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/batch_headers/data_types.cl" UNIT_TYPE FUNC(find_max_value)(__local UNIT_TYPE* partial_max, const int global_id, const int idx, const int batch_offset, const int data_sets_count, const __global UNIT_TYPE* input) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_items_class_optimized.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_items_class_optimized.cl index cb75d83e9e2..625e2aa050e 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_items_class_optimized.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_items_class_optimized.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/batch_headers/common.cl" +#include "include/acc_type.cl" #define DATA_PER_WORKITEM ( (INPUT0_CLASS_NUM + (WORKITEMS_PER_CLASSES - 1) ) / WORKITEMS_PER_CLASSES) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_ref.cl index c0f40bf0977..4a0e89c5a1e 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/softmax_gpu_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/common.cl" -#include "include/data_types.cl" +#include "include/acc_type.cl" +#include "include/batch_headers/common.cl" __attribute__((intel_reqd_sub_group_size(16))) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/space_to_batch_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/space_to_batch_ref.cl index e46efbb3a51..719688172aa 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/space_to_batch_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/space_to_batch_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(space_to_batch_ref)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/space_to_depth_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/space_to_depth_ref.cl index bebd431fcf2..0beabae108c 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/space_to_depth_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/space_to_depth_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" #if OUTPUT_DIMS == 5 #define SPATIAL_BLOCK_SIZE (BLOCK_SIZE*BLOCK_SIZE*BLOCK_SIZE) diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/strided_slice_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/strided_slice_ref.cl index fa019712bbe..e69634e6be2 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/strided_slice_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/strided_slice_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(strided_slice_ref)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/tile_ref.cl b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/tile_ref.cl index e7ad03fc514..e08247a270b 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/tile_ref.cl +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/tile_ref.cl @@ -2,8 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 // -#include "include/data_types.cl" -#include "include/fetch_data.cl" +#include "include/batch_headers/data_types.cl" +#include "include/batch_headers/fetch_data.cl" KERNEL(tile_ref)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output) { diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/common/jitter.cpp b/inference-engine/thirdparty/clDNN/kernel_selector/core/common/jitter.cpp index dc9d7b8dbf4..6b7244a6770 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/common/jitter.cpp +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/common/jitter.cpp @@ -484,6 +484,8 @@ class WeightTensorJitConstant : public TensorBaseTJitConstant; if (l == WeightsLayout::oiyx || diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db.cpp b/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db.cpp index c17d0898630..08d7dd382d9 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db.cpp +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db.cpp @@ -21,6 +21,9 @@ namespace cache { primitive_db::primitive_db() : primitives({ #include "ks_primitive_db.inc" + }), + batch_header_str({ +#include "ks_primitive_db_batch_headers.inc" }) { } @@ -58,7 +61,6 @@ std::vector primitive_db::get(const primitive_id& id) const { throw std::runtime_error("cannot find the kernel " + id + " in primitive database."); } } - } // namespace cache } // namespace gpu } // namespace kernel_selector diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db.h b/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db.h index d1f2be60761..04ba2d1e8dd 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db.h +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db.h @@ -22,6 +22,7 @@ struct primitive_db { primitive_db(); std::vector get(const primitive_id& id) const; + std::vector get_batch_header_str() const { return std::move(batch_header_str); } private: struct case_insensitive_compare { @@ -34,6 +35,7 @@ private: } }; std::multimap primitives; + std::vector batch_header_str; }; } // namespace cache diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db_gen.py b/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db_gen.py index 1052d9b1fa9..69c544298cd 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db_gen.py +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/common/primitive_db_gen.py @@ -1,8 +1,6 @@ #!/usr/bin/python3 - # Copyright (C) 2018-2021 Intel Corporation # SPDX-License-Identifier: Apache-2.0 - # To add new kernel please add a .cl file to kernels directory # the database name will be the part of the file name up to first '.' character # the trailing characters are a tag to allow multiple primitive implementations @@ -16,11 +14,45 @@ import re class OpenCL2CHeaders(object): - def __init__(self, kernels_folder, out_path, out_file_name): + def __init__(self, kernels_folder, out_path, out_file_name_prim_db, out_file_name_batch_headers): self.kernels_folder = os.path.abspath(kernels_folder) self.out_path = os.path.abspath(out_path) - self.out_file_name = out_file_name + self.out_file_name_prim_db = out_file_name_prim_db + self.out_file_name_batch_headers = out_file_name_batch_headers self.include_files = {} + self.batch_headers = [] + self.find_and_set_batch_headers() + + # NOTE: batch_headers are headers with macros on which the runtime jitter might depend on. + # For example, fetch_data.cl defines GET_DATA_INDEX which is to be used to define macros generated by jitter. + # These headers contains generally used macros and located under cl_kernels/include/batch_headers to be handled + # specially for improving the jit compilation performance, i.e., + # they are not to be included in each kernel, but to be included only once at the beginning of each batch. + def find_and_set_batch_headers(self): + batch_headers_list = [ntpath.basename(h) for h in glob.glob(os.path.join(self.kernels_folder, "include/batch_headers/*.cl"))] + deps = {} + for h in batch_headers_list: + header_file = os.path.abspath(os.path.join(self.kernels_folder, "include/batch_headers", h)) + f = open(header_file) + content = f.readlines() + deps[h] = {h} + for line in content: + if line.startswith('#include'): + include_file_name = line.strip().split('"')[1].strip() + deps[h].add(include_file_name) + else: + continue + while deps: + self.topological_sort(next(iter(deps)), deps, [], self.batch_headers) + + def topological_sort(self, cur_key, items, stack, res): + stack.append(cur_key) + deps = [dep for dep in items[cur_key] if dep not in stack and dep not in res] + for dep in deps: + self.topological_sort(dep, items, stack, res) + res.append(cur_key) + items.pop(cur_key) + stack.pop() def convert(self): res = '// This file is autogenerated by primitive_db_gen.py, all changes to this file will be undone\n\n' @@ -31,15 +63,17 @@ class OpenCL2CHeaders(object): res += self.cl_file_to_str(filename) #except: # pass - - out_file_name = os.path.join(self.out_path, self.out_file_name) - #with open(out_file_name, 'r') as out_file: - # old_content = out_file.read() - #if old_content != res: - #print('Replacing old DB') - with open(out_file_name, 'w') as out_file: + out_file_name_prim_db = os.path.join(self.out_path, self.out_file_name_prim_db) + with open(out_file_name_prim_db, 'w') as out_file: out_file.write(res) + # write batch_header_str + batch_headers = '// This file is autogenerated by primitive_db_gen.py, all changes to this file will be undone\n\n' + batch_headers += self.batch_headers_to_str() + out_file_name_batch_headers = os.path.join(self.out_path, self.out_file_name_batch_headers) + with open(out_file_name_batch_headers, 'w') as out_file: + out_file.write(batch_headers) + def append_undefs(self, filename): undefs = "" content = [] @@ -62,14 +96,83 @@ class OpenCL2CHeaders(object): undefs += include_file_undefs return undefs + # Conservative removal of macro checking with its potential users + # Potential users are determined by checking followings: + # - if macro is appearing on the line + # - if macro might be composed by concat + def found_concat_user(self, words, start_idx, macro): + potential_macro_user_exist = False + concat_len = 0 + iter_idx = start_idx + len_str1 = len_str2 = 0 + if words[iter_idx + 2] == "CAT": + user_exist, len_str1 = self.found_concat_user(words, iter_idx + 2, macro) + potential_macro_user_exist |= user_exist + else: + if macro.find(words[iter_idx + 2]) >= 0 : + potential_macro_user_exist = True + len_str1 = 1 + + if words[iter_idx + 3 + len_str1] == "CAT": + user_exist, len_str2 = self.found_concat_user(words, iter_idx + 3, macro) + potential_macro_user_exist |= user_exist + else: + if macro.find(words[iter_idx + 3 + len_str1]) >= 0: + potential_macro_user_exist = True + len_str2 = 1 + return potential_macro_user_exist, (len_str1 + len_str2 + 4) + + def found_potential_macro_user(self, macro, contents_list): + for line in contents_list: + if line.find(macro) >= 0: + return True + if line.find("CAT") >= 0: + words = ' '.join(re.split("(\W)", line)).split() + iter_w = 0 + while iter_w < len(words): + if words[iter_w] != "CAT": + iter_w += 1 + continue + user_exist, concat_len = self.found_concat_user(words, iter_w, macro) + if user_exist: + return True + iter_w += concat_len + return False + + def reduce_macros(self, contents): + new_contents = "" + contents_list = contents.split("\n") + idx = 0 + while idx < len(contents_list): + line = contents_list[idx] + is_macro = re.search('#\s*define', line) + macro = "" + + if is_macro: + words = ' '.join(re.split("(\W)", line)).split() + macro = words[words.index("define") + 1] + + if len(macro) == 0 or self.found_potential_macro_user(macro, contents_list): + new_contents += (line.rstrip() + "\n") + idx += 1 + else: + if line.rstrip()[-1] == '\\': + while contents_list[idx].rstrip()[-1] == '\\': + idx += 1 + idx += 1 + return new_contents + def append_file_content(self, filename, origin_file): res = "" content = [] with open(filename) as f: content += f.readlines() + for line in content: if line.startswith('#include'): include_file_name = line.strip().split('"')[1].strip() + if ntpath.basename(include_file_name) in self.batch_headers: + continue full_path_include = os.path.abspath(os.path.join(os.path.dirname(filename), include_file_name)) if full_path_include not in self.include_files[origin_file]: self.include_files[origin_file][full_path_include] = True @@ -77,12 +180,36 @@ class OpenCL2CHeaders(object): res += "\n" continue res += '{}\n'.format(line.rstrip()) + if filename == origin_file: + return self.reduce_macros(res) + else: + return res + + def batch_headers_to_str(self): + max_lines = 200 + max_characters = 16350 + characters = 1 # Newline character above + res = "" + for h in self.batch_headers: + res += '(std::string) R"(\n' + header_file = os.path.abspath(os.path.join(os.path.dirname(self.kernels_folder + "/include/batch_headers"), "batch_headers/" + h)) + content = [] + with open(header_file) as f: + content += f.readlines() + for i, line in enumerate(content): + if line.startswith('#include'): + continue + if (i + 1) % max_lines == 0 or characters + len(line) + 1 > max_characters: + res += ')",' + ' (std::string) R"(' + characters = 0 + res += '{}\n'.format(line.rstrip()) + characters += len(line) + 1 + res += ')",\n\n' return res def cl_file_to_str(self, filename): name = ntpath.basename(filename) self.include_files[filename] = {} - #kernel_name = name[:name.find('.')] kernel_name = name[:name.find('.cl')] res = '{{"{}",\n(std::string) R"__krnl(\n'.format(kernel_name) content = self.append_file_content(filename, filename) @@ -128,10 +255,11 @@ def main(): ap = argparse.ArgumentParser() ap.add_argument('-kernels', required=True, metavar='PATH', help='The absolute path to OpenCL kernels folder') ap.add_argument('-out_path', required=True, metavar='PATH', help='The absolute path to dump file') - ap.add_argument('-out_file_name', required=True, metavar='PATH', help='dump file name') + ap.add_argument('-out_file_name_prim_db', required=True, metavar='PATH', help='dump file name') + ap.add_argument('-out_file_name_batch_headers', required=True, metavar='PATH', help='dump file name') args = ap.parse_args() - converter = OpenCL2CHeaders(args.kernels, args.out_path, args.out_file_name) + converter = OpenCL2CHeaders(args.kernels, args.out_path, args.out_file_name_prim_db, args.out_file_name_batch_headers) converter.convert() if __name__ == '__main__': diff --git a/inference-engine/thirdparty/clDNN/runtime/kernels_cache.cpp b/inference-engine/thirdparty/clDNN/runtime/kernels_cache.cpp index 74cf5cf2816..9a8189098e0 100644 --- a/inference-engine/thirdparty/clDNN/runtime/kernels_cache.cpp +++ b/inference-engine/thirdparty/clDNN/runtime/kernels_cache.cpp @@ -193,20 +193,16 @@ void kernels_cache::get_program_source(const kernels_code& kernels_source_code, auto& current_bucket = program_buckets[key]; if (current_bucket.empty()) { // new bucket - const auto& bucket_id = program_buckets.size() - 1; - current_bucket.push_back(batch_program()); - current_bucket.back().bucket_id = static_cast(bucket_id); - current_bucket.back().batch_id = 0; - current_bucket.back().options = options; + const auto& batch_id = 0; + const auto& bucket_id = static_cast(program_buckets.size() - 1); + current_bucket.push_back(batch_program(bucket_id, batch_id, options, batch_header_str)); } // Create new kernels batch when the limit is reached if (current_bucket.back().kernels_counter >= get_max_kernels_per_batch()) { - const auto& batch_id = current_bucket.size(); - current_bucket.push_back(batch_program()); - current_bucket.back().bucket_id = static_cast(program_buckets.size()); - current_bucket.back().batch_id = static_cast(batch_id); - current_bucket.back().options = options; + const auto& bucket_id = static_cast(program_buckets.size()); + const auto& batch_id = static_cast(current_bucket.size()); + current_bucket.push_back(batch_program(bucket_id, batch_id, options, batch_header_str)); } auto& current_batch = current_bucket.back(); diff --git a/inference-engine/thirdparty/clDNN/runtime/kernels_cache.hpp b/inference-engine/thirdparty/clDNN/runtime/kernels_cache.hpp index df0d85f5b9d..35f8e62a8ab 100644 --- a/inference-engine/thirdparty/clDNN/runtime/kernels_cache.hpp +++ b/inference-engine/thirdparty/clDNN/runtime/kernels_cache.hpp @@ -93,14 +93,25 @@ class kernels_cache { public: using source_code = std::vector; struct batch_program { - int32_t bucket_id = 0; - int32_t batch_id = 0; - source_code source; + int32_t bucket_id; + int32_t batch_id; size_t hash_value; - uint32_t kernels_counter = 0; + uint32_t kernels_counter; + source_code source; std::string options; - bool dump_custom_program = false; + bool dump_custom_program; std::map entry_point_to_id; + + explicit batch_program(int32_t _bucket_id, int32_t _batch_id, std::string _options, const std::vector& batch_header_str) + : bucket_id(_bucket_id), + batch_id(_batch_id), + hash_value(0), + kernels_counter(0), + source(std::move(batch_header_str)), + options(_options), + dump_custom_program(false), + entry_point_to_id({}) { + } }; struct kernel_code { @@ -141,6 +152,7 @@ private: #elif(CLDNN_THREADING == CLDNN_THREADING_THREADPOOL) std::unique_ptr pool; #endif + std::vector batch_header_str; void get_program_source(const kernels_code& kernels_source_code, std::vector*) const; void build_batch(const engine& build_engine, const batch_program& batch); @@ -154,6 +166,9 @@ public: kernel_id set_kernel_source(const std::shared_ptr& kernel_string, bool dump_custom_program); kernel::ptr get_kernel(kernel_id id) const; + void set_batch_header_str(const std::vector &batch_headers) { + batch_header_str = std::move(batch_headers); + } // forces compilation of all pending kernels/programs void build_all(); void reset(); diff --git a/inference-engine/thirdparty/clDNN/src/program.cpp b/inference-engine/thirdparty/clDNN/src/program.cpp index 8a807710a4d..9ce691798e6 100644 --- a/inference-engine/thirdparty/clDNN/src/program.cpp +++ b/inference-engine/thirdparty/clDNN/src/program.cpp @@ -101,6 +101,7 @@ program::program(engine& engine_ref, set_options(); pm = std::unique_ptr(new pass_manager(*this)); prepare_nodes(topology); + _kernels_cache->set_batch_header_str(kernel_selector::KernelBase::get_db().get_batch_header_str()); if (no_optimizations) { init_graph(); } else { @@ -119,6 +120,7 @@ program::program(engine& engine_ref, tuning_cache(nullptr) { init_primitives(); set_options(); + _kernels_cache->set_batch_header_str(kernel_selector::KernelBase::get_db().get_batch_header_str()); pm = std::unique_ptr(new pass_manager(*this)); prepare_nodes(nodes); build_program(is_internal);