diff --git a/doc/primitives/eltwise.md b/doc/primitives/eltwise.md index ec190d938b3..1e30b584353 100644 --- a/doc/primitives/eltwise.md +++ b/doc/primitives/eltwise.md @@ -25,7 +25,6 @@ The following operations are supported: | Operation | oneDNN algorithm kind | Forward formula | Backward formula (from src) | Backward formula (from dst) | | :-- | :-- | :-- | :-- | :-- | | abs | #dnnl_eltwise_abs | \f$ d = \begin{cases} s & \text{if}\ s > 0 \\ -s & \text{if}\ s \leq 0 \end{cases} \f$ | \f$ ds = \begin{cases} dd & \text{if}\ s > 0 \\ -dd & \text{if}\ s < 0 \\ 0 & \text{if}\ s = 0 \end{cases} \f$ | -- | -| bounded_relu | #dnnl_eltwise_bounded_relu | \f$ d = \begin{cases} \alpha & \text{if}\ s > \alpha \geq 0 \\ s & \text{if}\ 0 < s \leq \alpha \\ 0 & \text{if}\ s \leq 0 \end{cases} \f$ | \f$ ds = \begin{cases} dd & \text{if}\ 0 < s \leq \alpha, \\ 0 & \text{otherwise}\ \end{cases} \f$ | -- | | clip | #dnnl_eltwise_clip | \f$ d = \begin{cases} \beta & \text{if}\ s > \beta \geq \alpha \\ s & \text{if}\ \alpha < s \leq \beta \\ \alpha & \text{if}\ s \leq \alpha \end{cases} \f$ | \f$ ds = \begin{cases} dd & \text{if}\ \alpha < s \leq \beta \\ 0 & \text{otherwise}\ \end{cases} \f$ | -- | | clip_v2 | #dnnl_eltwise_clip_v2
#dnnl_eltwise_clip_v2_use_dst_for_bwd | \f$ d = \begin{cases} \beta & \text{if}\ s \geq \beta \geq \alpha \\ s & \text{if}\ \alpha < s < \beta \\ \alpha & \text{if}\ s \leq \alpha \end{cases} \f$ | \f$ ds = \begin{cases} dd & \text{if}\ \alpha < s < \beta \\ 0 & \text{otherwise}\ \end{cases} \f$ | \f$ ds = \begin{cases} dd & \text{if}\ \alpha < d < \beta \\ 0 & \text{otherwise}\ \end{cases} \f$ | | elu | #dnnl_eltwise_elu
#dnnl_eltwise_elu_use_dst_for_bwd | \f$ d = \begin{cases} s & \text{if}\ s > 0 \\ \alpha (e^s - 1) & \text{if}\ s \leq 0 \end{cases} \f$ | \f$ ds = \begin{cases} dd & \text{if}\ s > 0 \\ dd \cdot \alpha e^s & \text{if}\ s \leq 0 \end{cases} \f$ | \f$ ds = \begin{cases} dd & \text{if}\ d > 0 \\ dd \cdot (d + \alpha) & \text{if}\ d \leq 0 \end{cases}. See\ (2). \f$ | @@ -54,6 +53,7 @@ The following operations are supported: \f$ (3)\ \text{where, } \omega = e^{3s} + 4 \cdot e^{2s} + e^{s} \cdot (4 \cdot s + 6) + 4 \cdot (s + 1) \text{ and } \delta = e^{2s} + 2 \cdot e^{s} + 2. \f$ Note that following equations hold: +* \f$ bounded\_relu(s, alpha) = clip(s, 0, alpha) \f$ * \f$ logsigmoid(s) = soft\_relu(s, -1) \f$ #### Difference Between Forward Training and Forward Inference diff --git a/include/oneapi/dnnl/dnnl.hpp b/include/oneapi/dnnl/dnnl.hpp index b3cb31fdba6..a4fdd710e4f 100644 --- a/include/oneapi/dnnl/dnnl.hpp +++ b/include/oneapi/dnnl/dnnl.hpp @@ -543,8 +543,6 @@ enum class algorithm { eltwise_swish = dnnl_eltwise_swish, /// Elementwise: linear eltwise_linear = dnnl_eltwise_linear, - /// Elementwise: bounded_relu - eltwise_bounded_relu = dnnl_eltwise_bounded_relu, /// Elementwise: soft_relu eltwise_soft_relu = dnnl_eltwise_soft_relu, /// Elementwise: mish diff --git a/include/oneapi/dnnl/dnnl_types.h b/include/oneapi/dnnl/dnnl_types.h index 9cb06f27432..963eba80d5d 100644 --- a/include/oneapi/dnnl/dnnl_types.h +++ b/include/oneapi/dnnl/dnnl_types.h @@ -1454,8 +1454,6 @@ typedef enum { dnnl_eltwise_sqrt = 0x6f, /// Eltwise: linear dnnl_eltwise_linear = 0x7f, - /// Eltwise: bounded_relu - dnnl_eltwise_bounded_relu = 0x8f, /// Eltwise: soft_relu dnnl_eltwise_soft_relu = 0xa0, /// Eltwise: hardsigmoid diff --git a/src/common/c_types_map.hpp b/src/common/c_types_map.hpp index 5264c343f55..36de3b4aab8 100644 --- a/src/common/c_types_map.hpp +++ b/src/common/c_types_map.hpp @@ -76,7 +76,6 @@ const alg_kind_t eltwise_abs = dnnl_eltwise_abs; const alg_kind_t eltwise_sqrt = dnnl_eltwise_sqrt; const alg_kind_t eltwise_swish = dnnl_eltwise_swish; const alg_kind_t eltwise_linear = dnnl_eltwise_linear; -const alg_kind_t eltwise_bounded_relu = dnnl_eltwise_bounded_relu; const alg_kind_t eltwise_soft_relu = dnnl_eltwise_soft_relu; const alg_kind_t eltwise_logistic = dnnl_eltwise_logistic; const alg_kind_t eltwise_mish = dnnl_eltwise_mish; diff --git a/src/common/dnnl_debug_autogenerated.cpp b/src/common/dnnl_debug_autogenerated.cpp index 17378ee87c6..462baa9e4fc 100644 --- a/src/common/dnnl_debug_autogenerated.cpp +++ b/src/common/dnnl_debug_autogenerated.cpp @@ -1130,7 +1130,6 @@ const char *dnnl_alg_kind2str(dnnl_alg_kind_t v) { if (v == dnnl_eltwise_abs) return "eltwise_abs"; if (v == dnnl_eltwise_sqrt) return "eltwise_sqrt"; if (v == dnnl_eltwise_linear) return "eltwise_linear"; - if (v == dnnl_eltwise_bounded_relu) return "eltwise_bounded_relu"; if (v == dnnl_eltwise_soft_relu) return "eltwise_soft_relu"; if (v == dnnl_eltwise_hardsigmoid) return "eltwise_hardsigmoid"; if (v == dnnl_eltwise_logistic) return "eltwise_logistic"; diff --git a/src/common/eltwise_pd.hpp b/src/common/eltwise_pd.hpp index 56060ec8f8f..aeecc9391b9 100644 --- a/src/common/eltwise_pd.hpp +++ b/src/common/eltwise_pd.hpp @@ -138,8 +138,8 @@ struct eltwise_fwd_pd_t : public eltwise_pd_t { using namespace utils; return one_of(alg, eltwise_relu, eltwise_tanh, eltwise_elu, eltwise_square, eltwise_abs, eltwise_sqrt, eltwise_swish, - eltwise_bounded_relu, eltwise_gelu_tanh, - eltwise_gelu_erf, eltwise_round, eltwise_hardswish) + eltwise_gelu_tanh, eltwise_gelu_erf, eltwise_round, + eltwise_hardswish) || one_of(alg, eltwise_relu_use_dst_for_bwd, eltwise_tanh_use_dst_for_bwd, eltwise_elu_use_dst_for_bwd, @@ -221,12 +221,12 @@ struct eltwise_bwd_pd_t : public eltwise_pd_t { // zero, so excluding all those algs from here. using namespace alg_kind; using namespace utils; - return one_of(alg, eltwise_abs, eltwise_bounded_relu, eltwise_clip, - eltwise_clip_v2, eltwise_elu, eltwise_exp, - eltwise_gelu_erf, eltwise_gelu_tanh, eltwise_hardsigmoid, - eltwise_linear, eltwise_logistic, eltwise_mish, - eltwise_relu, eltwise_soft_relu, eltwise_square, - eltwise_swish, eltwise_tanh) + return one_of(alg, eltwise_abs, eltwise_clip, eltwise_clip_v2, + eltwise_elu, eltwise_exp, eltwise_gelu_erf, + eltwise_gelu_tanh, eltwise_hardsigmoid, eltwise_linear, + eltwise_logistic, eltwise_mish, eltwise_relu, + eltwise_soft_relu, eltwise_square, eltwise_swish, + eltwise_tanh) || one_of(alg, eltwise_elu_use_dst_for_bwd, eltwise_exp_use_dst_for_bwd, eltwise_logistic_use_dst_for_bwd, diff --git a/src/common/math_utils.hpp b/src/common/math_utils.hpp index a1d7ba095b8..38a74a6e268 100644 --- a/src/common/math_utils.hpp +++ b/src/common/math_utils.hpp @@ -214,18 +214,6 @@ inline U linear_bwd(T dd, T s, A alpha, A beta) { return (U)(dd * alpha); } -template ::type> -inline U bounded_relu_fwd(T s, A alpha) { - s = s > 0 ? s : (U)0; - return s > alpha ? (U)(alpha) : s; -} -template ::type> -inline U bounded_relu_bwd(T dd, T s, A alpha) { - return dd * (0 < s && s <= alpha ? 1 : 0); -} - template ::type> inline U logistic_fwd(T s) { // Here we avoid division/inverse by infinity as some architectures have @@ -395,7 +383,7 @@ inline U hardsigmoid_bwd(T dd, T s, A alpha, A beta) { template ::type> inline U hardswish_fwd(T s) { - return (s / 6.f) * bounded_relu_fwd(s + 3.f, 6.f); + return (s / 6.f) * clip_fwd(s + 3.f, 0.f, 6.f); } template ::type> inline U hardswish_bwd(T dd, T s) { @@ -411,12 +399,11 @@ inline bool is_eltwise_ok( const bool eltwise_use_src = one_of(alg, eltwise_relu, eltwise_tanh, eltwise_elu, eltwise_square, eltwise_abs, eltwise_sqrt, eltwise_linear, - eltwise_bounded_relu, eltwise_soft_relu, eltwise_mish, - eltwise_logistic, eltwise_exp, eltwise_gelu_tanh, - eltwise_hardsigmoid, eltwise_hardswish, eltwise_swish, - eltwise_log, eltwise_clip, eltwise_clip_v2, eltwise_pow, + eltwise_soft_relu, eltwise_mish, eltwise_logistic, + eltwise_exp, eltwise_gelu_tanh, eltwise_hardsigmoid, + eltwise_hardswish, eltwise_swish, eltwise_log, + eltwise_clip, eltwise_clip_v2, eltwise_pow, eltwise_gelu_erf, eltwise_round) - && IMPLICATION(alg == eltwise_bounded_relu, alpha >= 0) && IMPLICATION( one_of(alg, eltwise_clip, eltwise_clip_v2), beta >= alpha) && IMPLICATION(alg == eltwise_round, dt == dnnl_f32) diff --git a/src/common/opdesc.hpp b/src/common/opdesc.hpp index a48dc0ababb..3c9d8c94ed1 100644 --- a/src/common/opdesc.hpp +++ b/src/common/opdesc.hpp @@ -195,7 +195,6 @@ struct eltwise_desc_t { // The kind of eltwise algorithm. Possible values: #dnnl_eltwise_relu, // #dnnl_eltwise_tanh, #dnnl_eltwise_elu, #dnnl_eltwise_square, // #dnnl_eltwise_abs, #dnnl_eltwise_sqrt, #dnnl_eltwise_linear, - // #dnnl_eltwise_bounded_relu, // #dnnl_eltwise_soft_relu, #dnnl_eltwise_logistic, #dnnl_eltwise_exp, // #dnnl_eltwise_gelu_tanh, #dnnl_eltwise_swish, #dnnl_eltwise_log, // #dnnl_eltwise_clip, #dnnl_eltwise_clip_v2, #dnnl_eltwise_pow, @@ -221,7 +220,6 @@ struct eltwise_desc_t { // - #dnnl_eltwise_abs: @p alpha and @p beta ignored // - #dnnl_eltwise_sqrt: @p alpha and @p beta ignored // - #dnnl_eltwise_linear: @p alpha -- scale, @p beta -- shift - // - #dnnl_eltwise_bounded_relu: @p alpha -- upper bound, @p beta ignored // - #dnnl_eltwise_soft_relu: @p alpha -- soft_relu arg scaling, @p beta ignored // - #dnnl_eltwise_logistic: @p alpha and @p beta ignored // - #dnnl_eltwise_exp: @p alpha and @p beta ignored diff --git a/src/cpu/aarch64/acl_utils.cpp b/src/cpu/aarch64/acl_utils.cpp index 79ea775d6db..b698657b4f4 100644 --- a/src/cpu/aarch64/acl_utils.cpp +++ b/src/cpu/aarch64/acl_utils.cpp @@ -86,9 +86,6 @@ status_t convert_to_acl_act(alg_kind_t eltwise_alg, float alpha, float beta, case eltwise_linear: act_info = ActivationLayerInfo(act_func::LINEAR, alpha, beta); break; - case eltwise_bounded_relu: - act_info = ActivationLayerInfo(act_func::BOUNDED_RELU, alpha, beta); - break; case eltwise_soft_relu: act_info = ActivationLayerInfo(act_func::SOFT_RELU, alpha, beta); break; diff --git a/src/cpu/aarch64/injectors/jit_uni_eltwise_injector.cpp b/src/cpu/aarch64/injectors/jit_uni_eltwise_injector.cpp index e3945d12ffa..29976fdee60 100644 --- a/src/cpu/aarch64/injectors/jit_uni_eltwise_injector.cpp +++ b/src/cpu/aarch64/injectors/jit_uni_eltwise_injector.cpp @@ -40,7 +40,7 @@ bool is_alg_supported(alg_kind_t alg) { using namespace alg_kind; return utils::one_of(alg, eltwise_relu, eltwise_tanh, eltwise_elu, eltwise_square, eltwise_abs, eltwise_sqrt, eltwise_linear, - eltwise_bounded_relu, /*eltwise_soft_relu,*/ + /*eltwise_soft_relu,*/ eltwise_logistic, /*eltwise_mish,*/ eltwise_exp, eltwise_gelu_tanh, /*eltwise_hardswish,*/ eltwise_swish, eltwise_log, eltwise_clip, /*eltwise_clip_v2, eltwise_pow,*/ eltwise_gelu_erf, eltwise_round, @@ -256,7 +256,6 @@ void jit_uni_eltwise_injector_f32::set_coef_to_regs() { table_val(alpha, z_tmp); table_val(beta, vmm_aux0); break; - case eltwise_bounded_relu: table_val(alpha, z_tmp); break; // case eltwise_soft_relu: // TODO: enable me. case eltwise_logistic_use_dst_for_bwd: case eltwise_logistic: @@ -286,7 +285,6 @@ void jit_uni_eltwise_injector_f32::set_coef_to_regs() { case eltwise_sqrt_use_dst_for_bwd: case eltwise_sqrt: case eltwise_linear: - case eltwise_bounded_relu: // case eltwise_soft_relu: case eltwise_logistic_use_dst_for_bwd: case eltwise_logistic: @@ -630,13 +628,6 @@ void jit_uni_eltwise_injector_f32::linear_compute_vector_fwd( h->fmad(vmm_src, p_all / T_m, z_tmp, vmm_aux0); } -template -void jit_uni_eltwise_injector_f32::bounded_relu_compute_vector_fwd( - const TRegS &vmm_src) { - h->fmaxnm(vmm_src, p_all, 0.f); - h->fminnm(vmm_src, p_all, z_tmp); -} - template void jit_uni_eltwise_injector_f32::clip_compute_vector_fwd( const TRegS &vmm_src) { @@ -1090,21 +1081,6 @@ void jit_uni_eltwise_injector_f32::linear_compute_vector_bwd( h->mov(ZRegD(IDX(vmm_src)), ZRegD(IDX(table_val(alpha, z_tmp)))); } -template -void jit_uni_eltwise_injector_f32::bounded_relu_compute_vector_bwd( - const TRegS &vmm_src) { - // get mask of values > alpha and blend with 0.f - compute_cmp_mask(vmm_src, table_val(alpha, z_tmp), _cmp_gt_os); - blend_with_mask(vmm_src, table_val(zero, z_tmp)); - // make all negative values zeros - h->fmov(z_tmp, 0.f); - h->fmaxnm(vmm_src, p_all, z_tmp); - - // everything bigger than 0.f should be 1.f - compute_cmp_mask(vmm_src, table_val(zero, z_tmp), _cmp_gt_os); - blend_with_mask(vmm_src, table_val(one, z_tmp)); -} - template void jit_uni_eltwise_injector_f32::soft_relu_compute_vector_bwd( const TRegS &vmm_src) { @@ -1296,7 +1272,6 @@ size_t jit_uni_eltwise_injector_f32::aux_vecs_count() { case eltwise_sqrt_use_dst_for_bwd: case eltwise_sqrt: return 0; case eltwise_linear: return 2; - case eltwise_bounded_relu: return 1; // case eltwise_soft_relu: return 5; case eltwise_logistic_use_dst_for_bwd: case eltwise_logistic: return 5; /* = exp + 1 */ @@ -1323,7 +1298,6 @@ size_t jit_uni_eltwise_injector_f32::aux_vecs_count() { case eltwise_sqrt_use_dst_for_bwd: case eltwise_sqrt: return 2; case eltwise_linear: return 1; - case eltwise_bounded_relu: return 1; // case eltwise_soft_relu: return 5; /* = logistic */ case eltwise_logistic_use_dst_for_bwd: return 2; case eltwise_logistic: return 5; /* = logistic */ @@ -1370,9 +1344,6 @@ void jit_uni_eltwise_injector_f32::compute_body( case eltwise_linear: linear_compute_vector_fwd(TRegS(idx)); break; - case eltwise_bounded_relu: - bounded_relu_compute_vector_fwd(TRegS(idx)); - break; // case eltwise_soft_relu: // soft_relu_compute_vector_fwd(TRegS(idx)); // break; @@ -1410,9 +1381,6 @@ void jit_uni_eltwise_injector_f32::compute_body( case eltwise_linear: linear_compute_vector_bwd(TRegS(idx)); break; - case eltwise_bounded_relu: - bounded_relu_compute_vector_bwd(TRegS(idx)); - break; // case eltwise_soft_relu: // soft_relu_compute_vector_bwd(TRegS(idx)); // break; diff --git a/src/cpu/aarch64/injectors/jit_uni_eltwise_injector.hpp b/src/cpu/aarch64/injectors/jit_uni_eltwise_injector.hpp index 79d772ab6ec..ee04bc42d0b 100644 --- a/src/cpu/aarch64/injectors/jit_uni_eltwise_injector.hpp +++ b/src/cpu/aarch64/injectors/jit_uni_eltwise_injector.hpp @@ -120,13 +120,12 @@ struct jit_uni_eltwise_injector_f32 { assert(is_superset(isa, sve_128)); assert(utils::one_of(alg_, eltwise_relu, eltwise_tanh, eltwise_elu, eltwise_square, eltwise_abs, eltwise_sqrt, eltwise_linear, - eltwise_bounded_relu, eltwise_logistic, eltwise_exp, - eltwise_gelu_tanh, eltwise_swish, eltwise_log, eltwise_clip, - eltwise_clip_v2, eltwise_gelu_erf, eltwise_round, - eltwise_relu_use_dst_for_bwd, eltwise_tanh_use_dst_for_bwd, - eltwise_elu_use_dst_for_bwd, eltwise_sqrt_use_dst_for_bwd, - eltwise_logistic_use_dst_for_bwd, eltwise_exp_use_dst_for_bwd, - eltwise_clip_v2_use_dst_for_bwd)); + eltwise_logistic, eltwise_exp, eltwise_gelu_tanh, eltwise_swish, + eltwise_log, eltwise_clip, eltwise_clip_v2, eltwise_gelu_erf, + eltwise_round, eltwise_relu_use_dst_for_bwd, + eltwise_tanh_use_dst_for_bwd, eltwise_elu_use_dst_for_bwd, + eltwise_sqrt_use_dst_for_bwd, eltwise_logistic_use_dst_for_bwd, + eltwise_exp_use_dst_for_bwd, eltwise_clip_v2_use_dst_for_bwd)); register_table_entries(); } @@ -227,7 +226,6 @@ struct jit_uni_eltwise_injector_f32 { void abs_compute_vector_fwd(const TRegS &vmm_src); void sqrt_compute_vector_fwd(const TRegS &vmm_src); void linear_compute_vector_fwd(const TRegS &vmm_src); - void bounded_relu_compute_vector_fwd(const TRegS &vmm_src); void soft_relu_compute_vector_fwd(const TRegS &vmm_src); void logistic_compute_vector_fwd(const TRegS &vmm_src); void gelu_tanh_compute_vector_fwd(const TRegS &vmm_src); @@ -245,7 +243,6 @@ struct jit_uni_eltwise_injector_f32 { void abs_compute_vector_bwd(const TRegS &vmm_src); void sqrt_compute_vector_bwd(const TRegS &vmm_src); void linear_compute_vector_bwd(const TRegS &vmm_src); - void bounded_relu_compute_vector_bwd(const TRegS &vmm_src); void soft_relu_compute_vector_bwd(const TRegS &vmm_src); void logistic_compute_vector_bwd(const TRegS &vmm_src); void gelu_tanh_compute_vector_bwd(const TRegS &vmm_src); diff --git a/src/cpu/aarch64/jit_uni_eltwise.cpp b/src/cpu/aarch64/jit_uni_eltwise.cpp index be69abbc666..56d58ea16a1 100644 --- a/src/cpu/aarch64/jit_uni_eltwise.cpp +++ b/src/cpu/aarch64/jit_uni_eltwise.cpp @@ -203,11 +203,10 @@ status_t jit_uni_eltwise_fwd_t::pd_t::init(engine_t *engine) { eltwise_relu, eltwise_elu_use_dst_for_bwd, eltwise_elu, eltwise_tanh_use_dst_for_bwd, eltwise_tanh, eltwise_square, eltwise_abs, eltwise_sqrt_use_dst_for_bwd, eltwise_sqrt, - eltwise_linear, eltwise_bounded_relu, eltwise_soft_relu, - eltwise_logistic_use_dst_for_bwd, eltwise_logistic, - eltwise_exp_use_dst_for_bwd, eltwise_exp, eltwise_gelu_tanh, - eltwise_swish, eltwise_log, eltwise_clip, eltwise_gelu_erf, - eltwise_round); + eltwise_linear, eltwise_soft_relu, eltwise_logistic_use_dst_for_bwd, + eltwise_logistic, eltwise_exp_use_dst_for_bwd, eltwise_exp, + eltwise_gelu_tanh, eltwise_swish, eltwise_log, eltwise_clip, + eltwise_gelu_erf, eltwise_round); return ok ? status::success : status::unimplemented; } @@ -279,10 +278,10 @@ status_t jit_uni_eltwise_bwd_t::pd_t::init(engine_t *engine) { eltwise_relu, eltwise_elu_use_dst_for_bwd, eltwise_elu, eltwise_tanh_use_dst_for_bwd, eltwise_tanh, eltwise_square, eltwise_abs, eltwise_sqrt_use_dst_for_bwd, eltwise_sqrt, - eltwise_linear, eltwise_bounded_relu, eltwise_soft_relu, - eltwise_logistic_use_dst_for_bwd, eltwise_logistic, - eltwise_exp_use_dst_for_bwd, eltwise_exp, eltwise_gelu_tanh, - eltwise_swish, eltwise_log, eltwise_clip, eltwise_gelu_erf); + eltwise_linear, eltwise_soft_relu, eltwise_logistic_use_dst_for_bwd, + eltwise_logistic, eltwise_exp_use_dst_for_bwd, eltwise_exp, + eltwise_gelu_tanh, eltwise_swish, eltwise_log, eltwise_clip, + eltwise_gelu_erf); return ok ? status::success : status::unimplemented; } diff --git a/src/cpu/primitive_attr_postops.cpp b/src/cpu/primitive_attr_postops.cpp index e8752e51c1f..bd3e109d18a 100644 --- a/src/cpu/primitive_attr_postops.cpp +++ b/src/cpu/primitive_attr_postops.cpp @@ -55,7 +55,6 @@ float compute_eltwise_scalar_fwd( case eltwise_abs: d = abs_fwd(s); break; case eltwise_sqrt: d = sqrt_fwd(s); break; case eltwise_linear: d = linear_fwd(s, alpha, beta); break; - case eltwise_bounded_relu: d = bounded_relu_fwd(s, alpha); break; case eltwise_soft_relu: d = soft_relu_fwd(s, alpha); break; case eltwise_logistic: d = logistic_fwd(s); break; case eltwise_exp: d = exp_fwd(s); break; @@ -96,7 +95,6 @@ float compute_eltwise_scalar_bwd( case eltwise_abs: ds = abs_bwd(dd, s); break; case eltwise_sqrt: ds = sqrt_bwd(dd, s); break; case eltwise_linear: ds = linear_bwd(dd, s, alpha, beta); break; - case eltwise_bounded_relu: ds = bounded_relu_bwd(dd, s, alpha); break; case eltwise_soft_relu: ds = soft_relu_bwd(dd, s, alpha); break; case eltwise_logistic: ds = logistic_bwd(dd, s); break; case eltwise_exp: ds = exp_bwd(dd, s); break; @@ -154,14 +152,14 @@ ref_eltwise_scalar_fwd_t::ref_eltwise_scalar_fwd_t( : alg_(alg), alpha_(alpha), beta_(beta), scale_(scale) { assert(utils::one_of(alg_, eltwise_relu, eltwise_tanh, eltwise_elu, eltwise_square, eltwise_abs, eltwise_sqrt, eltwise_linear, - eltwise_bounded_relu, eltwise_soft_relu, eltwise_mish, - eltwise_logistic, eltwise_exp, eltwise_gelu_tanh, eltwise_swish, - eltwise_log, eltwise_clip, eltwise_clip_v2, eltwise_pow, - eltwise_gelu_erf, eltwise_round, eltwise_hardsigmoid, - eltwise_hardswish, eltwise_relu_use_dst_for_bwd, - eltwise_tanh_use_dst_for_bwd, eltwise_elu_use_dst_for_bwd, - eltwise_sqrt_use_dst_for_bwd, eltwise_logistic_use_dst_for_bwd, - eltwise_exp_use_dst_for_bwd, eltwise_clip_v2_use_dst_for_bwd)); + eltwise_soft_relu, eltwise_mish, eltwise_logistic, eltwise_exp, + eltwise_gelu_tanh, eltwise_swish, eltwise_log, eltwise_clip, + eltwise_clip_v2, eltwise_pow, eltwise_gelu_erf, eltwise_round, + eltwise_hardsigmoid, eltwise_hardswish, + eltwise_relu_use_dst_for_bwd, eltwise_tanh_use_dst_for_bwd, + eltwise_elu_use_dst_for_bwd, eltwise_sqrt_use_dst_for_bwd, + eltwise_logistic_use_dst_for_bwd, eltwise_exp_use_dst_for_bwd, + eltwise_clip_v2_use_dst_for_bwd)); } ref_eltwise_scalar_fwd_t::ref_eltwise_scalar_fwd_t( diff --git a/src/cpu/x64/injectors/jit_uni_eltwise_injector.cpp b/src/cpu/x64/injectors/jit_uni_eltwise_injector.cpp index 9d40ef54b42..8e6305e0922 100644 --- a/src/cpu/x64/injectors/jit_uni_eltwise_injector.cpp +++ b/src/cpu/x64/injectors/jit_uni_eltwise_injector.cpp @@ -36,10 +36,10 @@ bool is_alg_supported(alg_kind_t alg) { using namespace alg_kind; return utils::one_of(alg, eltwise_relu, eltwise_tanh, eltwise_elu, eltwise_square, eltwise_abs, eltwise_sqrt, eltwise_linear, - eltwise_bounded_relu, eltwise_soft_relu, eltwise_logistic, - eltwise_mish, eltwise_exp, eltwise_gelu_tanh, eltwise_hardsigmoid, - eltwise_hardswish, eltwise_swish, eltwise_log, eltwise_clip, - eltwise_clip_v2, eltwise_pow, eltwise_gelu_erf, eltwise_round, + eltwise_soft_relu, eltwise_logistic, eltwise_mish, eltwise_exp, + eltwise_gelu_tanh, eltwise_hardsigmoid, eltwise_hardswish, + eltwise_swish, eltwise_log, eltwise_clip, eltwise_clip_v2, + eltwise_pow, eltwise_gelu_erf, eltwise_round, eltwise_relu_use_dst_for_bwd, eltwise_tanh_use_dst_for_bwd, eltwise_elu_use_dst_for_bwd, eltwise_sqrt_use_dst_for_bwd, eltwise_logistic_use_dst_for_bwd, eltwise_exp_use_dst_for_bwd, @@ -601,13 +601,6 @@ void jit_uni_eltwise_injector_f32::linear_compute_vector_fwd( h->uni_vfmadd213ps(vmm_src, vmm_aux0, table_val(beta)); } -template -void jit_uni_eltwise_injector_f32::bounded_relu_compute_vector_fwd( - const Vmm &vmm_src) { - h->uni_vmaxps(vmm_src, vmm_src, table_val(zero)); - h->uni_vminps(vmm_src, vmm_src, table_val(alpha)); -} - template void jit_uni_eltwise_injector_f32::clip_compute_vector_fwd( const Vmm &vmm_src) { @@ -1329,19 +1322,6 @@ void jit_uni_eltwise_injector_f32::linear_compute_vector_bwd( h->uni_vmovups(vmm_src, table_val(alpha)); } -template -void jit_uni_eltwise_injector_f32::bounded_relu_compute_vector_bwd( - const Vmm &vmm_src) { - // get mask of values > alpha and blend with 0.f - compute_cmp_mask(vmm_src, table_val(alpha), _cmp_gt_os); - blend_with_mask(vmm_src, table_val(zero)); - // make all negative values zeros - h->uni_vmaxps(vmm_src, vmm_src, table_val(zero)); - // everything bigger than 0.f should be 1.f - compute_cmp_mask(vmm_src, table_val(zero), _cmp_gt_os); - blend_with_mask(vmm_src, table_val(one)); -} - template void jit_uni_eltwise_injector_f32::soft_relu_compute_vector_bwd( const Vmm &vmm_src) { @@ -1618,7 +1598,6 @@ size_t jit_uni_eltwise_injector_f32::aux_vecs_count() { case eltwise_sqrt_use_dst_for_bwd: case eltwise_sqrt: return 0; case eltwise_linear: return 1; - case eltwise_bounded_relu: return 0; case eltwise_soft_relu: return 4; case eltwise_mish: return 4; case eltwise_logistic_use_dst_for_bwd: @@ -1651,7 +1630,6 @@ size_t jit_uni_eltwise_injector_f32::aux_vecs_count() { case eltwise_sqrt_use_dst_for_bwd: case eltwise_sqrt: return 1; case eltwise_linear: return 0; - case eltwise_bounded_relu: return 1; case eltwise_soft_relu: return 4; case eltwise_mish: return 4; case eltwise_logistic_use_dst_for_bwd: return 1; @@ -1700,9 +1678,6 @@ void jit_uni_eltwise_injector_f32::compute_body( case eltwise_sqrt: sqrt_compute_vector_fwd(Vmm(idx)); break; case eltwise_swish: swish_compute_vector_fwd(Vmm(idx)); break; case eltwise_linear: linear_compute_vector_fwd(Vmm(idx)); break; - case eltwise_bounded_relu: - bounded_relu_compute_vector_fwd(Vmm(idx)); - break; case eltwise_soft_relu: soft_relu_compute_vector_fwd(Vmm(idx)); break; @@ -1746,9 +1721,6 @@ void jit_uni_eltwise_injector_f32::compute_body( case eltwise_sqrt_use_dst_for_bwd: case eltwise_sqrt: sqrt_compute_vector_bwd(Vmm(idx)); break; case eltwise_linear: linear_compute_vector_bwd(Vmm(idx)); break; - case eltwise_bounded_relu: - bounded_relu_compute_vector_bwd(Vmm(idx)); - break; case eltwise_soft_relu: soft_relu_compute_vector_bwd(Vmm(idx)); break; diff --git a/src/cpu/x64/injectors/jit_uni_eltwise_injector.hpp b/src/cpu/x64/injectors/jit_uni_eltwise_injector.hpp index b93d932a758..314b8238a61 100644 --- a/src/cpu/x64/injectors/jit_uni_eltwise_injector.hpp +++ b/src/cpu/x64/injectors/jit_uni_eltwise_injector.hpp @@ -208,7 +208,6 @@ struct jit_uni_eltwise_injector_f32 { void abs_compute_vector_fwd(const Vmm &vmm_src); void sqrt_compute_vector_fwd(const Vmm &vmm_src); void linear_compute_vector_fwd(const Vmm &vmm_src); - void bounded_relu_compute_vector_fwd(const Vmm &vmm_src); void soft_relu_compute_vector_fwd(const Vmm &vmm_src); void mish_compute_vector_fwd(const Vmm &vmm_src); void logistic_compute_vector_fwd(const Vmm &vmm_src); @@ -230,7 +229,6 @@ struct jit_uni_eltwise_injector_f32 { void abs_compute_vector_bwd(const Vmm &vmm_src); void sqrt_compute_vector_bwd(const Vmm &vmm_src); void linear_compute_vector_bwd(const Vmm &vmm_src); - void bounded_relu_compute_vector_bwd(const Vmm &vmm_src); void soft_relu_compute_vector_bwd(const Vmm &vmm_src); void logistic_compute_vector_bwd(const Vmm &vmm_src); void mish_compute_vector_bwd(const Vmm &vmm_src); diff --git a/src/gpu/amd/miopen_eltwise.hpp b/src/gpu/amd/miopen_eltwise.hpp index f7876a2e4f4..71336232d77 100644 --- a/src/gpu/amd/miopen_eltwise.hpp +++ b/src/gpu/amd/miopen_eltwise.hpp @@ -44,8 +44,8 @@ struct miopen_eltwise_fwd_t : public primitive_t { prop_kind::forward_inference) // Supported algorithms && utils::one_of(desc()->alg_kind, eltwise_relu, - eltwise_bounded_relu, eltwise_tanh, eltwise_elu, - eltwise_soft_relu, eltwise_abs, eltwise_logistic) + eltwise_tanh, eltwise_elu, eltwise_soft_relu, + eltwise_abs, eltwise_logistic) && IMPLICATION(desc()->alg_kind == eltwise_soft_relu, desc()->alpha == 1.f) // Supported data types @@ -81,8 +81,8 @@ struct miopen_eltwise_bwd_t : public primitive_t { bool ok = true && desc()->prop_kind == prop_kind::backward_data // Supported algorithms - && utils::one_of(desc()->alg_kind, eltwise_relu, - eltwise_bounded_relu, eltwise_soft_relu) + && utils::one_of( + desc()->alg_kind, eltwise_relu, eltwise_soft_relu) && IMPLICATION(desc()->alg_kind == eltwise_soft_relu, desc()->alpha == 1.f) // Supported data types diff --git a/src/gpu/amd/miopen_eltwise_impl.hpp b/src/gpu/amd/miopen_eltwise_impl.hpp index 5e249686a75..f6190b6da2d 100644 --- a/src/gpu/amd/miopen_eltwise_impl.hpp +++ b/src/gpu/amd/miopen_eltwise_impl.hpp @@ -52,10 +52,6 @@ struct miopen_eltwise_impl_base_t { *miopen_alg_kind = miopenActivationMode_t::miopenActivationLEAKYRELU; break; - case alg_kind::eltwise_bounded_relu: - *miopen_alg_kind - = miopenActivationMode_t::miopenActivationCLIPPEDRELU; - break; case alg_kind::eltwise_tanh: *miopen_alg_kind = miopenActivationMode_t::miopenActivationTANH; break; diff --git a/src/gpu/jit/ir/eltwise.hpp b/src/gpu/jit/ir/eltwise.hpp index 5cd1d09c3f7..79d448e300b 100644 --- a/src/gpu/jit/ir/eltwise.hpp +++ b/src/gpu/jit/ir/eltwise.hpp @@ -59,7 +59,6 @@ class eltwise_t : public func_impl_t { case alg_kind::eltwise_sqrt: return "sqrt"; case alg_kind::eltwise_swish: return "swish"; case alg_kind::eltwise_linear: return "linear"; - case alg_kind::eltwise_bounded_relu: return "bounded_relu"; case alg_kind::eltwise_soft_relu: return "soft_relu"; case alg_kind::eltwise_logistic: return "logistic"; case alg_kind::eltwise_mish: return "mish"; diff --git a/src/gpu/jit/jit_eltwise_injector.cpp b/src/gpu/jit/jit_eltwise_injector.cpp index c236186b501..ddca503da1a 100644 --- a/src/gpu/jit/jit_eltwise_injector.cpp +++ b/src/gpu/jit/jit_eltwise_injector.cpp @@ -53,7 +53,6 @@ int jit_eltwise_injector_f32::min_scratch_regs() { case eltwise_tanh_use_dst_for_bwd: return 2; case eltwise_round: return 0; case eltwise_linear: return 0; - case eltwise_bounded_relu: case eltwise_clip: case eltwise_clip_v2: case eltwise_clip_v2_use_dst_for_bwd: return 0; @@ -68,7 +67,6 @@ int jit_eltwise_injector_f32::min_scratch_regs() { case eltwise_abs: return 1; case eltwise_square: return 0; case eltwise_linear: return 0; - case eltwise_bounded_relu: case eltwise_clip: return 1; case eltwise_gelu_tanh: return 2; default: assert(!"unsupported eltwise algorithm"); @@ -165,7 +163,6 @@ int jit_eltwise_injector_f32::phase_count(alg_kind_t alg) { case eltwise_tanh_use_dst_for_bwd: return (use_tanh_compat()) ? 9 : 6; case eltwise_linear: return 2; - case eltwise_bounded_relu: case eltwise_clip: case eltwise_clip_v2: case eltwise_clip_v2_use_dst_for_bwd: return 2; @@ -177,7 +174,6 @@ int jit_eltwise_injector_f32::phase_count(alg_kind_t alg) { } else { switch (alg) { case eltwise_abs: return 2; - case eltwise_bounded_relu: case eltwise_clip: return 4; case eltwise_gelu_tanh: return 14; default: break; @@ -704,9 +700,6 @@ void jit_eltwise_injector_f32::compute(const ngen::GRFRange ®s) { else relu_compute_fwd(simd, base, phase, ii); break; - case eltwise_bounded_relu: - clip_compute_fwd(simd, base, phase, 0, alpha_); - break; case eltwise_abs: abs_compute_fwd(simd, base); break; case eltwise_soft_relu: soft_relu_compute_fwd(simd, base, phase, ii); @@ -752,9 +745,6 @@ void jit_eltwise_injector_f32::compute(const ngen::GRFRange ®s) { } else { switch (alg_) { case eltwise_relu: relu_compute_bwd(simd, base); break; - case eltwise_bounded_relu: - clip_compute_bwd(simd, base, phase, 0, alpha_); - break; case eltwise_abs: abs_compute_bwd(simd, base, phase); break; @@ -803,7 +793,6 @@ void jit_eltwise_injector_f32::prepare() { switch (alg_) { case eltwise_relu: relu_prepare_bwd(); break; case eltwise_abs: abs_prepare_bwd(); break; - case eltwise_bounded_relu: case eltwise_clip: clip_prepare_bwd(); break; default: break; } diff --git a/src/gpu/jit/jit_eltwise_injector.hpp b/src/gpu/jit/jit_eltwise_injector.hpp index 0a96bfff580..2f79e527aa0 100644 --- a/src/gpu/jit/jit_eltwise_injector.hpp +++ b/src/gpu/jit/jit_eltwise_injector.hpp @@ -34,11 +34,10 @@ inline bool jit_eltwise_injector_f32_is_supported(alg_kind_t alg) { eltwise_exp, eltwise_exp_use_dst_for_bwd, eltwise_gelu_tanh, eltwise_gelu_erf, eltwise_hardsigmoid, eltwise_hardswish, eltwise_log, eltwise_mish, eltwise_pow, eltwise_relu, - eltwise_relu_use_dst_for_bwd, eltwise_bounded_relu, - eltwise_soft_relu, eltwise_sqrt, eltwise_sqrt_use_dst_for_bwd, - eltwise_square, eltwise_swish, eltwise_tanh, - eltwise_tanh_use_dst_for_bwd, eltwise_abs, eltwise_round, - eltwise_linear, eltwise_clip, eltwise_clip_v2, + eltwise_relu_use_dst_for_bwd, eltwise_soft_relu, eltwise_sqrt, + eltwise_sqrt_use_dst_for_bwd, eltwise_square, eltwise_swish, + eltwise_tanh, eltwise_tanh_use_dst_for_bwd, eltwise_abs, + eltwise_round, eltwise_linear, eltwise_clip, eltwise_clip_v2, eltwise_clip_v2_use_dst_for_bwd, eltwise_logistic, eltwise_logistic_use_dst_for_bwd); } diff --git a/src/gpu/nvidia/README.md b/src/gpu/nvidia/README.md index 9747634d333..e2f250dd66a 100644 --- a/src/gpu/nvidia/README.md +++ b/src/gpu/nvidia/README.md @@ -174,10 +174,10 @@ eltwise forward and eltwise backward in oneDNN respectively. There are some limitations when using Nvidia backend for eltwise primitive: * cuDNN only supports the following operations - `RELU`, `ELU`, `TANH`, - `LOGISTIC` and `BRELU`. + `LOGISTIC`. * `RELU` is only supported with alpha = 0. * cuDNN expects `x`, `y` and `dy` as inputs to the backward pass, hence, only - `RELU` and `BRELU` operations are supported in the backward pass. + `RELU` operation supports backward proragation kind. TODO: add `ELU_DST`, `TANH_DST` and `LOGISTIC_DST` support which require `dy`. * Forward pass supports `f32`, `f16` and `s8` data types. Although blocking is not supported for `s8`. diff --git a/src/gpu/nvidia/cudnn_conv_inner_product_impl.hpp b/src/gpu/nvidia/cudnn_conv_inner_product_impl.hpp index f43120af72b..15c06759f44 100644 --- a/src/gpu/nvidia/cudnn_conv_inner_product_impl.hpp +++ b/src/gpu/nvidia/cudnn_conv_inner_product_impl.hpp @@ -401,9 +401,6 @@ struct cudnn_conv_inner_product_fwd_impl_t case alg_kind::eltwise_logistic: no_relu_mode = CUDNN_ACTIVATION_SIGMOID; break; - case alg_kind::eltwise_bounded_relu: - no_relu_mode = CUDNN_ACTIVATION_CLIPPED_RELU; - break; default: return status::unimplemented; } CHECK(CUDNN_EXECUTE_FUNC_S(cudnnSetActivationDescriptor, diff --git a/src/gpu/nvidia/cudnn_convolution_impl.hpp b/src/gpu/nvidia/cudnn_convolution_impl.hpp index 7e36a4191cf..c8d99332784 100644 --- a/src/gpu/nvidia/cudnn_convolution_impl.hpp +++ b/src/gpu/nvidia/cudnn_convolution_impl.hpp @@ -696,9 +696,6 @@ struct cudnn_convolution_impl_fwd_t : public cudnn_convolution_impl_base_t { case alg_kind::eltwise_logistic: act_mode = CUDNN_ACTIVATION_SIGMOID; break; - case alg_kind::eltwise_bounded_relu: - act_mode = CUDNN_ACTIVATION_CLIPPED_RELU; - break; default: return status::unimplemented; } CHECK(CUDNN_EXECUTE_FUNC_S(cudnnSetActivationDescriptor, eltwise_desc, diff --git a/src/gpu/nvidia/cudnn_eltwise.hpp b/src/gpu/nvidia/cudnn_eltwise.hpp index 7d25f817a12..a8f10e12c80 100644 --- a/src/gpu/nvidia/cudnn_eltwise.hpp +++ b/src/gpu/nvidia/cudnn_eltwise.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2020-2021 Intel Corporation +* Copyright 2020-2022 Intel Corporation * Copyright 2020 Codeplay Software Limited * * Licensed under the Apache License, Version 2.0 (the "License"); @@ -44,8 +44,7 @@ struct cudnn_eltwise_fwd_t : public primitive_t { prop_kind::forward_inference) // Supported algorithms && utils::one_of(desc()->alg_kind, eltwise_relu, - eltwise_bounded_relu, eltwise_tanh, eltwise_elu, - eltwise_logistic) + eltwise_tanh, eltwise_elu, eltwise_logistic) // Supported data types && utils::one_of(desc()->data_desc.data_type, data_type::f32, data_type::f16, data_type::s8) @@ -81,8 +80,7 @@ struct cudnn_eltwise_bwd_t : public primitive_t { bool ok = true && desc()->prop_kind == prop_kind::backward_data // Supported algorithms - && utils::one_of(desc()->alg_kind, eltwise_bounded_relu, - eltwise_relu) + && utils::one_of(desc()->alg_kind, eltwise_relu) // Supported data types && desc()->data_desc.data_type == data_type::f32 && IMPLICATION(desc()->alg_kind == eltwise_relu, diff --git a/src/gpu/nvidia/cudnn_eltwise_impl.hpp b/src/gpu/nvidia/cudnn_eltwise_impl.hpp index 24c3b0331ff..aa20406bdc4 100644 --- a/src/gpu/nvidia/cudnn_eltwise_impl.hpp +++ b/src/gpu/nvidia/cudnn_eltwise_impl.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2020 Intel Corporation +* Copyright 2020-2022 Intel Corporation * Copyright 2020 Codeplay Software Limited * * Licensed under the Apache License, Version 2.0 (the "License"); @@ -51,10 +51,6 @@ struct cudnn_eltwise_impl_base_t { case alg_kind::eltwise_relu: *cuda_alg_kind = cudnnActivationMode_t::CUDNN_ACTIVATION_RELU; break; - case alg_kind::eltwise_bounded_relu: - *cuda_alg_kind - = cudnnActivationMode_t::CUDNN_ACTIVATION_CLIPPED_RELU; - break; case alg_kind::eltwise_tanh: *cuda_alg_kind = cudnnActivationMode_t::CUDNN_ACTIVATION_TANH; break; @@ -83,8 +79,7 @@ struct cudnn_eltwise_impl_base_t { // alpha and beta are post operation scaling parameters used by cuDNN float alpha = 1; float beta = 0; - // coef in cuDNN is use for Relu (is equal to zero) and BRelu (represents - // the bound) + // coef in cuDNN is use for Relu (is equal to zero) double coef = 0; }; diff --git a/src/gpu/nvidia/cudnn_gemm_inner_product_impl.hpp b/src/gpu/nvidia/cudnn_gemm_inner_product_impl.hpp index 1201191f11b..2e0abff6d6c 100644 --- a/src/gpu/nvidia/cudnn_gemm_inner_product_impl.hpp +++ b/src/gpu/nvidia/cudnn_gemm_inner_product_impl.hpp @@ -245,9 +245,6 @@ struct cudnn_gemm_inner_product_fwd_impl_t case alg_kind::eltwise_logistic: act_mode = CUDNN_ACTIVATION_SIGMOID; break; - case alg_kind::eltwise_bounded_relu: - act_mode = CUDNN_ACTIVATION_CLIPPED_RELU; - break; default: return status::unimplemented; } CHECK(CUDNN_EXECUTE_FUNC_S(cudnnSetActivationDescriptor, act_desc_, diff --git a/src/gpu/nvidia/cudnn_matmul.hpp b/src/gpu/nvidia/cudnn_matmul.hpp index 76f7dd5cc8a..9d30984a549 100644 --- a/src/gpu/nvidia/cudnn_matmul.hpp +++ b/src/gpu/nvidia/cudnn_matmul.hpp @@ -97,8 +97,8 @@ struct cudnn_matmul_t : public primitive_t { if (eltwise_idx != -1) { using namespace alg_kind; const bool ok = utils::one_of(p.entry_[eltwise_idx].eltwise.alg, - eltwise_relu, eltwise_bounded_relu, eltwise_tanh, - eltwise_elu, eltwise_logistic); + eltwise_relu, eltwise_tanh, eltwise_elu, + eltwise_logistic); if (!ok) return false; } diff --git a/src/gpu/nvidia/cudnn_matmul_impl.hpp b/src/gpu/nvidia/cudnn_matmul_impl.hpp index c7694f5587f..0934fe8fcf4 100644 --- a/src/gpu/nvidia/cudnn_matmul_impl.hpp +++ b/src/gpu/nvidia/cudnn_matmul_impl.hpp @@ -77,9 +77,6 @@ struct cudnn_matmul_impl_t { case alg_kind::eltwise_relu: mode = cudnnActivationMode_t::CUDNN_ACTIVATION_RELU; break; - case alg_kind::eltwise_bounded_relu: - mode = cudnnActivationMode_t::CUDNN_ACTIVATION_CLIPPED_RELU; - break; case alg_kind::eltwise_tanh: mode = cudnnActivationMode_t::CUDNN_ACTIVATION_TANH; break; diff --git a/src/gpu/ocl/gen9_eltwise.hpp b/src/gpu/ocl/gen9_eltwise.hpp index 47f5a84c4d4..f15d67aab59 100644 --- a/src/gpu/ocl/gen9_eltwise.hpp +++ b/src/gpu/ocl/gen9_eltwise.hpp @@ -49,14 +49,14 @@ struct gen9_eltwise_fwd_t : public gpu_primitive_t { prop_kind::forward_training, prop_kind::forward_inference) && utils::one_of(desc()->alg_kind, eltwise_relu, - eltwise_linear, eltwise_bounded_relu, eltwise_abs, - eltwise_tanh, eltwise_elu, eltwise_square, - eltwise_sqrt, eltwise_soft_relu, eltwise_logistic, - eltwise_mish, eltwise_exp, eltwise_gelu_tanh, - eltwise_hardswish, eltwise_swish, eltwise_log, - eltwise_clip, eltwise_clip_v2, eltwise_pow, - eltwise_gelu_erf, eltwise_round, - eltwise_hardsigmoid, eltwise_relu_use_dst_for_bwd, + eltwise_linear, eltwise_abs, eltwise_tanh, + eltwise_elu, eltwise_square, eltwise_sqrt, + eltwise_soft_relu, eltwise_logistic, eltwise_mish, + eltwise_exp, eltwise_gelu_tanh, eltwise_hardswish, + eltwise_swish, eltwise_log, eltwise_clip, + eltwise_clip_v2, eltwise_pow, eltwise_gelu_erf, + eltwise_round, eltwise_hardsigmoid, + eltwise_relu_use_dst_for_bwd, eltwise_logistic_use_dst_for_bwd, eltwise_tanh_use_dst_for_bwd, eltwise_elu_use_dst_for_bwd, @@ -123,14 +123,13 @@ struct gen9_eltwise_bwd_t : public gpu_primitive_t { using namespace alg_kind; bool ok = desc()->prop_kind == backward_data && utils::one_of(desc()->alg_kind, eltwise_relu, - eltwise_linear, eltwise_bounded_relu, eltwise_abs, - eltwise_tanh, eltwise_elu, eltwise_square, - eltwise_sqrt, eltwise_soft_relu, eltwise_mish, - eltwise_logistic, eltwise_exp, eltwise_hardswish, - eltwise_gelu_tanh, eltwise_swish, eltwise_log, - eltwise_clip, eltwise_clip_v2, eltwise_pow, - eltwise_gelu_erf, eltwise_hardsigmoid, - eltwise_relu_use_dst_for_bwd, + eltwise_linear, eltwise_abs, eltwise_tanh, + eltwise_elu, eltwise_square, eltwise_sqrt, + eltwise_soft_relu, eltwise_mish, eltwise_logistic, + eltwise_exp, eltwise_hardswish, eltwise_gelu_tanh, + eltwise_swish, eltwise_log, eltwise_clip, + eltwise_clip_v2, eltwise_pow, eltwise_gelu_erf, + eltwise_hardsigmoid, eltwise_relu_use_dst_for_bwd, eltwise_logistic_use_dst_for_bwd, eltwise_tanh_use_dst_for_bwd, eltwise_elu_use_dst_for_bwd, diff --git a/src/gpu/ocl/ocl_eltwise.h b/src/gpu/ocl/ocl_eltwise.h index 9aa39250a33..d213ff5004f 100644 --- a/src/gpu/ocl/ocl_eltwise.h +++ b/src/gpu/ocl/ocl_eltwise.h @@ -52,14 +52,6 @@ float linear_bwd(float dd, float alpha) { return dd * alpha; } -float bounded_relu_fwd(float s, float alpha) { - s = s > 0 ? s : 0; - return s > alpha ? alpha : s; -} -float bounded_relu_bwd(float dd, float s, float alpha) { - return dd * (0 < s && s <= alpha ? 1 : 0); -} - float soft_relu_fwd(float s, float alpha) { s = alpha * s; float v = (s < log((float)DATA_MAX) ? log1p(exp(s)) : s); @@ -233,7 +225,7 @@ float hardsigmoid_bwd(float dd, float s, float alpha, float beta) { } float hardswish_fwd(float s) { - return (s / 6.f) * bounded_relu_fwd(s + 3.f, 6.f); + return (s / 6.f) * clip_fwd(s + 3.f, 0.f, 6.f); } float hardswish_bwd(float dd, float s) { return (s < 3.f && s > -3.f ? dd * (2 * s + 3.f) / 6.f @@ -245,7 +237,6 @@ float fwd_eltwise_common( switch (eltwise_alg) { case RELU: return scale_ * relu_fwd(x, alpha_); break; case LINEAR: return scale_ * linear_fwd(x, alpha_, beta_); break; - case BOUNDED_RELU: return scale_ * bounded_relu_fwd(x, alpha_); break; case SOFT_RELU: return scale_ * soft_relu_fwd(x, alpha_); break; case MISH: return scale_ * mish_fwd(x); break; case LOGISTIC: return scale_ * logistic_fwd(x); break; @@ -292,7 +283,6 @@ float bwd_eltwise(float x, float y, float alpha_, float beta_) { switch (ELTWISE_ALG) { case RELU: return relu_bwd(x, y, alpha_); break; case LINEAR: return linear_bwd(x, alpha_); break; - case BOUNDED_RELU: return bounded_relu_bwd(x, y, alpha_); break; case SOFT_RELU: return soft_relu_bwd(x, y, alpha_); break; case MISH: return mish_bwd(x, y); break; case LOGISTIC: return logistic_bwd(x, y); break; diff --git a/src/gpu/ocl/ref_eltwise.hpp b/src/gpu/ocl/ref_eltwise.hpp index 09ea1a4ee98..97457047fd6 100644 --- a/src/gpu/ocl/ref_eltwise.hpp +++ b/src/gpu/ocl/ref_eltwise.hpp @@ -51,14 +51,14 @@ struct ref_eltwise_fwd_t : public gpu_primitive_t { prop_kind::forward_training, prop_kind::forward_inference) && utils::one_of(desc()->alg_kind, eltwise_relu, - eltwise_linear, eltwise_bounded_relu, eltwise_abs, - eltwise_tanh, eltwise_elu, eltwise_square, - eltwise_sqrt, eltwise_soft_relu, eltwise_logistic, - eltwise_mish, eltwise_exp, eltwise_gelu_tanh, - eltwise_hardswish, eltwise_swish, eltwise_log, - eltwise_clip, eltwise_clip_v2, eltwise_pow, - eltwise_gelu_erf, eltwise_round, - eltwise_hardsigmoid, eltwise_relu_use_dst_for_bwd, + eltwise_linear, eltwise_abs, eltwise_tanh, + eltwise_elu, eltwise_square, eltwise_sqrt, + eltwise_soft_relu, eltwise_logistic, eltwise_mish, + eltwise_exp, eltwise_gelu_tanh, eltwise_hardswish, + eltwise_swish, eltwise_log, eltwise_clip, + eltwise_clip_v2, eltwise_pow, eltwise_gelu_erf, + eltwise_round, eltwise_hardsigmoid, + eltwise_relu_use_dst_for_bwd, eltwise_logistic_use_dst_for_bwd, eltwise_tanh_use_dst_for_bwd, eltwise_elu_use_dst_for_bwd, @@ -134,14 +134,13 @@ struct ref_eltwise_bwd_t : public gpu_primitive_t { using namespace alg_kind; const bool ok = desc()->prop_kind == backward_data && utils::one_of(desc()->alg_kind, eltwise_relu, - eltwise_linear, eltwise_bounded_relu, eltwise_abs, - eltwise_tanh, eltwise_elu, eltwise_square, - eltwise_sqrt, eltwise_soft_relu, eltwise_mish, - eltwise_logistic, eltwise_exp, eltwise_gelu_tanh, - eltwise_hardswish, eltwise_swish, eltwise_log, - eltwise_clip, eltwise_clip_v2, eltwise_pow, - eltwise_gelu_erf, eltwise_hardsigmoid, - eltwise_relu_use_dst_for_bwd, + eltwise_linear, eltwise_abs, eltwise_tanh, + eltwise_elu, eltwise_square, eltwise_sqrt, + eltwise_soft_relu, eltwise_mish, eltwise_logistic, + eltwise_exp, eltwise_gelu_tanh, eltwise_hardswish, + eltwise_swish, eltwise_log, eltwise_clip, + eltwise_clip_v2, eltwise_pow, eltwise_gelu_erf, + eltwise_hardsigmoid, eltwise_relu_use_dst_for_bwd, eltwise_logistic_use_dst_for_bwd, eltwise_tanh_use_dst_for_bwd, eltwise_elu_use_dst_for_bwd, diff --git a/src/gpu/primitive_conf.hpp b/src/gpu/primitive_conf.hpp index af52fe3ce55..229bfa94730 100644 --- a/src/gpu/primitive_conf.hpp +++ b/src/gpu/primitive_conf.hpp @@ -1036,7 +1036,6 @@ inline void def_binary_alg_kinds(compute::kernel_ctx_t &kernel_ctx) { inline void def_eltwise_alg_kinds(compute::kernel_ctx_t &kernel_ctx) { kernel_ctx.define_int("RELU", alg_kind::eltwise_relu); kernel_ctx.define_int("LINEAR", alg_kind::eltwise_linear); - kernel_ctx.define_int("BOUNDED_RELU", alg_kind::eltwise_bounded_relu); kernel_ctx.define_int("SOFT_RELU", alg_kind::eltwise_soft_relu); kernel_ctx.define_int("MISH", alg_kind::eltwise_mish); kernel_ctx.define_int("LOGISTIC", alg_kind::eltwise_logistic); diff --git a/tests/benchdnn/conv/conv.cpp b/tests/benchdnn/conv/conv.cpp index 1a7688916ad..4b0d2c6502b 100644 --- a/tests/benchdnn/conv/conv.cpp +++ b/tests/benchdnn/conv/conv.cpp @@ -45,7 +45,7 @@ double get_non_zero_trust_percent(const prb_t *prb, data_kind_t kind) { int count = 0; // Check for all post-ops that convert negative to zero - std::vector non_neg_po {pk::ABS, pk::BRELU}; + std::vector non_neg_po {pk::ABS}; std::vector non_neg_alpha_0_po { pk::CLIP, pk::CLIP_V2, pk::ELU, pk::RELU}; for (int i = 0; i < po.len(); ++i) { diff --git a/tests/benchdnn/deconv/deconv.cpp b/tests/benchdnn/deconv/deconv.cpp index 6bc82fa0bbf..b67ebf6538c 100644 --- a/tests/benchdnn/deconv/deconv.cpp +++ b/tests/benchdnn/deconv/deconv.cpp @@ -57,7 +57,7 @@ double get_non_zero_trust_percent(const prb_t *prb, data_kind_t kind) { int count = 0; // Check for all post-ops that convert negative to zero - std::vector non_neg_po {pk::ABS, pk::BRELU}; + std::vector non_neg_po {pk::ABS}; std::vector non_neg_alpha_0_po { pk::CLIP, pk::CLIP_V2, pk::ELU, pk::RELU}; for (int i = 0; i < po.len(); ++i) { diff --git a/tests/benchdnn/dnn_types.cpp b/tests/benchdnn/dnn_types.cpp index fa74f7a3754..1d422529107 100644 --- a/tests/benchdnn/dnn_types.cpp +++ b/tests/benchdnn/dnn_types.cpp @@ -329,8 +329,6 @@ static po_table_entry_t kind_table[] = { // eltwise {pk_t::ELTWISE_START, {"eltwise_undef"}, dnnl_alg_kind_undef}, {pk_t::ABS, {"abs", "eltwise_abs"}, dnnl_eltwise_abs}, - {pk_t::BRELU, {"bounded_relu", "eltwise_bounded_relu", "brelu"}, - dnnl_eltwise_bounded_relu}, {pk_t::CLIP, {"clip", "eltwise_clip"}, dnnl_eltwise_clip}, {pk_t::CLIP_V2, {"clip_v2", "eltwise_clip_v2"}, dnnl_eltwise_clip_v2}, {pk_t::CLIP_V2_DST, {"clip_v2_dst", "eltwise_clip_v2_use_dst_for_bwd"}, @@ -1323,7 +1321,6 @@ float compute_eltwise_fwd( case pk_t::ABS: return scale * abs_fwd(src); case pk_t::SQRT: return scale * sqrt_fwd(src); case pk_t::LINEAR: return scale * linear_fwd(src, alpha, beta); - case pk_t::BRELU: return scale * bounded_relu_fwd(src, alpha); case pk_t::SRELU: return scale * soft_relu_fwd(src, alpha); case pk_t::MISH: return scale * mish_fwd(src); case pk_t::LOGISTIC: return scale * logistic_fwd(src); @@ -1364,7 +1361,6 @@ float compute_eltwise_bwd( case pk_t::ABS: return abs_bwd(d_dst, src); case pk_t::SQRT: return sqrt_bwd(d_dst, src); case pk_t::LINEAR: return linear_bwd(d_dst, src, alpha, beta); - case pk_t::BRELU: return bounded_relu_bwd(d_dst, src, alpha); case pk_t::SRELU: return soft_relu_bwd(d_dst, src, alpha); case pk_t::MISH: return mish_bwd(d_dst, src); case pk_t::LOGISTIC: return logistic_bwd(d_dst, src); diff --git a/tests/benchdnn/dnn_types.hpp b/tests/benchdnn/dnn_types.hpp index 1bd03d19093..9cb5ca533a7 100644 --- a/tests/benchdnn/dnn_types.hpp +++ b/tests/benchdnn/dnn_types.hpp @@ -215,7 +215,6 @@ struct attr_t { // eltwise ELTWISE_START, // a guard to check kind is eltwise ABS, - BRELU, CLIP, CLIP_V2, CLIP_V2_DST, diff --git a/tests/benchdnn/doc/driver_eltwise.md b/tests/benchdnn/doc/driver_eltwise.md index 3def11f3456..a87b46fdaa0 100644 --- a/tests/benchdnn/doc/driver_eltwise.md +++ b/tests/benchdnn/doc/driver_eltwise.md @@ -40,9 +40,9 @@ the end to specify fewer dimensions. ## Floating point arguments -Some operations support `alpha` argument such as `BRELU`, `CLIP`, `CLIP_V2`, - `ELU`, `LINEAR`, `POW` and `RELU`. `CLIP`, `CLIP_V2`, `LINEAR` and `POW` also -support `beta` argument. +Some operations support `alpha` argument. Some also support `beta` argument. +Refer to a [full list](knobs_attr.md) for details. + The `alpha` and `beta` parameters should meet algorithm requirements, otherwise the problem will be silently skipped. For instance: diff --git a/tests/benchdnn/doc/knobs_attr.md b/tests/benchdnn/doc/knobs_attr.md index 06bb3510064..e76215c2b4b 100644 --- a/tests/benchdnn/doc/knobs_attr.md +++ b/tests/benchdnn/doc/knobs_attr.md @@ -168,7 +168,6 @@ Operations may be called in any order, e.g. apply `SUM` at first and then apply - `tanh` - `tanh_dst` - Eltwise operations that support only alpha: - - `bounded_relu` - `elu` - `elu_dst` - `relu` diff --git a/tests/benchdnn/eltwise/bench_eltwise.cpp b/tests/benchdnn/eltwise/bench_eltwise.cpp index d4b6e048044..2cf6ea1249d 100644 --- a/tests/benchdnn/eltwise/bench_eltwise.cpp +++ b/tests/benchdnn/eltwise/bench_eltwise.cpp @@ -75,7 +75,6 @@ void check_correctness(const settings_t &s) { "WARNING: non-zero beta is ignored. " "Consider adding --beta=0 to a command line."); break; - case alg_t::BRELU: case alg_t::ELU: case alg_t::ELU_DST: case alg_t::RELU: diff --git a/tests/benchdnn/eltwise/eltwise.cpp b/tests/benchdnn/eltwise/eltwise.cpp index bd8076527fb..f1311761994 100644 --- a/tests/benchdnn/eltwise/eltwise.cpp +++ b/tests/benchdnn/eltwise/eltwise.cpp @@ -163,9 +163,6 @@ static float get_eltwise_zero_trust_percent(const prb_t *prb) { case alg_t::LINEAR: if (prb->alpha == 0) ztp = 100.f; break; - case alg_t::BRELU: - if ((prb->alpha == 0) || (prb->dir & FLAG_BWD)) ztp = 100.f; - break; case alg_t::CLIP: case alg_t::CLIP_V2: case alg_t::CLIP_V2_DST: @@ -268,7 +265,6 @@ void skip_invalid_prb(const prb_t *prb, res_t *res) { case alg_t::CLIP: case alg_t::CLIP_V2: case alg_t::CLIP_V2_DST: is_invalid = prb->beta < prb->alpha; break; - case alg_t::BRELU: case alg_t::ELU_DST: case alg_t::RELU_DST: is_invalid = prb->alpha < 0; break; case alg_t::ROUND: diff --git a/tests/benchdnn/inputs/conv/option_set_all_eltwise_postops b/tests/benchdnn/inputs/conv/option_set_all_eltwise_postops index e771e27da95..b77be90f2b5 100644 --- a/tests/benchdnn/inputs/conv/option_set_all_eltwise_postops +++ b/tests/benchdnn/inputs/conv/option_set_all_eltwise_postops @@ -1,7 +1,7 @@ # Option set to test all eltwise postops. # alpha and beta values taken from ../eltwise/option_set_all_algs_ci ---attr-post-ops=abs,bounded_relu:2,clip:-2:3,clip_v2:-2:3,elu:-2,exp,gelu_erf, \ +--attr-post-ops=abs,clip:-2:3,clip_v2:-2:3,elu:-2,exp,gelu_erf, \ gelu_tanh,hardsigmoid:-2:3,hardswish,linear:-2:3,log,logistic, \ mish,pow:1:-1,pow:1:0.5,pow:1:1.5,relu,relu:-2, \ round,soft_relu:1,soft_relu:-1,soft_relu:-2,square, \ diff --git a/tests/benchdnn/inputs/conv/option_set_combined_postops b/tests/benchdnn/inputs/conv/option_set_combined_postops index 18a121920ab..538c698d5bf 100644 --- a/tests/benchdnn/inputs/conv/option_set_combined_postops +++ b/tests/benchdnn/inputs/conv/option_set_combined_postops @@ -6,7 +6,7 @@ --attr-post-ops=sum+elu:0.5 --batch=shapes_tails --attr-post-ops=sum+abs --batch=shapes_tails --attr-post-ops=sum+linear:0.5:1.5 --batch=shapes_tails ---attr-post-ops=sum+bounded_relu:0.5 --batch=shapes_tails +--attr-post-ops=sum+clip:0:0.5 --batch=shapes_tails --attr-post-ops=sum+logistic:0:0:2.5 --batch=shapes_tails --attr-post-ops=sum+square --batch=shapes_tails --attr-post-ops=sum+soft_relu:1 --batch=shapes_tails diff --git a/tests/benchdnn/inputs/eltwise/option_set_all_algs b/tests/benchdnn/inputs/eltwise/option_set_all_algs index 70ec48c5fbf..f1bb71a7581 100644 --- a/tests/benchdnn/inputs/eltwise/option_set_all_algs +++ b/tests/benchdnn/inputs/eltwise/option_set_all_algs @@ -6,7 +6,7 @@ --batch=shapes_eltwise --alpha=0,0.25,-0.25 --beta=0 ---alg=bounded_relu,elu,elu_dst,relu,relu_dst,swish +--alg=elu,elu_dst,relu,relu_dst,swish --batch=shapes_eltwise ### logsigmoid and soft_relu diff --git a/tests/benchdnn/inputs/eltwise/option_set_all_algs_ci b/tests/benchdnn/inputs/eltwise/option_set_all_algs_ci index 692d2c531e7..7a26f55683e 100644 --- a/tests/benchdnn/inputs/eltwise/option_set_all_algs_ci +++ b/tests/benchdnn/inputs/eltwise/option_set_all_algs_ci @@ -19,7 +19,7 @@ ## algs which do not support negative alpha --alpha=2 --beta=0 ---alg=bounded_relu,elu_dst,relu_dst +--alg=elu_dst,relu_dst --batch=shapes_ci ## algs which support alpha and beta diff --git a/tests/gtests/internals/test_comparison_operators.cpp b/tests/gtests/internals/test_comparison_operators.cpp index 8c29ce3f8d5..c4739a36d22 100644 --- a/tests/gtests/internals/test_comparison_operators.cpp +++ b/tests/gtests/internals/test_comparison_operators.cpp @@ -101,7 +101,7 @@ TEST(comparison_operators_t, TestEltwisePostOp) { dnnl::primitive_attr attr; dnnl::post_ops ops; - ops.append_eltwise(NAN, algorithm::eltwise_bounded_relu, 2.5f, 3.5f); + ops.append_eltwise(NAN, algorithm::eltwise_clip, 2.5f, 3.5f); attr.set_post_ops(ops); TEST_SELF_COMPARISON(attr); } diff --git a/tests/gtests/test_convolution_eltwise_forward_common.hpp b/tests/gtests/test_convolution_eltwise_forward_common.hpp index d0e1e57e173..28d79793f9e 100644 --- a/tests/gtests/test_convolution_eltwise_forward_common.hpp +++ b/tests/gtests/test_convolution_eltwise_forward_common.hpp @@ -91,8 +91,8 @@ void compute_ref_conv_eltwise_fwd(const test_convolution_sizes_t &c, case algorithm::eltwise_linear: d = linear_fwd(d, elt_alpha, elt_beta); break; - case algorithm::eltwise_bounded_relu: - d = bounded_relu_fwd(d, elt_alpha); + case algorithm::eltwise_clip: + d = clip_fwd(d, elt_alpha, elt_beta); break; case algorithm::eltwise_soft_relu: d = soft_relu_fwd(d, elt_alpha); @@ -145,7 +145,6 @@ class convolution_eltwise_test memory::format_tag::odhwi))), "Format is not supported."); SKIP_IF_CUDA(p.alg != algorithm::eltwise_relu - && p.alg != algorithm::eltwise_bounded_relu && p.alg != algorithm::eltwise_tanh && p.alg != algorithm::eltwise_elu && p.alg != algorithm::eltwise_logistic, diff --git a/tests/gtests/test_convolution_eltwise_forward_f32.cpp b/tests/gtests/test_convolution_eltwise_forward_f32.cpp index 89de760fc05..6fc8eac786d 100644 --- a/tests/gtests/test_convolution_eltwise_forward_f32.cpp +++ b/tests/gtests/test_convolution_eltwise_forward_f32.cpp @@ -64,8 +64,7 @@ TEST_P(convolution_test, TestConvolutionEltwise) {} EXPAND_ARGS(PARAMS_CONV(algorithm::eltwise_square, __VA_ARGS__)), \ EXPAND_ARGS(PARAMS_CONV(algorithm::eltwise_abs, __VA_ARGS__)), \ EXPAND_ARGS(PARAMS_CONV(algorithm::eltwise_linear, __VA_ARGS__)), \ - EXPAND_ARGS(PARAMS_CONV( \ - algorithm::eltwise_bounded_relu, __VA_ARGS__)), \ + EXPAND_ARGS(PARAMS_CONV(algorithm::eltwise_clip, __VA_ARGS__)), \ EXPAND_ARGS( \ PARAMS_CONV(algorithm::eltwise_soft_relu, __VA_ARGS__)), \ EXPAND_ARGS( \ diff --git a/tests/gtests/test_convolution_eltwise_forward_x8s8f32s32.cpp b/tests/gtests/test_convolution_eltwise_forward_x8s8f32s32.cpp index 52cc7551f1f..19d68e036ae 100644 --- a/tests/gtests/test_convolution_eltwise_forward_x8s8f32s32.cpp +++ b/tests/gtests/test_convolution_eltwise_forward_x8s8f32s32.cpp @@ -63,8 +63,7 @@ using convolution_test_s8s8s32f32 EXPAND_ARGS(PARAMS_CONV(algorithm::eltwise_square, __VA_ARGS__)), \ EXPAND_ARGS(PARAMS_CONV(algorithm::eltwise_abs, __VA_ARGS__)), \ EXPAND_ARGS(PARAMS_CONV(algorithm::eltwise_linear, __VA_ARGS__)), \ - EXPAND_ARGS(PARAMS_CONV( \ - algorithm::eltwise_bounded_relu, __VA_ARGS__)), \ + EXPAND_ARGS(PARAMS_CONV(algorithm::eltwise_clip, __VA_ARGS__)), \ EXPAND_ARGS( \ PARAMS_CONV(algorithm::eltwise_soft_relu, __VA_ARGS__)), \ EXPAND_ARGS( \ @@ -74,7 +73,7 @@ using convolution_test_s8s8s32f32 // EXPAND_ARGS(PARAMS_CONV(algorithm::eltwise_exp, __VA_ARGS__)) #define ELTWISE_ALPHA 0.5f -#define ELTWISE_BETA 0.f +#define ELTWISE_BETA 2.f #define PARAMS_CONV(alg, src, weights, bias, dst, ...) \ test_convolution_eltwise_params_t { \ diff --git a/tests/gtests/test_eltwise.cpp b/tests/gtests/test_eltwise.cpp index 851dc0cd96f..df3771ec353 100644 --- a/tests/gtests/test_eltwise.cpp +++ b/tests/gtests/test_eltwise.cpp @@ -101,14 +101,14 @@ T linear_bwd(T dd, T s, A alpha, A beta) { } template -T bounded_relu_fwd(T s, A alpha) { - s = s > 0 ? s : T(0); - return s > alpha ? T(alpha) : s; +T clip_fwd(T s, A alpha, A beta) { + s = s > alpha ? s : T(alpha); + return s > beta ? T(beta) : s; } template -T bounded_relu_bwd(T dd, T s, A alpha) { - return dd * ((0 < s && s < alpha) ? 1 : 0); +T clip_bwd(T dd, T s, A alpha, A beta) { + return dd * ((alpha < s && s < beta) ? 1 : 0); } template @@ -210,8 +210,8 @@ void check_eltwise_fwd(const eltwise_test_params_t &p, const memory::desc &md, case algorithm::eltwise_linear: ref_d = linear_fwd(s, p.alpha, p.beta); break; - case algorithm::eltwise_bounded_relu: - ref_d = bounded_relu_fwd(s, p.alpha); + case algorithm::eltwise_clip: + ref_d = clip_fwd(s, p.alpha, p.beta); break; case algorithm::eltwise_soft_relu: ref_d = soft_relu_fwd(s, p.alpha); @@ -297,8 +297,8 @@ void check_eltwise_bwd(const eltwise_test_params_t &p, const memory::desc &md, case algorithm::eltwise_linear: ref_ds = linear_bwd(ref_dd, ref_s, p.alpha, p.beta); break; - case algorithm::eltwise_bounded_relu: - ref_ds = bounded_relu_bwd(ref_dd, ref_s, p.alpha); + case algorithm::eltwise_clip: + ref_ds = clip_bwd(ref_dd, ref_s, p.alpha, p.beta); break; case algorithm::eltwise_soft_relu: ref_ds = soft_relu_bwd(ref_dd, ref_s, p.alpha); @@ -358,7 +358,6 @@ class eltwise_test_t : public ::testing::TestWithParam { SKIP_IF_HIP(data_type == memory::data_type::s8, "HIP backend doesn't support eltwise with s8"); SKIP_IF_CUDA(p.alg_kind != algorithm::eltwise_relu - && p.alg_kind != algorithm::eltwise_bounded_relu && p.alg_kind != algorithm::eltwise_tanh && p.alg_kind != algorithm::eltwise_elu && p.alg_kind != algorithm::eltwise_logistic, @@ -371,7 +370,6 @@ class eltwise_test_t : public ::testing::TestWithParam { "Unsupported format tag"); SKIP_IF_HIP(p.alg_kind != algorithm::eltwise_relu - && p.alg_kind != algorithm::eltwise_bounded_relu && p.alg_kind != algorithm::eltwise_tanh && p.alg_kind != algorithm::eltwise_elu && p.alg_kind != algorithm::eltwise_soft_relu @@ -464,15 +462,13 @@ class eltwise_test_t : public ::testing::TestWithParam { } void Backward() { - SKIP_IF_CUDA(p.alg_kind != algorithm::eltwise_relu - && p.alg_kind != algorithm::eltwise_bounded_relu, - "Unsupported algorithm"); + SKIP_IF_CUDA( + p.alg_kind != algorithm::eltwise_relu, "Unsupported algorithm"); SKIP_IF_CUDA(p.diff_format != p.data_format, "CUDA does not support different data formats for data and " "diff vectors"); SKIP_IF_HIP(p.alg_kind != algorithm::eltwise_relu - && p.alg_kind != algorithm::eltwise_soft_relu - && p.alg_kind != algorithm::eltwise_bounded_relu, + && p.alg_kind != algorithm::eltwise_soft_relu, "Unsupported algorithm"); SKIP_IF_HIP(p.diff_format != p.data_format, "HIP does not support different data formats for data and " @@ -580,7 +576,7 @@ TEST_P(eltwise_test_s8, TestsEltwise) {} #define PARAMS_ALL_ALG_SDPART(...) \ EXPAND(PARAMS(eltwise_linear, __VA_ARGS__)), \ - EXPAND(PARAMS(eltwise_bounded_relu, __VA_ARGS__)), \ + EXPAND(PARAMS(eltwise_clip, __VA_ARGS__)), \ EXPAND(PARAMS(eltwise_logistic, __VA_ARGS__)) #define _CPU_INST_TEST_CASE(str, data_t, ...) \ diff --git a/tests/gtests/test_iface_attr.cpp b/tests/gtests/test_iface_attr.cpp index fe1310622d6..33f399d25cf 100644 --- a/tests/gtests/test_iface_attr.cpp +++ b/tests/gtests/test_iface_attr.cpp @@ -345,7 +345,7 @@ HANDLE_EXCEPTIONS_FOR_TEST_F(attr_test_t, TestPostOps) { ASSERT_EQ(1, sum_zp); ASSERT_EQ(data_type::f32, dt); - ops.append_eltwise(2.2f, algorithm::eltwise_bounded_relu, 3.3f, 4.4f); + ops.append_eltwise(2.2f, algorithm::eltwise_clip, 3.3f, 4.4f); attr.set_post_ops(ops); ASSERT_EQ(attr.get_post_ops().len(), 2); @@ -353,7 +353,7 @@ HANDLE_EXCEPTIONS_FOR_TEST_F(attr_test_t, TestPostOps) { ASSERT_EQ(attr.get_post_ops().kind(1), primitive::kind::eltwise); attr.get_post_ops().get_params_eltwise(1, scale, alg, alpha, beta); ASSERT_FLOAT_EQ(scale, 2.2f); - ASSERT_EQ(alg, algorithm::eltwise_bounded_relu); + ASSERT_EQ(alg, algorithm::eltwise_clip); ASSERT_FLOAT_EQ(alpha, 3.3f); ASSERT_FLOAT_EQ(beta, 4.4f);