gpu: bnorm: nhwc-reusable: rename some kernels

Author: skazakov1

src/gpu/ocl/bnorm/bnorm_utils.hpp

@@ -74,14 +74,14 @@ constexpr int aux_fwd = 1;
 constexpr int aux_bwd = 0;
 
 namespace kernel_id {
-constexpr size_t update_fwd = 0;
+constexpr size_t norm_fwd = 0;
 constexpr size_t calc_mean = 1;
 constexpr size_t calc_var = 2;
 constexpr size_t reduce_fwd_reg = 3;
 constexpr size_t calc_mean_var = 4;
 constexpr size_t reduce_fwd_1pass = 5;
 constexpr size_t reduce_aux = 6;
-constexpr size_t update_bwd = 7;
+constexpr size_t norm_bwd = 7;
 constexpr size_t calc_stat = 8;
 constexpr size_t reduce_stat = 9;
 } // namespace kernel_id

src/gpu/ocl/bnorm/nhwc_reusable.cl

@@ -336,7 +336,7 @@ nhwc_reusable_calc_mean_var(__global DATA_T *src, __global float *reduce_temp,
 
 // Main FWD kernel, common for regular and 1pass algorithms
 __attribute__((intel_reqd_sub_group_size(16))) __kernel void
-nhwc_reusable_update_fwd(__global DATA_T *src, __global float *mean,
+nhwc_reusable_norm_fwd(__global DATA_T *src, __global float *mean,
         __global float *variance, __global DATA_T *dst,
         __global float *scaleshift, __global float *shift, __global char *ws,
         float eps, __global DATA_T *src_add, float relu_alpha, off_t ic_size,
@@ -632,7 +632,7 @@ nhwc_reusable_calc_stat(__global DATA_T *src, __global float *mean,
 
 // Main BWD pass kernel
 __attribute__((intel_reqd_sub_group_size(16))) __kernel void
-nhwc_reusable_update_bwd(__global DATA_T *src, __global float *mean,
+nhwc_reusable_norm_bwd(__global DATA_T *src, __global float *mean,
         __global float *variance, __global DATA_T *diff_dst,
         __global float *scaleshift, __global char *ws,
         __global DATA_T *diff_src, __global float *diff_scale,

src/gpu/ocl/bnorm/nhwc_reusable.cpp

@@ -495,7 +495,7 @@ status_t nhwc_reusable_batch_normalization_fwd_t::execute_forward(
     arg_list.append(rt_conf.update_sp_block);
 
     auto nd_range = pd()->dispatch.nd_range();
-    return parallel_for(ctx, nd_range, kernels_[update_fwd], arg_list);
+    return parallel_for(ctx, nd_range, kernels_[norm_fwd], arg_list);
 }
 
 status_t nhwc_reusable_batch_normalization_bwd_t::pd_t::init_conf(
@@ -639,7 +639,7 @@ status_t nhwc_reusable_batch_normalization_bwd_t::execute_backward(
     arg_list.append(rt_conf.update_sp_block);
 
     auto nd_range = pd()->dispatch.nd_range();
-    return parallel_for(ctx, nd_range, kernels_[update_bwd], arg_list);
+    return parallel_for(ctx, nd_range, kernels_[norm_bwd], arg_list);
 }
 
 } // namespace ocl

src/gpu/ocl/bnorm/nhwc_reusable.hpp

@@ -50,10 +50,10 @@ struct nhwc_reusable_bnorm_compile_params_t {
 
     const std::vector<const char *> &get_kernel_names() const {
         static const std::vector<const char *> kernel_names = {
-                "nhwc_reusable_update_fwd", "nhwc_reusable_calc_mean",
+                "nhwc_reusable_norm_fwd", "nhwc_reusable_calc_mean",
                 "nhwc_reusable_calc_var", "nhwc_reusable_reduce_fwd_reg",
                 "nhwc_reusable_calc_mean_var", "nhwc_reusable_reduce_fwd_1pass",
-                "nhwc_reusable_reduce_aux", "nhwc_reusable_update_bwd",
+                "nhwc_reusable_reduce_aux", "nhwc_reusable_norm_bwd",
                 "nhwc_reusable_calc_stat", "nhwc_reusable_reduce_stat"};
         return kernel_names;
     }