From e4b289031ff445d1de61c84e0552adc0789900fe Mon Sep 17 00:00:00 2001 From: Nikolay Shchegolev Date: Tue, 7 Jul 2026 07:14:53 +0000 Subject: [PATCH] [GPU] Fix for Interpolate func tests on PTL --- .../cl_kernels/resample_bfyx_cubic_opt.cl | 64 ++++++-- .../cl_kernels/resample_onnx.cl | 14 +- .../cl_kernels/resample_opt.cl | 6 +- .../cl_kernels/resample_ref.cl | 111 ++++++++++++-- .../kernels/resample/resample_kernel_base.cpp | 20 ++- .../resample_kernel_bfyx_cubic_opt.cpp | 12 +- .../kernels/resample/resample_kernel_opt.cpp | 139 +++++++++++++++++ .../resample/resample_kernel_pil_ref.cpp | 4 +- .../kernels/resample/resample_kernel_ref.cpp | 143 ++++++++++++++++++ .../intel_gpu/src/plugin/ops/interpolate.cpp | 7 +- .../single_layer_tests/interpolate.cpp | 4 +- .../dynamic/interpolate.cpp | 4 + .../shared/src/single_op/interpolate.cpp | 14 +- 13 files changed, 495 insertions(+), 47 deletions(-) diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_bfyx_cubic_opt.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_bfyx_cubic_opt.cl index ffdfdaee49da..396c0aa5cb9f 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_bfyx_cubic_opt.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_bfyx_cubic_opt.cl @@ -4,26 +4,41 @@ #include "include/fetch_utils.cl" +#pragma OPENCL FP_CONTRACT OFF + #ifdef RTE_OUTPUT #define TO_OUTPUT_TYPE(x) CAT(CAT(convert_, OUTPUT_TYPE), _rte)(x) #else #define TO_OUTPUT_TYPE(x) CAT(convert_, OUTPUT_TYPE)(x) #endif +inline float FUNC(ref_divide)(float numerator, float denominator) +{ + volatile float numerator_value = numerator; + volatile float denominator_value = denominator; + volatile float quotient = numerator_value / denominator_value; + volatile float residual = numerator_value - quotient * denominator_value; + return quotient + residual / denominator_value; +} + inline float FUNC(get_original_coordinate)(float num, float scale, int length_resized, int length_original) { if (scale == 1.0f) return num; #if defined(COORD_TRANS_MODE_HALF_PIXEL) - return (num + 0.5f) * scale - 0.5f; + return FUNC_CALL(ref_divide)(num + 0.5f, scale) - 0.5f; #elif defined(COORD_TRANS_MODE_PYTORCH_HALF_PIXEL) - return (length_resized > 1) ? (num + 0.5f) * scale - 0.5f : 0.f; + return (length_resized > 1) ? FUNC_CALL(ref_divide)(num + 0.5f, scale) - 0.5f : 0.f; #elif defined(COORD_TRANS_MODE_ASYMMETRIC) - return num * scale; + return FUNC_CALL(ref_divide)(num, scale); #elif defined(COORD_TRANS_MODE_TF_HALF_PIXEL_FOR_NN) - return (num + 0.5f) * scale; + return FUNC_CALL(ref_divide)(num + 0.5f, scale); #elif defined(COORD_TRANS_MODE_ALIGN_CORNERS) - return (length_resized != 1) ? num * (length_original - 1) / (length_resized - 1) : 0.f; + if (length_resized == 1) + return 0.f; + if (num == 0.f || num == (float)(length_resized - 1)) + return num == 0.f ? 0.f : (float)(length_original - 1); + return FUNC_CALL(ref_divide)((float)((int)num * (length_original - 1)), (float)(length_resized - 1)); #else #error [clDNN resample_bfyx_cubic_opt.cl]: coordinate transformation mode - not supported #endif @@ -32,10 +47,32 @@ inline float FUNC(get_original_coordinate)(float num, float scale, int length_re inline void FUNC(get_cubic_coeff)(float* cubic_coef, float coord, float coef) { float abs_num = fabs(coord); - cubic_coef[0] = coef * (abs_num - 1.0f) * (abs_num - 1.0f) * abs_num; - cubic_coef[1] = ((coef + 2.0f) * abs_num - (coef + 3.0f)) * abs_num * abs_num + 1.0f; - cubic_coef[2] = (((-coef - 2.0f) * abs_num + (2.0f * coef + 3.0f)) * abs_num - coef) * abs_num; - cubic_coef[3] = -coef * abs_num * abs_num * (abs_num - 1.0f); + float x0 = abs_num + 1.0f; + float x1 = abs_num; + float x2 = 1.0f - abs_num; + float x3 = 2.0f - abs_num; + float t0 = coef * x0; + t0 = t0 - 5.0f * coef; + t0 = t0 * x0; + t0 = t0 + 8.0f * coef; + t0 = t0 * x0; + cubic_coef[0] = t0 - 4.0f * coef; + float t1 = (coef + 2.0f) * x1; + t1 = t1 - (coef + 3.0f); + t1 = t1 * x1; + t1 = t1 * x1; + cubic_coef[1] = t1 + 1.0f; + float t2 = (coef + 2.0f) * x2; + t2 = t2 - (coef + 3.0f); + t2 = t2 * x2; + t2 = t2 * x2; + cubic_coef[2] = t2 + 1.0f; + float t3 = coef * x3; + t3 = t3 - 5.0f * coef; + t3 = t3 * x3; + t3 = t3 + 8.0f * coef; + t3 = t3 * x3; + cubic_coef[3] = t3 - 4.0f * coef; } KERNEL (resample_bfyx_cubic_opt)( @@ -57,7 +94,7 @@ KERNEL (resample_bfyx_cubic_opt)( // Compute Y coordinate mapping once (shared for all X in the block) const float orig_y = FUNC_CALL(get_original_coordinate)((float)out_y, SCALES[3], OUTPUT_SIZE_Y, - INPUT0_SIZE_Y + PADS_BEGIN[3] + PADS_END[3]) - PADS_BEGIN[3]; + INPUT0_SIZE_Y + PADS_BEGIN[3] + PADS_END[3]) - PADS_BEGIN[3]; const int iy = (int)floor(orig_y); const float y_frac = orig_y - (float)iy; float cy[4]; @@ -76,7 +113,7 @@ KERNEL (resample_bfyx_cubic_opt)( // Compute X coordinate mapping const float orig_x = FUNC_CALL(get_original_coordinate)((float)ox, SCALES[4], OUTPUT_SIZE_X, - INPUT0_SIZE_X + PADS_BEGIN[4] + PADS_END[4]) - PADS_BEGIN[4]; + INPUT0_SIZE_X + PADS_BEGIN[4] + PADS_END[4]) - PADS_BEGIN[4]; const int ix = (int)floor(orig_x); const float x_frac = orig_x - (float)ix; float cx[4]; @@ -98,9 +135,8 @@ KERNEL (resample_bfyx_cubic_opt)( x_idx[dx] >= 0 && x_idx[dx] < INPUT0_SIZE_X) #endif { - interp_val = fma((ACCUMULATOR_TYPE)(cy[dy] * cx[dx]), - (ACCUMULATOR_TYPE)input[INPUT0_GET_INDEX(out_b, out_f, y_idx[dy], x_idx[dx])], - interp_val); + const float term = cy[dy] * cx[dx] * (float)input[INPUT0_GET_INDEX(out_b, out_f, y_idx[dy], x_idx[dx])]; + interp_val = interp_val + (ACCUMULATOR_TYPE)term; } } } diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_onnx.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_onnx.cl index 91869aba8f0f..0db0f11346d0 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_onnx.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_onnx.cl @@ -26,15 +26,19 @@ inline float FUNC(get_original_coordinate)(float num, float scale, int length_re if (scale == 1.0f) return num; #if defined(COORD_TRANS_MODE_HALF_PIXEL) - return (num + 0.5f) * scale - 0.5f; + return (num + 0.5f) / scale - 0.5f; #elif defined(COORD_TRANS_MODE_PYTORCH_HALF_PIXEL) - return (length_resized > 1) ? (num + 0.5f) * scale - 0.5f : 0.f; + return (length_resized > 1) ? (num + 0.5f) / scale - 0.5f : 0.f; #elif defined(COORD_TRANS_MODE_ASYMMETRIC) - return num * scale; + return num / scale; #elif defined(COORD_TRANS_MODE_TF_HALF_PIXEL_FOR_NN) - return (num + 0.5f) * scale; + return (num + 0.5f) / scale; #elif defined(COORD_TRANS_MODE_ALIGN_CORNERS) - return (length_resized != 1) ? num * (length_original - 1) / (length_resized - 1) : 0.f; + if (length_resized == 1) + return 0.f; + if (num == 0.f || num == (float)(length_resized - 1)) + return num == 0.f ? 0.f : (float)(length_original - 1); + return (float)((int)num * (length_original - 1)) / (float)(length_resized - 1); #else #error [clDNN resample_onnx.cl]: coordinate transformation mode - not supported #endif diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_opt.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_opt.cl index 533db6da6410..e5b4634e3b8c 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_opt.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_opt.cl @@ -41,7 +41,11 @@ inline float FUNC(get_original_coordinate)(float num, float scale, int length_re #elif defined(COORD_TRANS_MODE_TF_HALF_PIXEL_FOR_NN) return (num + 0.5f) * scale; #elif defined(COORD_TRANS_MODE_ALIGN_CORNERS) - return (length_resized != 1) ? num * (length_original - 1) / (length_resized - 1) : 0.f; + if (length_resized == 1) + return 0.f; + if (num == 0.f || num == (float)(length_resized - 1)) + return num == 0.f ? 0.f : (float)(length_original - 1); + return (float)((int)num * (length_original - 1)) / (float)(length_resized - 1); #else #error [clDNN resample_opt.cl]: coordinate transformation mode - not supported #endif diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_ref.cl index 6bc2e26bb436..e555b5d3fbbf 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_ref.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_ref.cl @@ -4,6 +4,8 @@ #include "include/fetch_utils.cl" +#pragma OPENCL FP_CONTRACT OFF + #ifdef RTE_OUTPUT #define TO_OUTPUT_TYPE(x) CAT(CAT(convert_, OUTPUT_TYPE), _rte)(x) #endif @@ -25,20 +27,49 @@ inline int FUNC(get_nearest_val)(float num, bool is_downsample) #endif } +inline float FUNC(ref_divide)(float numerator, float denominator) +{ + volatile float numerator_value = numerator; + volatile float denominator_value = denominator; + volatile float quotient = numerator_value / denominator_value; + volatile float residual = numerator_value - quotient * denominator_value; + return quotient + residual / denominator_value; +} + inline float FUNC(get_original_coordinate)(float num, float scale, int length_resized, int length_original) { if (scale == 1.0f) return num; #if defined(COORD_TRANS_MODE_HALF_PIXEL) +#if RESAMPLE_USE_LEGACY_SCALE == 1 return (num + 0.5f) * scale - 0.5f; +#else + return FUNC_CALL(ref_divide)(num + 0.5f, scale) - 0.5f; +#endif #elif defined(COORD_TRANS_MODE_PYTORCH_HALF_PIXEL) +#if RESAMPLE_USE_LEGACY_SCALE == 1 return (length_resized > 1) ? (num + 0.5f) * scale - 0.5f : 0.f; +#else + return (length_resized > 1) ? FUNC_CALL(ref_divide)(num + 0.5f, scale) - 0.5f : 0.f; +#endif #elif defined(COORD_TRANS_MODE_ASYMMETRIC) +#if RESAMPLE_USE_LEGACY_SCALE == 1 return num * scale; +#else + return FUNC_CALL(ref_divide)(num, scale); +#endif #elif defined(COORD_TRANS_MODE_TF_HALF_PIXEL_FOR_NN) +#if RESAMPLE_USE_LEGACY_SCALE == 1 return (num + 0.5f) * scale; +#else + return FUNC_CALL(ref_divide)(num + 0.5f, scale); +#endif #elif defined(COORD_TRANS_MODE_ALIGN_CORNERS) - return (length_resized != 1) ? num * (length_original - 1) / (length_resized - 1) : 0.f; + if (length_resized == 1) + return 0.f; + if (num == 0.f || num == (float)(length_resized - 1)) + return num == 0.f ? 0.f : (float)(length_original - 1); + return FUNC_CALL(ref_divide)((float)((int)num * (length_original - 1)), (float)(length_resized - 1)); #else #error [clDNN resample_ref.cl]: coordinate transformation mode - not supported #endif @@ -47,10 +78,32 @@ inline float FUNC(get_original_coordinate)(float num, float scale, int length_re inline void FUNC(get_cubic_coeff)(float* cubic_coef, float coord, float coef) { float abs_num = fabs(coord); - cubic_coef[0] = coef * (abs_num - 1.0) * (abs_num - 1.0) * abs_num; - cubic_coef[1] = ((coef + 2.0) * abs_num - (coef + 3.0)) * abs_num * abs_num + 1.0; - cubic_coef[2] = (((-coef - 2.0) * abs_num + (2.0 * coef + 3.0)) * abs_num - coef) * abs_num; - cubic_coef[3] = -coef * abs_num * abs_num * (abs_num - 1.0); + float x0 = abs_num + 1.0f; + float x1 = abs_num; + float x2 = 1.0f - abs_num; + float x3 = 2.0f - abs_num; + float t0 = coef * x0; + t0 = t0 - 5.0f * coef; + t0 = t0 * x0; + t0 = t0 + 8.0f * coef; + t0 = t0 * x0; + cubic_coef[0] = t0 - 4.0f * coef; + float t1 = (coef + 2.0f) * x1; + t1 = t1 - (coef + 3.0f); + t1 = t1 * x1; + t1 = t1 * x1; + cubic_coef[1] = t1 + 1.0f; + float t2 = (coef + 2.0f) * x2; + t2 = t2 - (coef + 3.0f); + t2 = t2 * x2; + t2 = t2 * x2; + cubic_coef[2] = t2 + 1.0f; + float t3 = coef * x3; + t3 = t3 - 5.0f * coef; + t3 = t3 * x3; + t3 = t3 + 8.0f * coef; + t3 = t3 * x3; + cubic_coef[3] = t3 - 4.0f * coef; } #define TRIANGLE_COEFF(x) (ACCUMULATOR_MAX_FUNC(ACCUMULATOR_VAL_ZERO, ACCUMULATOR_VAL_ONE - ACCUMULATOR_ABS_FUNC(x))) @@ -80,18 +133,28 @@ KERNEL (resample_gpu_ref)(__global INPUT0_TYPE* input, out_coords[1] = ((int)get_global_id(2) * PACK_SIZE) % OUTPUT_FEATURE_NUM; out_coords[0] = ((int)get_global_id(2) * PACK_SIZE) / OUTPUT_FEATURE_NUM; int in_coords[5]; + int safe_in_coords[5]; bool isOutOfBounds = false; +#if RESAMPLE_FAST_NEAREST == 1 + safe_in_coords[4] = (int)floor(out_coords[4] * SCALES[4]); + safe_in_coords[3] = (int)floor(out_coords[3] * SCALES[3]); + safe_in_coords[2] = (int)floor(out_coords[2] * SCALES[2]); + safe_in_coords[1] = out_coords[1]; + safe_in_coords[0] = out_coords[0]; +#else unroll_for (int i = 0; i < 5; ++i) { const float orig_coord = FUNC_CALL(get_original_coordinate)(out_coords[i], SCALES[i], out_size[i], in_size[i] + PADS_BEGIN[i] + PADS_END[i]); - const int nearest_pixel = FUNC_CALL(get_nearest_val)(orig_coord, SCALES[i] > 1) - PADS_BEGIN[i]; - in_coords[i] = max(-PADS_BEGIN[0], min(nearest_pixel, in_size[i] + PADS_END[i] - 1)); + const int nearest_pixel = FUNC_CALL(get_nearest_val)(orig_coord, SCALES[i] < 1) - PADS_BEGIN[i]; + in_coords[i] = max(-PADS_BEGIN[i], min(nearest_pixel, in_size[i] + PADS_END[i] - 1)); + safe_in_coords[i] = clamp(in_coords[i], 0, in_size[i] - 1); #if PADDING_USED == 1 if (in_coords[i] < 0 || in_coords[i] >= in_size[i]) isOutOfBounds = true; #endif } +#endif - uint input_idx = FUNC_CALL(get_input_index)(in_coords[0], in_coords[1], 0, in_coords[2], in_coords[3], in_coords[4]); + uint input_idx = FUNC_CALL(get_input_index)(safe_in_coords[0], safe_in_coords[1], 0, safe_in_coords[2], safe_in_coords[3], safe_in_coords[4]); uint output_idx = FUNC_CALL(get_output_index)(out_coords[0], out_coords[1], 0, out_coords[2], out_coords[3], out_coords[4]); in_pack_t interp_val_pack = ((const __global in_pack_t*)(input + input_idx))[0]; @@ -134,17 +197,27 @@ KERNEL (resample_gpu_ref)(__global INPUT0_TYPE* input, out_coords[1] = (int)get_global_id(2) % OUTPUT_FEATURE_NUM; out_coords[0] = (int)get_global_id(2) / OUTPUT_FEATURE_NUM; int in_coords[5]; + int safe_in_coords[5]; bool isOutOfBounds = false; +#if RESAMPLE_FAST_NEAREST == 1 + safe_in_coords[4] = (int)floor(out_coords[4] * SCALES[4]); + safe_in_coords[3] = (int)floor(out_coords[3] * SCALES[3]); + safe_in_coords[2] = (int)floor(out_coords[2] * SCALES[2]); + safe_in_coords[1] = out_coords[1]; + safe_in_coords[0] = out_coords[0]; +#else unroll_for (int i = 0; i < 5; ++i) { const float orig_coord = FUNC_CALL(get_original_coordinate)(out_coords[i], SCALES[i], out_size[i], in_size[i] + PADS_BEGIN[i] + PADS_END[i]); - int nearest_pixel = FUNC_CALL(get_nearest_val)(orig_coord, SCALES[i] > 1) - PADS_BEGIN[i]; + int nearest_pixel = FUNC_CALL(get_nearest_val)(orig_coord, SCALES[i] < 1) - PADS_BEGIN[i]; in_coords[i] = max(-PADS_BEGIN[i], min(nearest_pixel, in_size[i] + PADS_END[i] - 1)); + safe_in_coords[i] = clamp(in_coords[i], 0, in_size[i] - 1); #if PADDING_USED == 1 if (in_coords[i] < 0 || in_coords[i] >= in_size[i]) isOutOfBounds = true; #endif } - INPUT0_TYPE interp_val = input[FUNC_CALL(get_input_index)(in_coords[0], in_coords[1], 0, in_coords[2], in_coords[3], in_coords[4])]; +#endif + INPUT0_TYPE interp_val = input[FUNC_CALL(get_input_index)(safe_in_coords[0], safe_in_coords[1], 0, safe_in_coords[2], safe_in_coords[3], safe_in_coords[4])]; #if PADDING_USED == 1 if (isOutOfBounds) interp_val = INPUT0_VAL_ZERO; @@ -182,12 +255,22 @@ KERNEL (resample_gpu_ref)(__global INPUT0_TYPE* input, float cubic_coeff[5][4]; unroll_for (int i = 0; i < 5; ++i) { float orig_coord = FUNC_CALL(get_original_coordinate)(out_coords[i], SCALES[i], out_size[i], in_size[i] + PADS_BEGIN[i] + PADS_END[i]) - PADS_BEGIN[i]; + #if SHAPE_CALC_MODE_SIZES && PADDING_USED == 1 && defined(COORD_TRANS_MODE_TF_HALF_PIXEL_FOR_NN) + if ((PADS_BEGIN[i] == 0) != (PADS_END[i] == 0)) + orig_coord = ((float)out_coords[i] + 0.5f) / (float)out_size[i] * (float)(in_size[i] + PADS_BEGIN[i] + PADS_END[i]) - PADS_BEGIN[i]; + else if (PADS_BEGIN[i] != 0 && PADS_END[i] != 0) { + volatile float inv_scale = 1.0f / SCALES[i]; + orig_coord = ((float)out_coords[i] + 0.5f) * inv_scale - PADS_BEGIN[i]; + } + #elif SHAPE_CALC_MODE_SIZES && PADDING_USED == 1 && defined(COORD_TRANS_MODE_ASYMMETRIC) + orig_coord = (float)out_coords[i] / (float)out_size[i] * (float)(in_size[i] + PADS_BEGIN[i] + PADS_END[i]) - PADS_BEGIN[i]; + #endif in_coords[i] = floor(orig_coord); orig_coord = (orig_coord - in_coords[i]) * AXES_USED[i]; FUNC_CALL(get_cubic_coeff)(cubic_coeff[i], orig_coord, CUBE_COEFF); } - INPUT0_TYPE interp_val = INPUT0_VAL_ZERO; + ACCUMULATOR_TYPE interp_val = ACCUMULATOR_VAL_ZERO; int index[5]; unroll_for (index[0] = 0; index[0] <= 3; ++index[0]) { unroll_for (index[1] = 0; index[1] <= 3; ++index[1]) { @@ -208,7 +291,11 @@ KERNEL (resample_gpu_ref)(__global INPUT0_TYPE* input, #if PADDING_USED == 1 if (!isOutOfBounds) #endif - interp_val += coeff_prod * input[FUNC_CALL(get_input_index)(coords_sum[0], coords_sum[1], 0, coords_sum[2], coords_sum[3], coords_sum[4])]; + { + interp_val = fma((ACCUMULATOR_TYPE)coeff_prod, + (ACCUMULATOR_TYPE)input[FUNC_CALL(get_input_index)(coords_sum[0], coords_sum[1], 0, coords_sum[2], coords_sum[3], coords_sum[4])], + interp_val); + } } } } diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_base.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_base.cpp index 6be76c4d0ed3..2e8230735e7c 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_base.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_base.cpp @@ -146,17 +146,17 @@ JitConstants ResampleKernelBase::GetJitConstants(const resample_params& params) paddingUsed |= (pads_begin[i] != 0 || pads_end[i] != 0); } - scales[0] = static_cast(b_size_padded) / static_cast(out_b_size_padded); - scales[1] = static_cast(f_size_padded) / static_cast(out_f_size_padded); - scales[4] = static_cast(x_size_padded) / static_cast(out_x_size_padded); - scales[3] = static_cast(y_size_padded) / static_cast(out_y_size_padded); - scales[2] = static_cast(z_size_padded) / static_cast(out_z_size_padded); + scales[0] = static_cast(out_b_size_padded) / static_cast(b_size_padded); + scales[1] = static_cast(out_f_size_padded) / static_cast(f_size_padded); + scales[4] = static_cast(out_x_size_padded) / static_cast(x_size_padded); + scales[3] = static_cast(out_y_size_padded) / static_cast(y_size_padded); + scales[2] = static_cast(out_z_size_padded) / static_cast(z_size_padded); for (std::size_t i = 0; i < params.axes.size(); i++) { int idx = getAxisIndex(params.axes[i]); axesUsed[idx] = 1; if (params.shapeCalculationMode == kernel_selector::ShapeCalculationMode::SCALES) - scales[idx] = 1.f / params.scales[i]; + scales[idx] = params.scales[i]; } for (size_t i = 0; i < scales.size(); ++i) { if (scales[i] != 1.f) @@ -167,6 +167,7 @@ JitConstants ResampleKernelBase::GetJitConstants(const resample_params& params) MakeJitConstant(toString(params.resampleType), ""), MakeJitConstant(toString(params.nearestMode), ""), MakeJitConstant(toString(params.coordTransMode), ""), + MakeJitConstant("SHAPE_CALC_MODE_SIZES", params.shapeCalculationMode == ShapeCalculationMode::SIZES), MakeJitConstant("SCALES", scales), MakeJitConstant("PADS_BEGIN", pads_begin), MakeJitConstant("PADS_END", pads_end), @@ -234,6 +235,13 @@ KernelsData ResampleKernelBase::GetCommonKernelsData(const Params& params) const auto& kernel = kd.kernels[0]; FillCLKernelData(kernel, dispatchData, params.engineInfo, kernelName, jit, entry_point, EXE_MODE_DEFAULT, false, false, 1, GetFusedPrimitiveInputsCount(params)); + if (newParams.resampleType == ResampleType::CUBIC && kernel.code.kernelString) { + auto& options = kernel.code.kernelString->options; + const std::string mad_option = " -cl-mad-enable"; + const auto mad_pos = options.find(mad_option); + if (mad_pos != std::string::npos) + options.erase(mad_pos, mad_option.size()); + } return {kd}; } diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_bfyx_cubic_opt.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_bfyx_cubic_opt.cpp index 048d9958f818..6b7075711e53 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_bfyx_cubic_opt.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_bfyx_cubic_opt.cpp @@ -3,6 +3,7 @@ // #include "resample_kernel_bfyx_cubic_opt.h" +#include #include #include @@ -30,6 +31,7 @@ ParamsKey ResampleKernelBfyxCubicOpt::GetSupportedKey() const { k.EnableTensorOffset(); k.EnableTensorPitches(); k.EnableBatching(); + k.EnableDynamicShapesSupport(); k.EnableResampleType(ResampleType::CUBIC); return k; } @@ -70,9 +72,15 @@ bool ResampleKernelBfyxCubicOpt::Validate(const Params& p) const { if (params.inputs[0].Dimentions() != 4) DO_NOT_USE_THIS_KERNEL(p.layerID); - // Only spatial axes (Y, X) may be resized. + if (std::any_of(params.pads_begin.begin(), params.pads_begin.end(), [](const auto pad) { return pad != 0; }) || + std::any_of(params.pads_end.begin(), params.pads_end.end(), [](const auto pad) { return pad != 0; })) + DO_NOT_USE_THIS_KERNEL(p.layerID); + + // Explicit axes may include B/F with unit scale. The optimized kernel is still valid + // as long as only spatial dimensions actually change. for (const auto& axis : params.axes) { - if (axis != InterpolateAxis::Y && axis != InterpolateAxis::X) + if (axis != InterpolateAxis::BATCH && axis != InterpolateAxis::FEATURE && + axis != InterpolateAxis::Y && axis != InterpolateAxis::X) DO_NOT_USE_THIS_KERNEL(p.layerID); } diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_opt.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_opt.cpp index 1317bb44eb19..070df402be3d 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_opt.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_opt.cpp @@ -3,6 +3,8 @@ // #include "resample_kernel_opt.h" + +#include #include #include @@ -93,6 +95,113 @@ static int get_feature_slice_size(const resample_params ¶ms) { return static_cast(16 * get_vec_size(params)); } +static bool is_integral_ratio(size_t lhs, size_t rhs) { + return lhs != 0 && rhs != 0 && (lhs % rhs == 0 || rhs % lhs == 0); +} + +static bool is_integral_upsampling_ratio(size_t output, size_t input) { + return input != 0 && output >= input && output % input == 0; +} + +static bool is_asymmetric_simple_optimized_case(const resample_params& params) { + const auto& input = params.inputs[0]; + const auto& output = params.outputs[0]; + + if (params.coordTransMode != CoordinateTransformationMode::ASYMMETRIC || params.nearestMode != NearestMode::SIMPLE) { + return false; + } + + if (!is_integral_upsampling_ratio(output.X().v, input.X().v) || + !is_integral_upsampling_ratio(output.Y().v, input.Y().v)) { + return false; + } + + return input.Dimentions() != 5 || is_integral_upsampling_ratio(output.Z().v, input.Z().v); +} + +static bool is_tf_half_pixel_for_nn_floor_optimized_case(const resample_params& params) { + const auto& input = params.inputs[0]; + const auto& output = params.outputs[0]; + + if (params.coordTransMode != CoordinateTransformationMode::TF_HALF_PIXEL_FOR_NN || + params.nearestMode != NearestMode::FLOOR) { + return false; + } + + if (!is_integral_upsampling_ratio(output.X().v, input.X().v) || + !is_integral_upsampling_ratio(output.Y().v, input.Y().v)) { + return false; + } + + return input.Dimentions() != 5 || is_integral_upsampling_ratio(output.Z().v, input.Z().v); +} + +static bool is_half_pixel_round_prefer_floor_optimized_case(const resample_params& params) { + const auto& input = params.inputs[0]; + const auto& output = params.outputs[0]; + + if (params.coordTransMode != CoordinateTransformationMode::HALF_PIXEL || + params.nearestMode != NearestMode::ROUND_PREFER_FLOOR) { + return false; + } + + if (!is_integral_ratio(output.X().v, input.X().v) || !is_integral_ratio(output.Y().v, input.Y().v)) { + return false; + } + + return input.Dimentions() != 5 || is_integral_ratio(output.Z().v, input.Z().v); +} + +static int get_axis_index(InterpolateAxis axis) { + switch (axis) { + case InterpolateAxis::BATCH: + return 0; + case InterpolateAxis::FEATURE: + return 1; + case InterpolateAxis::Z: + return 2; + case InterpolateAxis::Y: + return 3; + case InterpolateAxis::X: + return 4; + default: + return 0; + } +} + +static std::vector get_legacy_scales(const resample_params& params) { + const auto& input = params.inputs[0]; + const auto& output = params.outputs[0]; + auto pads_begin = params.pads_begin; + auto pads_end = params.pads_end; + if (pads_begin.size() == 4) + pads_begin.insert(pads_begin.begin() + 2, 0); + if (pads_end.size() == 4) + pads_end.insert(pads_end.begin() + 2, 0); + + const auto b_size_padded = pads_begin[0] + input.Batch().v + pads_end[0]; + const auto f_size_padded = pads_begin[1] + input.Feature().v + pads_end[1]; + const auto z_size_padded = pads_begin[2] + input.Z().v + pads_end[2]; + const auto y_size_padded = pads_begin[3] + input.Y().v + pads_end[3]; + const auto x_size_padded = pads_begin[4] + input.X().v + pads_end[4]; + + std::vector scales = { + static_cast(b_size_padded) / static_cast(output.Batch().v), + static_cast(f_size_padded) / static_cast(output.Feature().v), + static_cast(z_size_padded) / static_cast(output.Z().v), + static_cast(y_size_padded) / static_cast(output.Y().v), + static_cast(x_size_padded) / static_cast(output.X().v), + }; + + for (std::size_t i = 0; i < params.axes.size(); i++) { + const int idx = get_axis_index(params.axes[i]); + if (params.shapeCalculationMode == kernel_selector::ShapeCalculationMode::SCALES) + scales[idx] = 1.f / params.scales[i]; + } + + return scales; +} + ResampleKernelBase::DispatchData ResampleKernelOpt::SetDefault(const kernel_selector::resample_params &arg) const { DispatchData dispatchData; auto in_layout = arg.inputs[0].GetLayout(); @@ -152,6 +261,10 @@ bool ResampleKernelOpt::Validate(const Params& p) const { DO_NOT_USE_THIS_KERNEL(p.layerID); const auto& input = params.inputs[0]; + const auto& output = params.outputs[0]; + + const auto has_padding = std::any_of(params.pads_begin.begin(), params.pads_begin.end(), [](const auto pad) { return pad != 0; }) || + std::any_of(params.pads_end.begin(), params.pads_end.end(), [](const auto pad) { return pad != 0; }); if ((input.GetDType() == Datatype::UINT8 || input.GetDType() == Datatype::INT8) && params.resampleType != ResampleType::NEAREST_NEIGHBOR && @@ -162,11 +275,37 @@ bool ResampleKernelOpt::Validate(const Params& p) const { if (input.Dimentions() == 5 && params.resampleType != ResampleType::NEAREST_NEIGHBOR) DO_NOT_USE_THIS_KERNEL(p.layerID); + if (params.resampleType == ResampleType::NEAREST_NEIGHBOR) { + const auto optimized_nearest_case = + (params.coordTransMode == CoordinateTransformationMode::ASYMMETRIC && params.nearestMode == NearestMode::FLOOR) || + is_asymmetric_simple_optimized_case(params) || + is_tf_half_pixel_for_nn_floor_optimized_case(params) || + is_half_pixel_round_prefer_floor_optimized_case(params); + + if (has_padding || + !optimized_nearest_case || + input.Batch().v != output.Batch().v || + input.Feature().v != output.Feature().v) { + DO_NOT_USE_THIS_KERNEL(p.layerID); + } + } + + if (params.resampleType == ResampleType::BILINEAR_INTERP) { + if (has_padding || + params.coordTransMode != CoordinateTransformationMode::ASYMMETRIC || + input.Batch().v != output.Batch().v || + input.Feature().v != output.Feature().v) { + DO_NOT_USE_THIS_KERNEL(p.layerID); + } + } + return true; } JitConstants ResampleKernelOpt::GetJitConstants(const resample_params ¶ms) const { auto jit = Parent::GetJitConstants(params); + jit.RemoveConstant("SCALES"); + jit.AddConstant(MakeJitConstant("SCALES", get_legacy_scales(params))); auto opt_x_block_size = GetOptimalBlockSize(params); if (params.outputs[0].X().v > 32 && opt_x_block_size == 1) { diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_pil_ref.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_pil_ref.cpp index ff74b6a5e28b..9abfa2b382d9 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_pil_ref.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_pil_ref.cpp @@ -129,7 +129,7 @@ DataTensor GetIntermediateBufferSize(const resample_params& params) { OPENVINO_ASSERT(channelIndex >= 0, "Invalid layout channel index"); dims[channelIndex] = ybox_last - ybox_first; - DataTensor result{dims, Datatype::F16, layout}; + DataTensor result{dims, params.outputs[0].GetDType(), layout}; return result; } @@ -421,7 +421,7 @@ JitConstants ResampleKernelPilRef::GetJitConstantsForKernel(KernelId id, const r KernelsData ResampleKernelPilRef::GetKernelsData(const Params ¶ms) const { const resample_params& resample_parameters = static_cast(params); KernelData kd = KernelData::Default(params, GetKernelsNum(resample_parameters)); - kd.internalBufferDataType = Datatype::F16; + kd.internalBufferDataType = resample_parameters.outputs[0].GetDType(); int i = 0; for (ResampleKernelPilRef::KernelId id = eCalcHorizontalCoefficients; id < eEnd; ++id) { if (!NeedHorizontalPass(resample_parameters) && diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_ref.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_ref.cpp index 557390094a17..09706c3ddc81 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_ref.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_ref.cpp @@ -91,9 +91,152 @@ static bool use_packing(const resample_params& params) { return true; } +static bool has_padding(const resample_params& params) { + return std::any_of(params.pads_begin.begin(), params.pads_begin.end(), [](const auto pad) { return pad != 0; }) || + std::any_of(params.pads_end.begin(), params.pads_end.end(), [](const auto pad) { return pad != 0; }); +} + +static bool is_integral_ratio(size_t lhs, size_t rhs) { + return lhs != 0 && rhs != 0 && (lhs % rhs == 0 || rhs % lhs == 0); +} + +static bool is_integral_upsampling_ratio(size_t output, size_t input) { + return input != 0 && output >= input && output % input == 0; +} + +static bool is_fast_nearest_case(const resample_params& params) { + const auto& input = params.inputs[0]; + const auto& output = params.outputs[0]; + + if (params.resampleType != ResampleType::NEAREST_NEIGHBOR || has_padding(params) || + input.Batch().v != output.Batch().v || input.Feature().v != output.Feature().v) { + return false; + } + + const auto asymmetric_floor = params.coordTransMode == CoordinateTransformationMode::ASYMMETRIC && + params.nearestMode == NearestMode::FLOOR; + const auto asymmetric_simple_upsampling = + params.coordTransMode == CoordinateTransformationMode::ASYMMETRIC && + params.nearestMode == NearestMode::SIMPLE && + is_integral_upsampling_ratio(output.X().v, input.X().v) && + is_integral_upsampling_ratio(output.Y().v, input.Y().v) && + (input.Dimentions() != 5 || is_integral_upsampling_ratio(output.Z().v, input.Z().v)); + const auto tf_half_pixel_for_nn_floor_upsampling = + params.coordTransMode == CoordinateTransformationMode::TF_HALF_PIXEL_FOR_NN && + params.nearestMode == NearestMode::FLOOR && + is_integral_upsampling_ratio(output.X().v, input.X().v) && + is_integral_upsampling_ratio(output.Y().v, input.Y().v) && + (input.Dimentions() != 5 || is_integral_upsampling_ratio(output.Z().v, input.Z().v)); + const auto half_pixel_round_prefer_floor = + params.coordTransMode == CoordinateTransformationMode::HALF_PIXEL && + params.nearestMode == NearestMode::ROUND_PREFER_FLOOR && + is_integral_ratio(output.X().v, input.X().v) && + is_integral_ratio(output.Y().v, input.Y().v) && + (input.Dimentions() != 5 || is_integral_ratio(output.Z().v, input.Z().v)); + + return asymmetric_floor || asymmetric_simple_upsampling || tf_half_pixel_for_nn_floor_upsampling || + half_pixel_round_prefer_floor; +} + +static bool is_fast_linear_onnx_case(const resample_params& params) { + const auto& input = params.inputs[0]; + const auto& output = params.outputs[0]; + + if (params.resampleType != ResampleType::LINEAR_ONNX || has_padding(params) || + params.coordTransMode != CoordinateTransformationMode::HALF_PIXEL || + input.Batch().v != output.Batch().v || input.Feature().v != output.Feature().v) { + return false; + } + + if (!is_integral_upsampling_ratio(output.X().v, input.X().v) || + !is_integral_upsampling_ratio(output.Y().v, input.Y().v)) { + return false; + } + + return input.Dimentions() != 5 || is_integral_upsampling_ratio(output.Z().v, input.Z().v); +} + +static bool is_fast_caffe_bilinear_interp_case(const resample_params& params) { + const auto& input = params.inputs[0]; + const auto& output = params.outputs[0]; + + return params.resampleType == ResampleType::CAFFE_BILINEAR_INTERP && + !has_padding(params) && + input.Batch().v == output.Batch().v && + input.Feature().v == output.Feature().v; +} + +static int get_axis_index(InterpolateAxis axis) { + switch (axis) { + case InterpolateAxis::BATCH: + return 0; + case InterpolateAxis::FEATURE: + return 1; + case InterpolateAxis::Z: + return 2; + case InterpolateAxis::Y: + return 3; + case InterpolateAxis::X: + return 4; + default: + return 0; + } +} + +static std::vector get_legacy_scales(const resample_params& params) { + const auto& input = params.inputs[0]; + const auto& output = params.outputs[0]; + auto pads_begin = params.pads_begin; + auto pads_end = params.pads_end; + if (pads_begin.size() == 4) + pads_begin.insert(pads_begin.begin() + 2, 0); + if (pads_end.size() == 4) + pads_end.insert(pads_end.begin() + 2, 0); + + const auto b_size_padded = pads_begin[0] + input.Batch().v + pads_end[0]; + const auto f_size_padded = pads_begin[1] + input.Feature().v + pads_end[1]; + const auto z_size_padded = pads_begin[2] + input.Z().v + pads_end[2]; + const auto y_size_padded = pads_begin[3] + input.Y().v + pads_end[3]; + const auto x_size_padded = pads_begin[4] + input.X().v + pads_end[4]; + + std::vector scales = { + static_cast(b_size_padded) / static_cast(output.Batch().v), + static_cast(f_size_padded) / static_cast(output.Feature().v), + static_cast(z_size_padded) / static_cast(output.Z().v), + static_cast(y_size_padded) / static_cast(output.Y().v), + static_cast(x_size_padded) / static_cast(output.X().v), + }; + + for (std::size_t i = 0; i < params.axes.size(); i++) { + const int idx = get_axis_index(params.axes[i]); + if (params.shapeCalculationMode == kernel_selector::ShapeCalculationMode::SCALES) + scales[idx] = 1.f / params.scales[i]; + } + + return scales; +} + JitConstants ResampleKernelRef::GetJitConstants(const resample_params& params) const { JitConstants jit = ResampleKernelBase::GetJitConstants(params); + if (is_fast_nearest_case(params)) { + jit.RemoveConstant("SCALES"); + jit.AddConstant(MakeJitConstant("SCALES", get_legacy_scales(params))); + jit.AddConstant(MakeJitConstant("RESAMPLE_FAST_NEAREST", 1)); + } + + if (is_fast_linear_onnx_case(params)) { + jit.RemoveConstant("SCALES"); + jit.AddConstant(MakeJitConstant("SCALES", get_legacy_scales(params))); + jit.AddConstant(MakeJitConstant("RESAMPLE_USE_LEGACY_SCALE", 1)); + } + + if (is_fast_caffe_bilinear_interp_case(params)) { + jit.RemoveConstant("SCALES"); + jit.AddConstant(MakeJitConstant("SCALES", get_legacy_scales(params))); + jit.AddConstant(MakeJitConstant("RESAMPLE_USE_LEGACY_SCALE", 1)); + } + if (use_packing(params)) { jit.AddConstant(MakeJitConstant("PACK_SIZE", packing_factor(params))); jit.AddConstant(MakeJitConstant("FEATURE_PACKED_MODE", "1")); diff --git a/src/plugins/intel_gpu/src/plugin/ops/interpolate.cpp b/src/plugins/intel_gpu/src/plugin/ops/interpolate.cpp index b9f61bc10e8b..45a4513f022f 100644 --- a/src/plugins/intel_gpu/src/plugin/ops/interpolate.cpp +++ b/src/plugins/intel_gpu/src/plugin/ops/interpolate.cpp @@ -21,11 +21,14 @@ static std::vector ExtractAxes(const std::shared_ptrget_friendly_name(), " (", op->get_type_name(), ")"); axes = axes_constant->cast_vector(); - ov::util::try_normalize_axes(axes, inputRank, *op); - } else { + } + + if (axes.empty()) { for (size_t i = 0; i < inputRank; ++i) { axes.push_back(ov::util::try_normalize_axis(i, inputRank, *op)); } + } else { + ov::util::try_normalize_axes(axes, inputRank, *op); } return axes; } diff --git a/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/interpolate.cpp b/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/interpolate.cpp index 642c9b20ba82..7e4c2648a379 100644 --- a/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/interpolate.cpp +++ b/src/plugins/intel_gpu/tests/functional/shared_tests_instances/single_layer_tests/interpolate.cpp @@ -20,10 +20,10 @@ class GPUInterpolateLayerTest : public InterpolateLayerTest { // Some rounding float to integer types on GPU may differ from CPU, and as result, // the actual values may differ from reference ones on 1 when the float is very close to an integer, // e.g 6,0000023 calculated on CPU may be cast to 5 by OpenCL convert_uchar function. - // That is why the threshold is set 1.f for integer types. + // That is why the absolute threshold is set 1.f for integer types. if (targetDevice == "GPU" && (model_type == ov::element::u8 || model_type == ov::element::i8)) { - rel_threshold = 1.f; + abs_threshold = 1.f; } } }; diff --git a/src/plugins/intel_gpu/tests/functional/single_layer_tests/dynamic/interpolate.cpp b/src/plugins/intel_gpu/tests/functional/single_layer_tests/dynamic/interpolate.cpp index 271f09338f46..a98ebe2546be 100644 --- a/src/plugins/intel_gpu/tests/functional/single_layer_tests/dynamic/interpolate.cpp +++ b/src/plugins/intel_gpu/tests/functional/single_layer_tests/dynamic/interpolate.cpp @@ -152,6 +152,10 @@ class InterpolateLayerGPUTest : public testing::WithParamInterface(model_type, inputDynamicShapes.front()); - auto scales_orsizes_input = make_scales_or_sizes_input(shape_calc_mode, target_shape, scales); + auto target_sizes = target_shape; + if (!axes.empty()) { + target_sizes.clear(); + for (const auto axis : axes) { + target_sizes.push_back(target_shape[axis]); + } + } + + auto scales_orsizes_input = make_scales_or_sizes_input(shape_calc_mode, target_sizes, scales); ov::op::util::InterpolateBase::InterpolateAttrs interpolate_attributes{mode, shape_calc_mode, pad_begin, pad_end, coordinate_transform_mode, nearest_mode, antialias, cube_coef}; @@ -165,6 +173,10 @@ void Interpolate11LayerTest::SetUp() { auto result = std::make_shared(interpolate); function = std::make_shared(result, ov::ParameterVector{param}, "interpolate"); + + if (model_type == ov::element::f32) { + abs_threshold = 1e-6; + } } } // namespace test } // namespace ov