26 #if defined(DATA_TYPE) && defined(EPSILON)
105 int x = get_global_id(0);
106 int y = get_global_id(1);
107 int z = get_global_id(2);
112 #else // ! defined(DIM2)
116 #else // defined(NHWC)
118 #endif // defined(NHWC)
119 #endif // defined(DIM2)
121 int w_offset = x *
sizeof(DATA_TYPE) + y * w_stride_y + z * w_stride_z;
122 int v_offset = c1 *
sizeof(DATA_TYPE);
124 DATA_TYPE w_old = 0.0f;
125 DATA_TYPE b_old = 0.0f;
126 DATA_TYPE w_new = 0.0f;
127 DATA_TYPE b_new = 0.0f;
128 DATA_TYPE gamma = 1.0f;
129 DATA_TYPE mean = 0.0f;
130 DATA_TYPE var = 1.0f;
131 DATA_TYPE beta = 0.0f;
133 w_old = *((__global DATA_TYPE *)(w_ptr + w_offset + w_offset_first_element_in_bytes));
134 var = *((__global DATA_TYPE *)(var_ptr + v_offset + var_offset_first_element_in_bytes));
135 mean = *((__global DATA_TYPE *)(mean_ptr + v_offset + mean_offset_first_element_in_bytes));
138 gamma = *((__global DATA_TYPE *)(gamma_ptr + v_offset + gamma_offset_first_element_in_bytes));
139 #endif // defined(GAMMA)
142 w_new = (gamma * w_old) / (sqrt(var + EPSILON));
144 #if defined(IN_PLACE_W)
145 *((__global DATA_TYPE *)(w_ptr + w_offset + w_offset_first_element_in_bytes)) = w_new;
146 #else // defined(IN_PLACE_W)
147 *((__global DATA_TYPE *)(w_fused_ptr + w_offset + w_fused_offset_first_element_in_bytes)) = w_new;
148 #endif // defined(IN_PLACE_W)
151 #if !defined(DIM2) && defined(NHWC)
153 #else // !defined(DIM2) && defined(NHWC)
154 if(x == 0 && y == 0 && c0 == 0)
155 #endif // !defined(DIM2) && defined(NHWC)
158 b_old = *((__global DATA_TYPE *)(b_ptr + v_offset + b_offset_first_element_in_bytes));
159 #endif // defined(BIAS)
161 beta = *((__global DATA_TYPE *)(beta_ptr + v_offset + beta_offset_first_element_in_bytes));
162 #endif // defined(BETA)
164 b_new = ((gamma * (b_old - mean)) / (sqrt(var + EPSILON))) + beta;
168 #if defined(IN_PLACE_B)
169 *((__global DATA_TYPE *)(b_ptr + v_offset + b_offset_first_element_in_bytes)) = b_new;
170 #else // defined(IN_PLACE_B)
171 *((__global DATA_TYPE *)(b_fused_ptr + v_offset + b_fused_offset_first_element_in_bytes)) = b_new;
172 #endif // defined(IN_PLACE_B)
174 #else // defined(BIAS)
177 *((__global DATA_TYPE *)(b_fused_ptr + v_offset + b_fused_offset_first_element_in_bytes)) = b_new;
178 #endif // ifndef IN_PLACE_B
180 #endif // defined(BIAS)
183 #endif // defined(DATA_TYPE) && defined(EPSILON)