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

[GPU] Reduce unused macros to reduce loading time (#7435)

Browse Source 
* 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
This commit is contained in:
Taylor Yeonbok Lee 2021-10-01 15:18:15 +09:00 committed by GitHub
parent e0cea200e3
commit 61c97edd40
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
212 changed files with 740 additions and 593 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
inference-engine/thirdparty/clDNN/kernel_selector/CMakeLists.txt vendored
View File

@ -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}) ..."
)

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/activation_opt.cl vendored
View File

@ -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,

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/activation_ref.cl vendored
View File

@ -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)

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_axis.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_gpu_ref.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/arg_max_min_opt.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/average_unpooling_gpu_ref.cl vendored
View File

@ -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)
{

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/batch_to_space_ref.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_1x1.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_1x1_b_fs_yx_fsv16.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_generic.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/binary_convolution_gpu_ref.cl vendored
View File

@ -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,

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/border_gpu_ref.cl vendored
View File

@ -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)(

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/broadcast_gpu_ref.cl vendored
View File

@ -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)(

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_blocked.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_depth_bfyx_no_pitch.cl vendored
View File

@ -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".

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_fs_b_yx_fsv32.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_ref.cl vendored
View File

@ -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)

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/concatenation_gpu_simple_ref.cl vendored
View File

@ -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)

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv16_imad_1x1.cl vendored
View File

@ -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)

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_1x1.cl vendored
View File

@ -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:

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_dw.cl vendored
View File

@ -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:

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv4_int8.cl vendored
View File

@ -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

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_yx_fsv_16_32_imad_dw.cl vendored
View File

@ -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:

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_b_fs_zyx_fsv16_imad.cl vendored
View File

@ -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)

5
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1_hgemm_buf_16x1.cl vendored
View File

@ -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_) \

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_1x1_opt.cl vendored
View File

@ -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)))

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_depthwise_weights_lwg.cl vendored
View File

@ -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)))

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_direct_10_12_16.cl vendored
View File

@ -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);

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16.cl vendored
View File

@ -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)

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16_1x1.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_f16_depthwise.cl vendored
View File

@ -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

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_gemm_like_fp16.cl vendored
View File

@ -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);

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_gemm_like_fp32.cl vendored
View File

@ -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);

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_iyxo.cl vendored
View File

@ -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)))

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_os_iyx_osv16.cl vendored
View File

@ -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"
// ---------------------------------------------------------------------------------------------------------------------

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_bfyx_f16.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_bs_fs_yx_bsv16_fsv16.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_bfyx_to_fs_byx_fsv32.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32_1x1.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_fs_byx_fsv32_depthwise.cl vendored
View File

@ -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

5
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_imad_bs_fs_yx_bsv16_fsv16_1x1.cl vendored
View File

@ -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

5
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_imad_bs_fs_yx_bsv16_fsv16_3x3.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_b_fs_yx_fsv32.cl vendored
View File

@ -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)

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_b_fs_yx_fsv32_dw.cl vendored
View File

@ -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,

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_bfyx_to_b_fs_yx_fsv32.cl vendored
View File

@ -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))

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_mmad_bfyx_to_b_fs_yx_fsv4.cl vendored
View File

@ -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))

2
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_winograd_2x3_s1_fused.cl vendored
View File

@ -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) \

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_winograd_6x3_s1_fused.cl vendored
View File

@ -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) \

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_ref.cl vendored
View File

@ -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,

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b16_fp16.cl vendored
View File

@ -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)))

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b16_fp32.cl vendored
View File

@ -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)(

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b1_block_multiple_x_fp32.cl vendored
View File

@ -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)))

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/convolution_gpu_yxfb_yxio_b8_fp32.cl vendored
View File

@ -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)))

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/ctc_greedy_decoder_ref.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/cum_sum_partial_sum.cl vendored
View File

@ -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)

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/cum_sum_ref.cl vendored
View File

@ -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)

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_b_fs_zyx_fsv16_dw.cl vendored
View File

@ -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"

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_bfyx_opt.cl vendored
View File

@ -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

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_along_f_tile_bfx.cl vendored
View File

@ -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"

2
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_common.cl vendored
View File

@ -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))

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_imad_ref.cl vendored
View File

@ -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"

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deconvolution_gpu_ref.cl vendored
View File

@ -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,

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_conv.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_interp.cl vendored
View File

@ -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)(

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/deformable_convolution_gpu_bfyx_ref.cl vendored
View File

@ -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,

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/depth_to_space_block2_opt.cl vendored
View File

@ -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)
{

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/depth_to_space_ref.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/detection_output_gpu_ref.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_b_fs_yx_fsv16.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_b_fs_yx_fsv4.cl vendored
View File

@ -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)

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_fs_b_yx_fsv32.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_mixed_byxf_and_fs_b_yx_fsv32.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/eltwise_simple_vload8.cl vendored
View File

@ -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)

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/embedding_bag_ref.cl vendored
View File

@ -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)(

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/extract_image_patches_ref.cl vendored
View File

@ -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)

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_MMAD.cl vendored
View File

@ -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)

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_gemm.cl vendored
View File

@ -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)

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_input_spatial.cl vendored
View File

@ -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).

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_io_ref.cl vendored
View File

@ -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).

5
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bf_tiled.cl vendored
View File

@ -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};

6
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bfyx_ref.cl vendored
View File

@ -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,

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv16_af8_vload.cl vendored
View File

@ -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.

5
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv16_b1.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_bs_f_bsv8_af8_vload.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_b8_f8.cl vendored
View File

@ -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)))

5
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_b8_f8_vload.cl vendored
View File

@ -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

5
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_block_fp16.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_io_ref.cl vendored
View File

@ -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).

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_oi_b8_fp32_ref.cl vendored
View File

@ -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)))

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fb_oi_ref.cl vendored
View File

@ -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).

2
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_fs_byx_fsv32.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_imad.cl vendored
View File

@ -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

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fully_connected_gpu_yxfb_ref.cl vendored
View File

@ -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:

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_1x1_opt_fp32.cl vendored
View File

@ -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)))

5
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_iyxo.cl vendored
View File

@ -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)))

4
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_bfyx_os_iyx_osv16.cl vendored
View File

@ -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"
// ---------------------------------------------------------------------------------------------------------------------

2
inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/fused_conv_eltwise_gpu_imad.cl vendored
View File

@ -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

Some files were not shown because too many files have changed in this diff Show More