Skip to content

[GPU] Fix for Interpolate func tests on PTL #36752

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
nshchego wants to merge 1 commit into
base: master
Choose a base branch
from
+495 −47
Open

[GPU] Fix for Interpolate func tests on PTL #36752

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
64 changes: 50 additions & 14 deletions src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_bfyx_cubic_opt.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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)(
Expand All @@ -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];
Expand All @@ -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];
Expand All @@ -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;
}
}
}
Expand Down
14 changes: 9 additions & 5 deletions src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_onnx.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
6 changes: 5 additions & 1 deletion src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_opt.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
111 changes: 99 additions & 12 deletions src/plugins/intel_gpu/src/kernel_selector/cl_kernels/resample_ref.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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
Expand All @@ -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)))
Expand Down Expand Up @@ -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];
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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]) {
Expand All @@ -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);
}
}
}
}
Expand Down
20 changes: 14 additions & 6 deletions src/plugins/intel_gpu/src/kernel_selector/kernels/resample/resample_kernel_base.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -146,17 +146,17 @@ JitConstants ResampleKernelBase::GetJitConstants(const resample_params& params)
paddingUsed |= (pads_begin[i] != 0 || pads_end[i] != 0);
}

scales[0] = static_cast<float>(b_size_padded) / static_cast<float>(out_b_size_padded);
scales[1] = static_cast<float>(f_size_padded) / static_cast<float>(out_f_size_padded);
scales[4] = static_cast<float>(x_size_padded) / static_cast<float>(out_x_size_padded);
scales[3] = static_cast<float>(y_size_padded) / static_cast<float>(out_y_size_padded);
scales[2] = static_cast<float>(z_size_padded) / static_cast<float>(out_z_size_padded);
scales[0] = static_cast<float>(out_b_size_padded) / static_cast<float>(b_size_padded);
scales[1] = static_cast<float>(out_f_size_padded) / static_cast<float>(f_size_padded);
scales[4] = static_cast<float>(out_x_size_padded) / static_cast<float>(x_size_padded);
scales[3] = static_cast<float>(out_y_size_padded) / static_cast<float>(y_size_padded);
scales[2] = static_cast<float>(out_z_size_padded) / static_cast<float>(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)
Expand All @@ -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),
Expand Down Expand Up @@ -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};
}
Expand Down
Loading
Loading