73 const int num_of_slices = get_global_size(2);
74 const int current_slice = get_global_id(2);
76 const int left_slice = max(current_slice - (
int)radius, (
int)0);
77 const int right_slice = min(current_slice + (
int)radius, (
int)(num_of_slices - 1));
79 for(
int i = left_slice; i <= right_slice; i++)
84 const float normalized = pow(kappa + coeff * (
float)acc, beta);
86 const float normalized_pixel = (float) * ((__global
DATA_TYPE *)in.
ptr) / normalized;
139 const int current_pos = get_global_id(0) << 2;
141 const int left_pos = max(current_pos - (
int)radius, -3);
142 const int right_pos = min(current_pos + (
int)radius, (
int)((get_global_size(0) << 2) + 3 - 1));
144 for(
int i = left_pos; i <= right_pos; i += 1)
149 const float4 normalized = pow((float4)kappa + coeff * (float4)acc_vec, beta);
151 const float4 normalized_pixel =
CONVERT(vload4(0, (__global
DATA_TYPE *)in.
ptr), float4) / normalized;
#define CONVERT_TO_TENSOR3D_STRUCT(name)
__kernel void normalization_layer_in_map(__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_stride_z, uint input_step_z, uint input_offset_first_element_in_bytes, __global uchar *squared_input_ptr, uint squared_input_stride_x, uint squared_input_step_x, uint squared_input_stride_y, uint squared_input_step_y, uint squared_input_stride_z, uint squared_input_step_z, uint squared_input_offset_first_element_in_bytes, __global uchar *output_ptr, uint output_stride_x, uint output_step_x, uint output_stride_y, uint output_step_y, uint output_stride_z, uint output_step_z, uint output_offset_first_element_in_bytes, float coeff, float beta, float kappa, uint radius)
Apply in map normalization.
Structure to hold 3D tensor information.
#define VEC_DATA_TYPE(type, size)
__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)
Get the pointer position of a Tensor3D.
#define TENSOR3D_DECLARATION(name)
__global uchar * ptr
Pointer to the starting postion of the buffer.
__kernel void normalization_layer_cross_map(__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_stride_z, uint input_step_z, uint input_offset_first_element_in_bytes, __global uchar *squared_input_ptr, uint squared_input_stride_x, uint squared_input_step_x, uint squared_input_stride_y, uint squared_input_step_y, uint squared_input_stride_z, uint squared_input_step_z, uint squared_input_offset_first_element_in_bytes, __global uchar *output_ptr, uint output_stride_x, uint output_step_x, uint output_stride_y, uint output_step_y, uint output_stride_z, uint output_step_z, uint output_offset_first_element_in_bytes, float coeff, float beta, float kappa, uint radius)
Apply cross map normalization.