diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/pooling/pooling_kernel_gpu_b_fs_yx_fsv16.cpp b/inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/pooling/pooling_kernel_gpu_b_fs_yx_fsv16.cpp index aeb43373bd7..157430a6670 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/pooling/pooling_kernel_gpu_b_fs_yx_fsv16.cpp +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/pooling/pooling_kernel_gpu_b_fs_yx_fsv16.cpp @@ -53,11 +53,20 @@ size_t PoolingKernel_b_fs_yx_fsv16::GetBlockSize(const pooling_params& params) c return 1; } +size_t PoolingKernel_b_fs_yx_fsv16::GetSimdSize(const pooling_params& params) const { + auto& out = params.output; + // Use smaller simd size in case of global pooling and small channels count to have more threads + if (out.X().v == 1 && out.Y().v == 1 && out.Feature().v < 64) + return 8; + else + return 16; +} + PoolingKernelBase::DispatchData PoolingKernel_b_fs_yx_fsv16::SetDefault(const pooling_params& params) const { DispatchData kd = PoolingKernelBase::SetDefault(params); const auto& out = params.output; - const size_t alignment = 16; + const size_t alignment = GetSimdSize(params); size_t x_block_size = GetBlockSize(params); auto x = out.X().v; auto y = out.Y().v; @@ -78,7 +87,7 @@ PoolingKernelBase::DispatchData PoolingKernel_b_fs_yx_fsv16::SetDefault(const po } JitConstants PoolingKernel_b_fs_yx_fsv16::GetJitConstants(const pooling_params& params, DispatchData runInfo) const { - const size_t alignment = 16; + const size_t alignment = GetSimdSize(params); size_t x_block_size = GetBlockSize(params); auto input = params.inputs[0]; auto output = params.output; @@ -86,13 +95,26 @@ JitConstants PoolingKernel_b_fs_yx_fsv16::GetJitConstants(const pooling_params& size_t input_line_size = params.poolStride.x * (x_block_size - 1) + params.poolSize.x; + auto acc_type = GetAccumulatorType(params); + jit.Merge(MakeTypeJitConstants(acc_type, "ACCUMULATOR")); + + auto can_preload_full_line = [&]() -> bool { + const float max_reg_bytes = 128 * 32 * 0.95f; + const size_t line_bytes = input_line_size * BytesPerElement(input.GetDType()); + const size_t acc_bytes = x_block_size * BytesPerElement(acc_type); + + const float req_bytes = static_cast((line_bytes + acc_bytes) * alignment); + + return req_bytes < max_reg_bytes; + }; + + jit.AddConstant(MakeJitConstant("CAN_PRELOAD_FULL_LINE", can_preload_full_line())); jit.AddConstant(MakeJitConstant("PADDED_INPUT", params.inputs[0].X().pad.Total() != 0)); jit.AddConstant(MakeJitConstant("X_BLOCK_SIZE", x_block_size)); jit.AddConstant(MakeJitConstant("INPUT_LINE_SIZE", input_line_size)); jit.AddConstant(MakeJitConstant("SUB_GROUP_SIZE", alignment)); jit.AddConstant(MakeJitConstant("X_BLOCKS", CeilDiv(output.X().v, x_block_size))); jit.Merge(MakeTypeJitConstants(GetActivationType(params), "ACTIVATION")); - jit.Merge(MakeTypeJitConstants(GetAccumulatorType(params), "ACCUMULATOR")); if (params.output.Feature().v % 16 != 0) { jit.AddConstant(MakeJitConstant("OUTPUT_LEFTOVERS", 1)); @@ -101,7 +123,7 @@ JitConstants PoolingKernel_b_fs_yx_fsv16::GetJitConstants(const pooling_params& if (!params.fused_ops.empty()) { auto input_dt = GetActivationType(params); FusedOpsConfiguration conf_vec = {"_VEC", - {"b", "(f_block*16)", "y", "x"}, + {"b", "(f_block*FEATURE_SLICE_SIZE + f_val*SUB_GROUP_SIZE)", "y", "x"}, "pool_result", input_dt, x_block_size, @@ -110,7 +132,7 @@ JitConstants PoolingKernel_b_fs_yx_fsv16::GetJitConstants(const pooling_params& IndexType::TENSOR_COORD, Tensor::DataChannelName::X}; FusedOpsConfiguration conf_scalar = {"_SCALAR", - {"b", "(f_block*16)", "y", "(x+i)"}, + {"b", "(f_block*FEATURE_SLICE_SIZE + f_val*SUB_GROUP_SIZE)", "y", "(x+i)"}, "pool_result[i]", input_dt, 1, diff --git a/inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/pooling/pooling_kernel_gpu_b_fs_yx_fsv16.h b/inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/pooling/pooling_kernel_gpu_b_fs_yx_fsv16.h index 4877d4f9d10..6b35c94cd1b 100644 --- a/inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/pooling/pooling_kernel_gpu_b_fs_yx_fsv16.h +++ b/inference-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/pooling/pooling_kernel_gpu_b_fs_yx_fsv16.h @@ -37,5 +37,6 @@ protected: } size_t GetBlockSize(const pooling_params& params) const; + size_t GetSimdSize(const pooling_params& params) const; }; } // namespace kernel_selector 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 c20dbc1775a..92ad4682d3d 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 @@ -39,7 +39,7 @@ #define INIT_VAL ACCUMULATOR_VAL_ZERO #endif -__attribute__((intel_reqd_sub_group_size(16))) +__attribute__((intel_reqd_sub_group_size(SUB_GROUP_SIZE))) KERNEL(pooling_gpu_blocked)( const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output @@ -49,7 +49,13 @@ KERNEL(pooling_gpu_blocked)( ) { const int lid = get_sub_group_local_id(); +#if SUB_GROUP_SIZE == 16 const int f_block = get_group_id(1); + const int f_val = 0; +#else + const int f_block = (uint)get_group_id(1) / (FEATURE_SLICE_SIZE / SUB_GROUP_SIZE); + const int f_val = (uint)get_group_id(1) % (FEATURE_SLICE_SIZE / SUB_GROUP_SIZE); +#endif const int b = get_global_id(2); const int xy = get_global_id(0); @@ -71,7 +77,8 @@ KERNEL(pooling_gpu_blocked)( const uint input_offset = b * input_b_pitch + (f_block + input_fs_pad_before) * input_fs_pitch + (INPUT0_PAD_BEFORE_SIZE_Y + input_y) * input_y_pitch + - (INPUT0_PAD_BEFORE_SIZE_X + input_x) * input_x_pitch; + (INPUT0_PAD_BEFORE_SIZE_X + input_x) * input_x_pitch + + f_val * SUB_GROUP_SIZE; // Output offset calculations: const uint output_x_pitch = FEATURE_SLICE_SIZE; @@ -85,7 +92,8 @@ KERNEL(pooling_gpu_blocked)( const uint output_offset = b * output_b_pitch + (f_block + output_fs_pad_before) * output_fs_pitch + (y + OUTPUT_PAD_BEFORE_SIZE_Y) * output_y_pitch + - (x + OUTPUT_PAD_BEFORE_SIZE_X) * output_x_pitch; + (x + OUTPUT_PAD_BEFORE_SIZE_X) * output_x_pitch + + f_val * SUB_GROUP_SIZE; ACCUMULATOR_VAR_TYPE dst = (ACCUMULATOR_VAR_TYPE)INIT_VAL; @@ -122,12 +130,13 @@ KERNEL(pooling_gpu_blocked)( if (input_y + kh < 0 || input_y + kh >= INPUT0_SIZE_Y) continue; +#if CAN_PRELOAD_FULL_LINE INPUT0_TYPE line_cache[INPUT_LINE_SIZE]; for (int i = 0; i < INPUT_LINE_SIZE; i++) { if ((input_x + i) >= 0 && (input_x + i) < INPUT0_SIZE_X) line_cache[i] = DT_INPUT_BLOCK_READ(input, input_offset + kh*input_y_pitch + i*input_x_pitch); else - #if defined MAX_POOLING + #if defined MAX_POOLING line_cache[i] = INPUT0_VAL_MIN; #elif defined AVG_POOLING line_cache[i] = INPUT0_VAL_ZERO; @@ -152,35 +161,56 @@ KERNEL(pooling_gpu_blocked)( dst += src; #endif } - } +#else // CAN_PRELOAD_FULL_LINE + // TODO: try partial preload + for (int kw = 0; kw < POOL_SIZE_X; kw++) + { + INPUT_VAR_TYPE src; +#if X_BLOCK_SIZE > 1 + for (int i = 0; i < X_BLOCK_SIZE; i++) { + if ((input_x + kw + STRIDE_SIZE_X*i) >= 0 && (input_x + kw + STRIDE_SIZE_X*i) < INPUT0_SIZE_X) + src[i] = DT_INPUT_BLOCK_READ(input, input_offset + kh*input_y_pitch + (kw + STRIDE_SIZE_X*i)*input_x_pitch); + else + #if defined MAX_POOLING + src[i] = INPUT0_VAL_MIN; + #elif defined AVG_POOLING + src[i] = INPUT0_VAL_ZERO; + #endif + } +#else + src = DT_INPUT_BLOCK_READ(input, input_offset + kh*input_y_pitch + kw*input_x_pitch); +#endif +#if defined MAX_POOLING + dst = ACCUMULATOR_MAX_FUNC(dst, src); +#elif defined AVG_POOLING + dst += TO_ACCUMULATOR_VAR_TYPE(src); +#endif + } +#endif // CAN_PRELOAD_FULL_LINE + } ACTIVATION_VAR_TYPE pool_result; #if defined MAX_POOLING - pool_result = TO_ACTIVATION_VAR_TYPE(dst); - #if !HAS_FUSED_OP - pool_result = ACTIVATION(pool_result, ACTIVATION_PARAMS); - #endif + pool_result = TO_ACTIVATION_VAR_TYPE(dst); #elif defined AVG_POOLING && (defined(DYNAMIC_KERNEL_DIVIDER) || defined(DYNAMIC_WITH_PADDING_KERNEL_DIVIDER)) - pool_result = TO_ACTIVATION_VAR_TYPE(dst*scale); - #if !HAS_FUSED_OP - pool_result = ACTIVATION(pool_result, ACTIVATION_PARAMS); - #endif + pool_result = TO_ACTIVATION_VAR_TYPE(dst*scale); #elif defined AVG_POOLING - pool_result = TO_ACTIVATION_VAR_TYPE(dst/(POOL_SIZE_X*POOL_SIZE_Y)); - #if !HAS_FUSED_OP - pool_result = ACTIVATION(pool_result, ACTIVATION_PARAMS); - #endif + pool_result = TO_ACTIVATION_VAR_TYPE(dst/(POOL_SIZE_X*POOL_SIZE_Y)); +#endif + +#if !HAS_FUSED_OPS + pool_result = ACTIVATION(pool_result, ACTIVATION_PARAMS); #endif OUTPUT_VAR_TYPE final_result; #if OUTPUT_LEFTOVERS - if ((f_block+1)*FEATURE_SLICE_SIZE >= OUTPUT_FEATURE_NUM) { + if (f_block*FEATURE_SLICE_SIZE + (f_val + 1)*SUB_GROUP_SIZE >= OUTPUT_FEATURE_NUM) { for (int i = 0; i < X_BLOCK_SIZE; i++) { - if ((f_block*FEATURE_SLICE_SIZE + lid < OUTPUT_FEATURE_NUM) && (x + i) < OUTPUT_SIZE_X) { + if ((f_block*FEATURE_SLICE_SIZE + f_val*SUB_GROUP_SIZE + lid < OUTPUT_FEATURE_NUM) && (x + i) < OUTPUT_SIZE_X) { #if X_BLOCK_SIZE > 1 - #if HAS_FUSED_OP + #if HAS_FUSED_OPS FUSED_OPS_SCALAR; final_result[i] = FUSED_OPS_RESULT_SCALAR; #else @@ -195,7 +225,6 @@ KERNEL(pooling_gpu_blocked)( final_result = TO_OUTPUT_VAR_TYPE(pool_result); #endif output[output_offset + i * output_x_pitch + lid] = final_result; - #endif } } @@ -211,6 +240,7 @@ KERNEL(pooling_gpu_blocked)( final_result = TO_OUTPUT_VAR_TYPE(pool_result); #endif +#if SUB_GROUP_SIZE == FEATURE_SLICE_SIZE #if X_BLOCK_SIZE == 8 DT_OUTPUT_BLOCK_WRITE8(output, output_offset, final_result); #elif X_BLOCK_SIZE == 4 @@ -220,10 +250,20 @@ KERNEL(pooling_gpu_blocked)( #elif X_BLOCK_SIZE == 1 DT_OUTPUT_BLOCK_WRITE(output, output_offset, final_result); #endif +#else + #if X_BLOCK_SIZE > 1 + __attribute__((opencl_unroll_hint(X_BLOCK_SIZE))) + for (int i = 0; i < X_BLOCK_SIZE; i++) { + DT_OUTPUT_BLOCK_WRITE(output, output_offset + i * output_x_pitch, final_result[i]); + } + #else + DT_OUTPUT_BLOCK_WRITE(output, output_offset, final_result); + #endif +#endif } else { - const int x_tail = OUTPUT_SIZE_X - x; + const int x_tail = OUTPUT_SIZE_X % X_BLOCK_SIZE; for (int i = 0; i < x_tail; i++){ #if X_BLOCK_SIZE > 1 #if HAS_FUSED_OPS