gen9_bnorm_kernel.cpp source code [oneDNN/build/src/gpu/ocl/gen9_bnorm_kernel.cpp]

1	namespace dnnl {
2	namespace impl {
3	namespace gpu {
4	namespace ocl {
5	const char gen9_bnorm_kernel = R"==(/****************************************************************************** )==""\n"
6	R"==(* Copyright 2020-2022 Intel Corporation )==""\n"
7	R"==(* )==""\n"
8	R"==(* Licensed under the Apache License, Version 2.0 (the "License"); )==""\n"
9	R"==(* you may not use this file except in compliance with the License. )==""\n"
10	R"==(* You may obtain a copy of the License at )==""\n"
11	R"==(* )==""\n"
12	R"==(* http: )==""\n"
13	R"==(* )==""\n"
14	R"==(* Unless required by applicable law or agreed to in writing, software )==""\n"
15	R"==(* distributed under the License is distributed on an "AS IS" BASIS, )==""\n"
16	R"==(* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. )==""\n"
17	R"==(* See the License for the specific language governing permissions and )==""\n"
18	R"==(* limitations under the License. )==""\n"
19	R"==(*******************************************************************************/ )==""\n"
20	R"==(#define VECT_DT_N VECT_SIZE )==""\n"
21	R"==(#include "gpu/ocl/ocl_types.h" )==""\n"
22	R"==(#define IS_IC_EQ_8 (IC == 8) )==""\n"
23	R"==(#define HAS_IC_TAIL (IC != IC16) )==""\n"
24	R"==(#define HAS_STAT_SP_BLOCK_TAIL (SP % STAT_SP_BLOCK) )==""\n"
25	R"==(#if NHWC_OPTIMIZED )==""\n"
26	R"==(#if HAS_IC_TAIL )==""\n"
27	R"==(#error IC tail processing not supported )==""\n"
28	R"==(#endif )==""\n"
29	R"==(#else )==""\n"
30	R"==(#if HAS_IC_TAIL && !USE_NHWC )==""\n"
31	R"==(#error IC tail processing not supported )==""\n"
32	R"==(#endif )==""\n"
33	R"==(#define HAS_STAT_SP_TAIL (STAT_SP_TAIL != STAT_SP_NBLOCKS) )==""\n"
34	R"==(#define HAS_SP_TAIL (SP != SP_TAIL) )==""\n"
35	R"==(#endif )==""\n"
36	R"==(#define IC_BLOCK_SGROUPS (IC_BLOCK / 16) )==""\n"
37	R"==(#define IC_TAIL_SGROUPS (IC_BLOCK_SGROUPS % VECT_SIZE) )==""\n"
38	R"==(#define IC_VECT_SGROUPS (IC_BLOCK_SGROUPS - IC_TAIL_SGROUPS) )==""\n"
39	R"==(#define HAS_IC_VECT_TAIL (IC_TAIL_SGROUPS > 0) )==""\n"
40	R"==(#define LOAD_FLOAT_1x16(ptr) \ )==""\n"
41	R"==(as_float(intel_sub_group_block_read((const __global uint *)(ptr))) )==""\n"
42	R"==(#define LOAD_UINT_1x16(ptr) \ )==""\n"
43	R"==(as_uint(intel_sub_group_block_read((const __global uint *)(ptr))) )==""\n"
44	R"==(#define LOAD_UINT_8x16(ptr) \ )==""\n"
45	R"==(convert_uint8(as_uint8( \ )==""\n"
46	R"==(intel_sub_group_block_read8((const __global uint *)(ptr)))) )==""\n"
47	R"==(#define LOAD_CHAR_1x16(ptr) \ )==""\n"
48	R"==(as_char(intel_sub_group_block_read_uc((const __global uchar *)(ptr))) )==""\n"
49	R"==(#define LOAD_CHAR_8x16(ptr) \ )==""\n"
50	R"==(convert_char8(as_char8( \ )==""\n"
51	R"==(intel_sub_group_block_read_uc8((const __global uchar *)(ptr)))) )==""\n"
52	R"==(#define LOAD_DATA_1x16(ptr) \ )==""\n"
53	R"==(CONVERT_FLOAT_T(AS_DATA_T(BLOCK_READ((const __global BLOCK_DATA_T *)(ptr)))) )==""\n"
54	R"==(#define LOAD_DATA_8x16(ptr) \ )==""\n"
55	R"==(CONVERT_FLOAT8_T( \ )==""\n"
56	R"==(AS_DATA8_T(BLOCK_READ8((const __global BLOCK_DATA_T *)(ptr)))) )==""\n"
57	R"==(#define LOAD_VECT_DATA(ptr) \ )==""\n"
58	R"==(CONVERT_VECT_FLOAT_T(AS_VECT_DATA_T( \ )==""\n"
59	R"==(VECT_BLOCK_READ((const __global BLOCK_DATA_T *)(ptr)))) )==""\n"
60	R"==(#define LOAD_VECT_CHAR(ptr) \ )==""\n"
61	R"==(CONVERT_VECT_CHAR_T( \ )==""\n"
62	R"==(AS_VECT_CHAR_T(VECT_UCHAR_READ((const __global uchar *)(ptr)))) )==""\n"
63	R"==(#define LOAD_VECT_FLOAT(ptr) \ )==""\n"
64	R"==(AS_VECT_FLOAT_T(VECT_UINT_READ((const __global uint *)(ptr))) )==""\n"
65	R"==(#define STORE_DATA_1x16(ptr, val) \ )==""\n"
66	R"==(BLOCK_WRITE((__global BLOCK_DATA_T *)(ptr), \ )==""\n"
67	R"==(AS_BLOCK_DATA_T(CONVERT_DATA_T(val))) )==""\n"
68	R"==(#define STORE_DATA_8x16(ptr, val) \ )==""\n"
69	R"==(BLOCK_WRITE8((__global BLOCK_DATA_T *)ptr, \ )==""\n"
70	R"==(AS_BLOCK_DATA8_T(CONVERT_DATA8_T(val))) )==""\n"
71	R"==(#define STORE_VECT_DATA(ptr, val) \ )==""\n"
72	R"==(VECT_BLOCK_WRITE((__global BLOCK_DATA_T *)(ptr), \ )==""\n"
73	R"==(AS_VECT_BLOCK_DATA_T(CONVERT_VECTOR_DATA_T(val))) )==""\n"
74	R"==(#define STORE_FLOAT_1x16(ptr, val) \ )==""\n"
75	R"==(intel_sub_group_block_write((__global uint *)(ptr), as_uint(val)) )==""\n"
76	R"==(#define STORE_FLOAT_8x16(ptr, val) \ )==""\n"
77	R"==(intel_sub_group_block_write8((__global uint *)(ptr), as_uint8(val)) )==""\n"
78	R"==(#define STORE_CHAR_1x16(ptr, val) \ )==""\n"
79	R"==(intel_sub_group_block_write_uc((__global uchar *)(ptr), as_uchar(val)) )==""\n"
80	R"==(#define STORE_CHAR_8x16(ptr, val) \ )==""\n"
81	R"==(intel_sub_group_block_write_uc8((__global uchar *)(ptr), as_uchar8(val)) )==""\n"
82	R"==(#define STORE_VECT_CHAR(ptr, val) \ )==""\n"
83	R"==(VECT_UCHAR_WRITE((__global uchar *)(ptr), \ )==""\n"
84	R"==(AS_VECT_UCHAR_T(CONVERT_VECT_CHAR_T(val))) )==""\n"
85	R"==(#if HAS_IC_TAIL )==""\n"
86	R"==(#define MAYBE_LAST_IC_LOAD_FLOAT_1x16(ptr, idx) \ )==""\n"
87	R"==((is_last_ic_block ? (simd_id < 8 ? ptr[(idx) + simd_id] : 0.0f) \ )==""\n"
88	R"==(: as_float(intel_sub_group_block_read( \ )==""\n"
89	R"==((const __global uint *)(&ptr[(idx)])))) )==""\n"
90	R"==(#else )==""\n"
91	R"==(#define MAYBE_LAST_IC_LOAD_FLOAT_1x16(ptr, idx) LOAD_FLOAT_1x16(&ptr[(idx)]) )==""\n"
92	R"==(#endif )==""\n"
93	R"==(#if USE_NHWC )==""\n"
94	R"==(#define IC_BLOCK_STRIDE IC )==""\n"
95	R"==(#else )==""\n"
96	R"==(#define IC_BLOCK_STRIDE 16 )==""\n"
97	R"==(#endif )==""\n"
98	R"==(#if NHWC_OPTIMIZED )==""\n"
99	R"==(#define REDUCE_NUM_SGROUPS IC_BLOCK_SGROUPS )==""\n"
100	R"==(#else )==""\n"
101	R"==(#define REDUCE_NUM_SGROUPS 1 )==""\n"
102	R"==(#endif )==""\n"
103	R"==(#define CALC_SLM_LINE_SIZE (REDUCE_NUM_SGROUPS * GWS_LWS0_CALC) )==""\n"
104	R"==(#define CALC_SLM_SIZE (CALC_SLM_LINE_SIZE * GWS_LWS1_CALC * GWS_LWS2_CALC) )==""\n"
105	R"==(NAMED_KERNEL_ATTR(AUX) )==""\n"
106	R"==(__kernel void gen9_fused_reduce_init( )==""\n"
107	R"==(#if IS_FWD )==""\n"
108	R"==(__global float mean, __global float variance )==""\n"
109	R"==(#else )==""\n"
110	R"==(__global float diff_scale, __global float diff_shift )==""\n"
111	R"==(#endif )==""\n"
112	R"==() { )==""\n"
113	R"==(const int c = GWS_GET_IC_AUX(); )==""\n"
114	R"==(#if IS_FWD )==""\n"
115	R"==(mean[c] = 0.0f; )==""\n"
116	R"==(variance[c] = 0.0f; )==""\n"
117	R"==(#else )==""\n"
118	R"==(diff_scale[c] = 0.0f; )==""\n"
119	R"==(#if DIFF_SHIFT == 1 )==""\n"
120	R"==(diff_shift[c] = 0.0f; )==""\n"
121	R"==(#else )==""\n"
122	R"==(diff_shift[IC + IC * REDUCE_STAT_NBLOCKS + c] = 0.0f; )==""\n"
123	R"==(#endif )==""\n"
124	R"==(#endif )==""\n"
125	R"==(return; )==""\n"
126	R"==(} )==""\n"
127	R"==(#if IS_FWD )==""\n"
128	R"==(#define LOAD_DATA_Nx16_USING_LOOP_IDX(n, dest, src, idx) \ )==""\n"
129	R"==({ \ )==""\n"
130	R"==(for (int k = 0; k < n; ++k) { \ )==""\n"
131	R"==(dest[k] = LOAD_DATA_1x16(&src[(k + idx) * IC]); \ )==""\n"
132	R"==(} \ )==""\n"
133	R"==(} )==""\n"
134	R"==(#define LOAD_DATA_Nx16_USING_LOOP_IDX_HALF(n, dest, src, idx) \ )==""\n"
135	R"==({ \ )==""\n"
136	R"==(for (int k = 0; k < n; k += 2) { \ )==""\n"
137	R"==(dest[k] = LOAD_DATA_1x16(&src[(k + idx) * IC]); \ )==""\n"
138	R"==(} \ )==""\n"
139	R"==(} )==""\n"
140	R"==(#if USE_STATS_ONE_PASS )==""\n"
141	R"==(#define ACCUM_DATA_T float )==""\n"
142	R"==(#define ACCUM_DATA8_T float8 )==""\n"
143	R"==(#define ACCUM_DATA2_T float2 )==""\n"
144	R"==(#define SUM_DATA_T ACCUM_DATA2_T )==""\n"
145	R"==(SUM_DATA_T summation(ACCUM_DATA_T input, SUM_DATA_T state) { )==""\n"
146	R"==(ACCUM_DATA2_T ret; )==""\n"
147	R"==(ACCUM_DATA_T y = input - state.s1; )==""\n"
148	R"==(ACCUM_DATA_T t = state.s0 + y; )==""\n"
149	R"==(ret.s1 = (t - state.s0) - y; )==""\n"
150	R"==(ret.s0 = t; )==""\n"
151	R"==(return ret; )==""\n"
152	R"==(} )==""\n"
153	R"==(#if FUSED_ATOMICS_REDUCTION )==""\n"
154	R"==(void gen9_mean_var_calc_fused_reduction(volatile __global atomic_float *mean, )==""\n"
155	R"==(volatile __global atomic_float *variance, int dst_offset, )==""\n"
156	R"==(SUM_DATA_T sum, SUM_DATA_T sum_sq, __local SUM_DATA_T *local_sum, )==""\n"
157	R"==(__local SUM_DATA_T *local_sum_sq) { )==""\n"
158	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
159	R"==(const int group_size = GWS_LWS1_CALC * GWS_LWS2_CALC; )==""\n"
160	R"==(const int sg_group_id = get_local_id(0) / 16; )==""\n"
161	R"==(const int local_id = get_local_id(1); )==""\n"
162	R"==(if (local_id > 0) { )==""\n"
163	R"==(for (int sg = 0; sg < REDUCE_NUM_SGROUPS; ++sg) { )==""\n"
164	R"==(const int slm_offset = CALC_SLM_LINE_SIZE * local_id )==""\n"
165	R"==(+ REDUCE_NUM_SGROUPS * 16 * sg_group_id + sg * 16 + simd_id; )==""\n"
166	R"==(local_sum[slm_offset] = sum[sg]; )==""\n"
167	R"==(local_sum_sq[slm_offset] = sum_sq[sg]; )==""\n"
168	R"==(} )==""\n"
169	R"==(} )==""\n"
170	R"==(barrier(CLK_LOCAL_MEM_FENCE); )==""\n"
171	R"==(if (local_id == 0) { )==""\n"
172	R"==(for (int sg = 0; sg < REDUCE_NUM_SGROUPS; ++sg) { )==""\n"
173	R"==(for (int gr_id = 1; gr_id < group_size; ++gr_id) { )==""\n"
174	R"==(const int off_local = CALC_SLM_LINE_SIZE * gr_id )==""\n"
175	R"==(+ REDUCE_NUM_SGROUPS * 16 * sg_group_id + sg * 16 )==""\n"
176	R"==(+ simd_id; )==""\n"
177	R"==(SUM_DATA_T tmp = local_sum[off_local]; )==""\n"
178	R"==(SUM_DATA_T tmp_sq = local_sum_sq[off_local]; )==""\n"
179	R"==(sum[sg] = summation(tmp.s1, sum[sg]); )==""\n"
180	R"==(sum_sq[sg] = summation(tmp_sq.s1, sum_sq[sg]); )==""\n"
181	R"==(sum[sg] = summation(tmp.s0, sum[sg]); )==""\n"
182	R"==(sum_sq[sg] = summation(tmp_sq.s0, sum_sq[sg]); )==""\n"
183	R"==(} )==""\n"
184	R"==(const int offset = dst_offset + sg * 16 + simd_id; )==""\n"
185	R"==(#if HAS_IC_TAIL )==""\n"
186	R"==(if (offset < IC) { )==""\n"
187	R"==(#endif )==""\n"
188	R"==(atomic_add_global(&mean[offset], sum[sg].s0); )==""\n"
189	R"==(atomic_add_global(&variance[offset], sum_sq[sg].s0); )==""\n"
190	R"==(#if HAS_IC_TAIL )==""\n"
191	R"==(} )==""\n"
192	R"==(#endif )==""\n"
193	R"==(} )==""\n"
194	R"==(} )==""\n"
195	R"==(} )==""\n"
196	R"==(#endif )==""\n"
197	R"==(#endif )==""\n"
198	R"==(#if FUSED_ATOMICS_REDUCTION )==""\n"
199	R"==(void gen9_calc_fused_reduction(volatile __global atomic_float *dst, )==""\n"
200	R"==(int dst_offset, float sum, __local float local_sum) { )==""\n"
201	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
202	R"==(const int group_size = GWS_LWS1_CALC * GWS_LWS2_CALC; )==""\n"
203	R"==(const int sg_group_id = get_local_id(0) / 16; )==""\n"
204	R"==(const int local_id = get_local_id(1); )==""\n"
205	R"==(if (local_id > 0) { )==""\n"
206	R"==(for (int sg = 0; sg < REDUCE_NUM_SGROUPS; ++sg) { )==""\n"
207	R"==(const int slm_offset = CALC_SLM_LINE_SIZE * local_id )==""\n"
208	R"==(+ REDUCE_NUM_SGROUPS * 16 * sg_group_id + sg * 16 + simd_id; )==""\n"
209	R"==(local_sum[slm_offset] = sum[sg]; )==""\n"
210	R"==(} )==""\n"
211	R"==(} )==""\n"
212	R"==(barrier(CLK_LOCAL_MEM_FENCE); )==""\n"
213	R"==(if (local_id == 0) { )==""\n"
214	R"==(for (int sg = 0; sg < REDUCE_NUM_SGROUPS; ++sg) { )==""\n"
215	R"==(for (int gr_id = 1; gr_id < group_size; ++gr_id) { )==""\n"
216	R"==(const int off_local = CALC_SLM_LINE_SIZE * gr_id )==""\n"
217	R"==(+ REDUCE_NUM_SGROUPS * 16 * sg_group_id + sg * 16 )==""\n"
218	R"==(+ simd_id; )==""\n"
219	R"==(sum[sg] += local_sum[off_local]; )==""\n"
220	R"==(} )==""\n"
221	R"==(const int offset = dst_offset + sg * 16 + simd_id; )==""\n"
222	R"==(#if HAS_IC_TAIL )==""\n"
223	R"==(if (offset < IC) )==""\n"
224	R"==(#endif )==""\n"
225	R"==(atomic_add_global(&dst[offset], sum[sg]); )==""\n"
226	R"==(} )==""\n"
227	R"==(} )==""\n"
228	R"==(return; )==""\n"
229	R"==(} )==""\n"
230	R"==(#endif )==""\n"
231	R"==(void gen9_reduce_common(__global float reduce_temp, __local float local_sum, )==""\n"
232	R"==(__global float *dst) { )==""\n"
233	R"==(const int ic_sub_group = get_global_id(0) / 16; )==""\n"
234	R"==(const int group_c = get_global_id(1); )==""\n"
235	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
236	R"==(const int c = group_c * 16 + simd_id; )==""\n"
237	R"==(const bool is_last_ic_block = (IC - group_c * 16) < 16; )==""\n"
238	R"==(float sum = 0.0f; )==""\n"
239	R"==(reduce_temp )==""\n"
240	R"==(+= REDUCE_STAT_NBLOCKS / REDUCE_IC_SUB_GROUPS * 16 * ic_sub_group )==""\n"
241	R"==(+ REDUCE_STAT_NBLOCKS * 16 * group_c + simd_id; )==""\n"
242	R"==(for (int i = 0; i < REDUCE_STAT_NBLOCKS / REDUCE_IC_SUB_GROUPS; i++) { )==""\n"
243	R"==(sum += reduce_temp[i * 16]; )==""\n"
244	R"==(} )==""\n"
245	R"==(if (ic_sub_group > 0) { local_sum[ic_sub_group * 16 + simd_id] = sum; } )==""\n"
246	R"==(barrier(CLK_LOCAL_MEM_FENCE); )==""\n"
247	R"==(if (ic_sub_group == 0) { )==""\n"
248	R"==(for (int i = 1; i < REDUCE_IC_SUB_GROUPS; i++) { )==""\n"
249	R"==(sum += local_sum[i * 16 + simd_id]; )==""\n"
250	R"==(} )==""\n"
251	R"==(#if HAS_IC_TAIL )==""\n"
252	R"==(if (!is_last_ic_block \|\| (is_last_ic_block && simd_id < 8)) )==""\n"
253	R"==(#endif )==""\n"
254	R"==(dst[c] = sum / (MB * ID * IH * IW); )==""\n"
255	R"==(} )==""\n"
256	R"==(} )==""\n"
257	R"==(#if USE_STATS_ONE_PASS )==""\n"
258	R"==(#if NHWC_OPTIMIZED )==""\n"
259	R"==(NAMED_KERNEL_ATTR(CALC) )==""\n"
260	R"==(__kernel void gen9_calc_mean_var(__global DATA_T *src, )==""\n"
261	R"==(__global ACCUM_DATA_T *reduce_temp, )==""\n"
262	R"==(volatile __global atomic_float *mean, )==""\n"
263	R"==(volatile __global atomic_float *variance) { )==""\n"
264	R"==(const int mb = GWS_GET_STAT_MB(); )==""\n"
265	R"==(const int c = GWS_GET_STAT_IC(); )==""\n"
266	R"==(const int sp_block_idx = GWS_GET_STAT_SP(); )==""\n"
267	R"==(const int ic_block_offset = (c / 16) * IC_BLOCK; )==""\n"
268	R"==(const int group_c_offset )==""\n"
269	R"==(= REDUCE_STAT_NBLOCKS * ic_block_offset + sp_block_idx * 16; )==""\n"
270	R"==(const int ver_offs = REDUCE_STAT_NBLOCKS * IC; )==""\n"
271	R"==(const int src_off = ic_block_offset + sp_block_idx * STAT_SP_BLOCK * IC; )==""\n"
272	R"==(src += src_off; )==""\n"
273	R"==(SUM_DATA_T sum[IC_BLOCK_SGROUPS] = {0.0f}; )==""\n"
274	R"==(SUM_DATA_T sum_sq[IC_BLOCK_SGROUPS] = {0.0f}; )==""\n"
275	R"==(#if HAS_STAT_SP_BLOCK_TAIL )==""\n"
276	R"==(for (int sp = 0; sp < min(STAT_SP_BLOCK, SP - sp_block_idx * STAT_SP_BLOCK); )==""\n"
277	R"==(++sp) { )==""\n"
278	R"==(#else )==""\n"
279	R"==(for (int sp = 0; sp < STAT_SP_BLOCK; ++sp) { )==""\n"
280	R"==(#endif )==""\n"
281	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS / VECT_SIZE; ++sg) { )==""\n"
282	R"==(VECT_FLOAT_T s_vect = LOAD_VECT_DATA(&src[sg * 16 * VECT_SIZE]); )==""\n"
283	R"==(for (int vect = 0; vect < VECT_SIZE; ++vect) { )==""\n"
284	R"==(const int sum_idx = sg * VECT_SIZE + vect; )==""\n"
285	R"==(#if VECT_SIZE > 1 )==""\n"
286	R"==(sum[sum_idx] = summation(s_vect[vect], sum[sum_idx]); )==""\n"
287	R"==(sum_sq[sum_idx] = summation( )==""\n"
288	R"==(s_vect[vect] * s_vect[vect], sum_sq[sum_idx]); )==""\n"
289	R"==(#else )==""\n"
290	R"==(sum[sum_idx] = summation(s_vect, sum[sum_idx]); )==""\n"
291	R"==(sum_sq[sum_idx] = summation(s_vect * s_vect, sum_sq[sum_idx]); )==""\n"
292	R"==(#endif )==""\n"
293	R"==(} )==""\n"
294	R"==(} )==""\n"
295	R"==(#if HAS_IC_VECT_TAIL )==""\n"
296	R"==(for (int sg = 0; sg < IC_TAIL_SGROUPS; ++sg) { )==""\n"
297	R"==(const int sg_idx = IC_VECT_SGROUPS + sg; )==""\n"
298	R"==(float s_tail = LOAD_DATA_1x16(&src[(IC_VECT_SGROUPS + sg) * 16]); )==""\n"
299	R"==(sum[sg_idx] = summation(s_tail, sum[sg_idx]); )==""\n"
300	R"==(sum_sq[sg_idx] = summation(s_tail * s_tail, sum_sq[sg_idx]); )==""\n"
301	R"==(} )==""\n"
302	R"==(#endif )==""\n"
303	R"==(src += IC; )==""\n"
304	R"==(} )==""\n"
305	R"==(#if FUSED_ATOMICS_REDUCTION )==""\n"
306	R"==(__local SUM_DATA_T local_sum[2 * CALC_SLM_SIZE]; )==""\n"
307	R"==(__local SUM_DATA_T *local_sum_sq = local_sum + CALC_SLM_SIZE; )==""\n"
308	R"==(gen9_mean_var_calc_fused_reduction(mean, variance, ic_block_offset, sum, )==""\n"
309	R"==(sum_sq, local_sum, local_sum_sq); )==""\n"
310	R"==(#else )==""\n"
311	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS; ++sg) { )==""\n"
312	R"==(const int reduce_off = group_c_offset + sg * 16 * REDUCE_STAT_NBLOCKS; )==""\n"
313	R"==(STORE_FLOAT_1x16(&reduce_temp[reduce_off], sum[sg].s0); )==""\n"
314	R"==(STORE_FLOAT_1x16(&reduce_temp[ver_offs + reduce_off], sum_sq[sg].s0); )==""\n"
315	R"==(} )==""\n"
316	R"==(#endif )==""\n"
317	R"==(} )==""\n"
318	R"==(#else )==""\n"
319	R"==(NAMED_KERNEL_ATTR(CALC) )==""\n"
320	R"==(__kernel void gen9_calc_mean_var(__global DATA_T *src, )==""\n"
321	R"==(__global ACCUM_DATA_T *reduce_temp, )==""\n"
322	R"==(volatile __global atomic_float *mean, )==""\n"
323	R"==(volatile __global atomic_float *variance) { )==""\n"
324	R"==(const int mb = GWS_GET_STAT_MB(); )==""\n"
325	R"==(const int c = GWS_GET_STAT_IC(); )==""\n"
326	R"==(const int sp_block_idx = GWS_GET_STAT_SP(); )==""\n"
327	R"==(const int mb_sp_idx = mb * STAT_SP_NBLOCKS + sp_block_idx; )==""\n"
328	R"==(const int group_c_offset = REDUCE_STAT_NBLOCKS * 16 * (int)(c / 16); )==""\n"
329	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
330	R"==(const int ver_offs = REDUCE_STAT_NBLOCKS * IC; )==""\n"
331	R"==(#if USE_NHWC )==""\n"
332	R"==(src += c + sp_block_idx * STAT_SP_BLOCK * IC; )==""\n"
333	R"==(#else )==""\n"
334	R"==(src += (c & 15) + sp_block_idx * STAT_SP_BLOCK * 16 + (c & ~15) * SP )==""\n"
335	R"==(+ mb * SP * IC; )==""\n"
336	R"==(#endif )==""\n"
337	R"==(SUM_DATA_T sum; )==""\n"
338	R"==(SUM_DATA_T sum_sq; )==""\n"
339	R"==(sum.s0 = 0; )==""\n"
340	R"==(sum.s1 = 0; )==""\n"
341	R"==(sum_sq.s0 = 0; )==""\n"
342	R"==(sum_sq.s1 = 0; )==""\n"
343	R"==(#if HAS_STAT_SP_TAIL )==""\n"
344	R"==(if (sp_block_idx == STAT_SP_TAIL) { )==""\n"
345	R"==(int sp = SP - STAT_SP_TAIL * STAT_SP_BLOCK; )==""\n"
346	R"==(while (sp >= 16) { )==""\n"
347	R"==(#if USE_NHWC )==""\n"
348	R"==(float8 s0, s1; )==""\n"
349	R"==(for (int k = 0; k < 8; ++k) )==""\n"
350	R"==(s0[k] = LOAD_DATA_1x16(&src[k * IC]); )==""\n"
351	R"==(for (int k = 0; k < 8; ++k) )==""\n"
352	R"==(s1[k] = LOAD_DATA_1x16(&src[(k + 8) * IC]); )==""\n"
353	R"==(#else )==""\n"
354	R"==(float8 s0 = LOAD_DATA_8x16(&src[0]); )==""\n"
355	R"==(float8 s1 = LOAD_DATA_8x16(&src[8 * 16]); )==""\n"
356	R"==(#endif )==""\n"
357	R"==(for (int i = 0; i < 8; i++) { )==""\n"
358	R"==(sum = summation(s0[i], sum); )==""\n"
359	R"==(sum = summation(s1[i], sum); )==""\n"
360	R"==(sum_sq = summation(s0[i] * s0[i], sum_sq); )==""\n"
361	R"==(sum_sq = summation(s1[i] * s1[i], sum_sq); )==""\n"
362	R"==(} )==""\n"
363	R"==(src += 16 * IC_BLOCK_STRIDE; )==""\n"
364	R"==(sp -= 16; )==""\n"
365	R"==(} )==""\n"
366	R"==(while (sp >= 1) { )==""\n"
367	R"==(float s0 = LOAD_DATA_1x16(&src[0]); )==""\n"
368	R"==(sum = summation(s0, sum); )==""\n"
369	R"==(sum_sq = summation(s0 * s0, sum_sq); )==""\n"
370	R"==(src += IC_BLOCK_STRIDE; )==""\n"
371	R"==(--sp; )==""\n"
372	R"==(} )==""\n"
373	R"==(} else )==""\n"
374	R"==(#endif )==""\n"
375	R"==({ )==""\n"
376	R"==(for (int sp = 0; sp < STAT_SP_BLOCK / 16; ++sp) { )==""\n"
377	R"==(#if USE_NHWC )==""\n"
378	R"==(float8 s0, s1; )==""\n"
379	R"==(for (int k = 0; k < 8; ++k) )==""\n"
380	R"==(s0[k] = LOAD_DATA_1x16(&src[k * IC]); )==""\n"
381	R"==(for (int k = 0; k < 8; ++k) )==""\n"
382	R"==(s1[k] = LOAD_DATA_1x16(&src[(k + 8) * IC]); )==""\n"
383	R"==(#else )==""\n"
384	R"==(float8 s0 = LOAD_DATA_8x16(&src[0]); )==""\n"
385	R"==(float8 s1 = LOAD_DATA_8x16(&src[8 * 16]); )==""\n"
386	R"==(#endif )==""\n"
387	R"==(for (int i = 0; i < 8; i++) { )==""\n"
388	R"==(sum = summation(s0[i], sum); )==""\n"
389	R"==(sum = summation(s1[i], sum); )==""\n"
390	R"==(sum_sq = summation(s0[i] * s0[i], sum_sq); )==""\n"
391	R"==(sum_sq = summation(s1[i] * s1[i], sum_sq); )==""\n"
392	R"==(} )==""\n"
393	R"==(src += 16 * IC_BLOCK_STRIDE; )==""\n"
394	R"==(} )==""\n"
395	R"==(} )==""\n"
396	R"==(#if FUSED_ATOMICS_REDUCTION )==""\n"
397	R"==(__local SUM_DATA_T local_sum[2 * CALC_SLM_SIZE]; )==""\n"
398	R"==(__local SUM_DATA_T *local_sum_sq = local_sum + CALC_SLM_SIZE; )==""\n"
399	R"==(gen9_mean_var_calc_fused_reduction( )==""\n"
400	R"==(mean, variance, c, &sum, &sum_sq, local_sum, local_sum_sq); )==""\n"
401	R"==(#else )==""\n"
402	R"==(STORE_FLOAT_1x16(&reduce_temp[group_c_offset + mb_sp_idx * 16], sum.s0); )==""\n"
403	R"==(STORE_FLOAT_1x16(&reduce_temp[ver_offs + group_c_offset + mb_sp_idx * 16], )==""\n"
404	R"==(sum_sq.s0); )==""\n"
405	R"==(#endif )==""\n"
406	R"==(} )==""\n"
407	R"==(#endif )==""\n"
408	R"==(NAMED_KERNEL_ATTR(REDUCE) )==""\n"
409	R"==(__kernel void gen9_reduce_mean_var(__global ACCUM_DATA_T *reduce_temp, )==""\n"
410	R"==(__global float mean, __global float variance) { )==""\n"
411	R"==(__local SUM_DATA_T local_sum[16 * REDUCE_IC_SUB_GROUPS]; )==""\n"
412	R"==(__local SUM_DATA_T local_sum_sq[16 * REDUCE_IC_SUB_GROUPS]; )==""\n"
413	R"==(const int ic_sub_group = get_global_id(0) / 16; )==""\n"
414	R"==(const int group_c = get_global_id(1); )==""\n"
415	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
416	R"==(const int c = group_c * 16 + simd_id; )==""\n"
417	R"==(SUM_DATA_T sum; )==""\n"
418	R"==(SUM_DATA_T sum_sq; )==""\n"
419	R"==(sum.s0 = 0; )==""\n"
420	R"==(sum.s1 = 0; )==""\n"
421	R"==(sum_sq.s0 = 0; )==""\n"
422	R"==(sum_sq.s1 = 0; )==""\n"
423	R"==(int offs_sq = REDUCE_STAT_NBLOCKS * IC; )==""\n"
424	R"==(int offs = REDUCE_STAT_NBLOCKS / REDUCE_IC_SUB_GROUPS * 16 * ic_sub_group )==""\n"
425	R"==(+ REDUCE_STAT_NBLOCKS * 16 * group_c + simd_id; )==""\n"
426	R"==(for (int i = 0; i < REDUCE_STAT_NBLOCKS / REDUCE_IC_SUB_GROUPS; i++) { )==""\n"
427	R"==(float tmp = reduce_temp[offs + i * 16]; )==""\n"
428	R"==(sum = summation(tmp, sum); )==""\n"
429	R"==(} )==""\n"
430	R"==(for (int i = 0; i < REDUCE_STAT_NBLOCKS / REDUCE_IC_SUB_GROUPS; i++) { )==""\n"
431	R"==(float tmp = reduce_temp[offs_sq + offs + i * 16]; )==""\n"
432	R"==(sum_sq = summation(tmp, sum_sq); )==""\n"
433	R"==(} )==""\n"
434	R"==(if (ic_sub_group > 0) { )==""\n"
435	R"==(local_sum[ic_sub_group * 16 + simd_id] = sum; )==""\n"
436	R"==(local_sum_sq[ic_sub_group * 16 + simd_id] = sum_sq; )==""\n"
437	R"==(} )==""\n"
438	R"==(barrier(CLK_LOCAL_MEM_FENCE); )==""\n"
439	R"==(if (ic_sub_group == 0) { )==""\n"
440	R"==(for (int i = 1; i < REDUCE_IC_SUB_GROUPS; i++) { )==""\n"
441	R"==(SUM_DATA_T tmp = local_sum[i * 16 + simd_id]; )==""\n"
442	R"==(SUM_DATA_T tmp_sq = local_sum_sq[i * 16 + simd_id]; )==""\n"
443	R"==(sum = summation(tmp.s1, sum); )==""\n"
444	R"==(sum_sq = summation(tmp_sq.s1, sum_sq); )==""\n"
445	R"==(sum = summation(tmp.s0, sum); )==""\n"
446	R"==(sum_sq = summation(tmp_sq.s0, sum_sq); )==""\n"
447	R"==(} )==""\n"
448	R"==(float tmp_mean = sum.s0 / (MB * ID * IH * IW); )==""\n"
449	R"==(mean[c] = tmp_mean; )==""\n"
450	R"==(float tmp_var = max(0.0f, )==""\n"
451	R"==((sum_sq.s0 / (MB * ID * IH * IW)) - (tmp_mean * tmp_mean)); )==""\n"
452	R"==(variance[c] = tmp_var; )==""\n"
453	R"==(} )==""\n"
454	R"==(} )==""\n"
455	R"==(#endif )==""\n"
456	R"==(NAMED_KERNEL_ATTR(AUX) )==""\n"
457	R"==(__kernel void gen9_fused_reduce_final( )==""\n"
458	R"==(#if USE_STATS_ONE_PASS )==""\n"
459	R"==(__global float mean, __global float variance )==""\n"
460	R"==(#else )==""\n"
461	R"==(__global float *data_reduce )==""\n"
462	R"==(#endif )==""\n"
463	R"==() { )==""\n"
464	R"==(const int c = GWS_GET_IC_AUX(); )==""\n"
465	R"==(#if USE_STATS_ONE_PASS )==""\n"
466	R"==(mean[c] = mean[c] / (MB * ID * IH * IW); )==""\n"
467	R"==(float tmp_var = max( )==""\n"
468	R"==(0.0f, (variance[c] / (MB * ID * IH * IW)) - (mean[c] * mean[c])); )==""\n"
469	R"==(variance[c] = tmp_var; )==""\n"
470	R"==(#else )==""\n"
471	R"==(data_reduce[c] = data_reduce[c] / (MB * ID * IH * IW); )==""\n"
472	R"==(#endif )==""\n"
473	R"==(return; )==""\n"
474	R"==(} )==""\n"
475	R"==(#if NHWC_OPTIMIZED )==""\n"
476	R"==(NAMED_KERNEL_ATTR(CALC) )==""\n"
477	R"==(__kernel void gen9_calc_mean(__global DATA_T src, __global float reduce_temp, )==""\n"
478	R"==(volatile __global atomic_float *mean) { )==""\n"
479	R"==(const int mb = GWS_GET_STAT_MB(); )==""\n"
480	R"==(const int c = GWS_GET_STAT_IC(); )==""\n"
481	R"==(const int sp_block_idx = GWS_GET_STAT_SP(); )==""\n"
482	R"==(const int ic_block_offset = (c / 16) * IC_BLOCK; )==""\n"
483	R"==(const int group_c_offset )==""\n"
484	R"==(= REDUCE_STAT_NBLOCKS * ic_block_offset + sp_block_idx * 16; )==""\n"
485	R"==(const int src_off = ic_block_offset + sp_block_idx * STAT_SP_BLOCK * IC; )==""\n"
486	R"==(src += src_off; )==""\n"
487	R"==(float v_mean[IC_BLOCK_SGROUPS] = {0.0f}; )==""\n"
488	R"==(#if HAS_STAT_SP_BLOCK_TAIL )==""\n"
489	R"==(for (int sp = 0; sp < min(STAT_SP_BLOCK, SP - sp_block_idx * STAT_SP_BLOCK); )==""\n"
490	R"==(++sp) { )==""\n"
491	R"==(#else )==""\n"
492	R"==(for (int sp = 0; sp < STAT_SP_BLOCK; ++sp) { )==""\n"
493	R"==(#endif )==""\n"
494	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS / VECT_SIZE; ++sg) { )==""\n"
495	R"==(VECT_FLOAT_T s_vect = LOAD_VECT_DATA(&src[sg * 16 * VECT_SIZE]); )==""\n"
496	R"==(for (int vect = 0; vect < VECT_SIZE; ++vect) { )==""\n"
497	R"==(v_mean[sg * VECT_SIZE + vect] )==""\n"
498	R"==(#if VECT_SIZE > 1 )==""\n"
499	R"==(+= s_vect[vect]; )==""\n"
500	R"==(#else )==""\n"
501	R"==(+= s_vect; )==""\n"
502	R"==(#endif )==""\n"
503	R"==(} )==""\n"
504	R"==(} )==""\n"
505	R"==(#if HAS_IC_VECT_TAIL )==""\n"
506	R"==(for (int sg = 0; sg < IC_TAIL_SGROUPS; ++sg) { )==""\n"
507	R"==(float s_tail = LOAD_DATA_1x16(&src[(IC_VECT_SGROUPS + sg) * 16]); )==""\n"
508	R"==(v_mean[IC_VECT_SGROUPS + sg] += s_tail; )==""\n"
509	R"==(} )==""\n"
510	R"==(#endif )==""\n"
511	R"==(src += IC; )==""\n"
512	R"==(} )==""\n"
513	R"==(#if FUSED_ATOMICS_REDUCTION )==""\n"
514	R"==(__local float local_sum[CALC_SLM_SIZE]; )==""\n"
515	R"==(gen9_calc_fused_reduction(mean, ic_block_offset, v_mean, local_sum); )==""\n"
516	R"==(#else )==""\n"
517	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS; ++sg) { )==""\n"
518	R"==(const int reduce_off = group_c_offset + sg * 16 * REDUCE_STAT_NBLOCKS; )==""\n"
519	R"==(STORE_FLOAT_1x16(&reduce_temp[reduce_off], v_mean[sg]); )==""\n"
520	R"==(} )==""\n"
521	R"==(#endif )==""\n"
522	R"==(} )==""\n"
523	R"==(#else )==""\n"
524	R"==(NAMED_KERNEL_ATTR(CALC) )==""\n"
525	R"==(__kernel void gen9_calc_mean(__global DATA_T src, __global float reduce_temp, )==""\n"
526	R"==(__global float *mean) { )==""\n"
527	R"==(const int mb = GWS_GET_STAT_MB(); )==""\n"
528	R"==(const int c = GWS_GET_STAT_IC(); )==""\n"
529	R"==(const int sp_block_idx = GWS_GET_STAT_SP(); )==""\n"
530	R"==(const int mb_sp_idx = mb * STAT_SP_NBLOCKS + sp_block_idx; )==""\n"
531	R"==(const int group_c_offset = REDUCE_STAT_NBLOCKS * 16 * (int)(c / 16); )==""\n"
532	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
533	R"==(#if HAS_IC_TAIL )==""\n"
534	R"==(const bool is_last_ic_block = c + 16 > IC; )==""\n"
535	R"==(const bool is_last_sp_block = (sp_block_idx == STAT_SP_NBLOCKS - 1); )==""\n"
536	R"==(#endif )==""\n"
537	R"==(#if USE_NHWC )==""\n"
538	R"==(src += c + sp_block_idx * STAT_SP_BLOCK * IC; )==""\n"
539	R"==(#else )==""\n"
540	R"==(src += (c & 15) + sp_block_idx * STAT_SP_BLOCK * 16 + (c & ~15) * SP )==""\n"
541	R"==(+ mb * SP * IC; )==""\n"
542	R"==(#endif )==""\n"
543	R"==(float8 res0 = 0.0f, res1 = 0.0f; )==""\n"
544	R"==(float v_mean = 0.0f; )==""\n"
545	R"==(#if HAS_STAT_SP_TAIL )==""\n"
546	R"==(if (sp_block_idx == STAT_SP_TAIL) { )==""\n"
547	R"==(int sp = SP - STAT_SP_TAIL * STAT_SP_BLOCK; )==""\n"
548	R"==(while (sp >= 16) { )==""\n"
549	R"==(#if USE_NHWC )==""\n"
550	R"==(float8 s0, s1; )==""\n"
551	R"==(#if IS_IC_EQ_8 )==""\n"
552	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX_HALF(8, s0, src, 0); )==""\n"
553	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX_HALF(8, s1, src, 8); )==""\n"
554	R"==(float8 t0 = intel_sub_group_shuffle_down(s0, s0, 8); )==""\n"
555	R"==(float8 t1 = intel_sub_group_shuffle_down(s1, s1, 8); )==""\n"
556	R"==(for (int k = 0; k < 7; k += 2) { )==""\n"
557	R"==(s0[k + 1] = t0[k]; )==""\n"
558	R"==(s1[k + 1] = t1[k]; )==""\n"
559	R"==(} )==""\n"
560	R"==(#elif HAS_IC_TAIL )==""\n"
561	R"==(const bool is_last_sp = sp == 16; )==""\n"
562	R"==(if (is_last_sp && is_last_ic_block) { )==""\n"
563	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(7, s0, src, 0); )==""\n"
564	R"==(s0[7] = simd_id < 8 ? CONVERT_FLOAT_T(src[7 * IC + simd_id]) )==""\n"
565	R"==(: 0.0f; )==""\n"
566	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(7, s1, src, 8); )==""\n"
567	R"==(s1[7] = simd_id < 8 ? CONVERT_FLOAT_T(src[15 * IC + simd_id]) )==""\n"
568	R"==(: 0.0f; )==""\n"
569	R"==(} else { )==""\n"
570	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s0, src, 0); )==""\n"
571	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s1, src, 8); )==""\n"
572	R"==(} )==""\n"
573	R"==(#else )==""\n"
574	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s0, src, 0); )==""\n"
575	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s1, src, 8); )==""\n"
576	R"==(#endif )==""\n"
577	R"==(#else )==""\n"
578	R"==(float8 s0 = LOAD_DATA_8x16(&src[0]); )==""\n"
579	R"==(float8 s1 = LOAD_DATA_8x16(&src[8 * 16]); )==""\n"
580	R"==(#endif )==""\n"
581	R"==(res0 += s0; )==""\n"
582	R"==(res1 += s1; )==""\n"
583	R"==(src += 16 * IC_BLOCK_STRIDE; )==""\n"
584	R"==(sp -= 16; )==""\n"
585	R"==(} )==""\n"
586	R"==(while (sp >= 1) { )==""\n"
587	R"==(#if HAS_IC_TAIL )==""\n"
588	R"==(float s0; )==""\n"
589	R"==(if (sp == 1 && is_last_ic_block) )==""\n"
590	R"==(s0 = simd_id < 8 ? CONVERT_FLOAT_T(src[simd_id]) : 0.0f; )==""\n"
591	R"==(else )==""\n"
592	R"==(s0 = LOAD_DATA_1x16(&src[0]); )==""\n"
593	R"==(#else )==""\n"
594	R"==(float s0 = LOAD_DATA_1x16(&src[0]); )==""\n"
595	R"==(#endif )==""\n"
596	R"==(v_mean += s0; )==""\n"
597	R"==(src += IC_BLOCK_STRIDE; )==""\n"
598	R"==(--sp; )==""\n"
599	R"==(} )==""\n"
600	R"==(} else )==""\n"
601	R"==(#endif )==""\n"
602	R"==({ )==""\n"
603	R"==(for (int sp = 0; sp < STAT_SP_BLOCK / 16; ++sp) { )==""\n"
604	R"==(#if USE_NHWC )==""\n"
605	R"==(float8 s0, s1; )==""\n"
606	R"==(#if IS_IC_EQ_8 )==""\n"
607	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX_HALF(8, s0, src, 0); )==""\n"
608	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX_HALF(8, s1, src, 8); )==""\n"
609	R"==(float8 t0 = intel_sub_group_shuffle_down(s0, s0, 8); )==""\n"
610	R"==(float8 t1 = intel_sub_group_shuffle_down(s1, s1, 8); )==""\n"
611	R"==(for (int k = 0; k < 7; k += 2) { )==""\n"
612	R"==(s0[k + 1] = t0[k]; )==""\n"
613	R"==(s1[k + 1] = t1[k]; )==""\n"
614	R"==(} )==""\n"
615	R"==(#elif HAS_IC_TAIL )==""\n"
616	R"==(const bool is_last_sp = sp == STAT_SP_BLOCK / 16 - 1; )==""\n"
617	R"==(if (is_last_sp && is_last_ic_block && is_last_sp_block) { )==""\n"
618	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(7, s0, src, 0); )==""\n"
619	R"==(s0[7] = simd_id < 8 ? CONVERT_FLOAT_T(src[7 * IC + simd_id]) )==""\n"
620	R"==(: 0.0f; )==""\n"
621	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(7, s1, src, 8); )==""\n"
622	R"==(s1[7] = simd_id < 8 ? CONVERT_FLOAT_T(src[15 * IC + simd_id]) )==""\n"
623	R"==(: 0.0f; )==""\n"
624	R"==(} else { )==""\n"
625	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s0, src, 0); )==""\n"
626	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s1, src, 8); )==""\n"
627	R"==(} )==""\n"
628	R"==(#else )==""\n"
629	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s0, src, 0); )==""\n"
630	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s1, src, 8); )==""\n"
631	R"==(#endif )==""\n"
632	R"==(#else )==""\n"
633	R"==(float8 s0 = LOAD_DATA_8x16(&src[0]); )==""\n"
634	R"==(float8 s1 = LOAD_DATA_8x16(&src[8 * 16]); )==""\n"
635	R"==(#endif )==""\n"
636	R"==(res0 += s0; )==""\n"
637	R"==(res1 += s1; )==""\n"
638	R"==(src += 16 * IC_BLOCK_STRIDE; )==""\n"
639	R"==(} )==""\n"
640	R"==(} )==""\n"
641	R"==(for (int i = 0; i < 8; i++) { )==""\n"
642	R"==(v_mean += res0[i] + res1[i]; )==""\n"
643	R"==(} )==""\n"
644	R"==(#if FUSED_ATOMICS_REDUCTION )==""\n"
645	R"==(__local float local_sum[CALC_SLM_SIZE]; )==""\n"
646	R"==(gen9_calc_fused_reduction(mean, c, &v_mean, local_sum); )==""\n"
647	R"==(#else )==""\n"
648	R"==(STORE_FLOAT_1x16(&reduce_temp[group_c_offset + mb_sp_idx * 16], v_mean); )==""\n"
649	R"==(#endif )==""\n"
650	R"==(} )==""\n"
651	R"==(#endif )==""\n"
652	R"==(NAMED_KERNEL_ATTR(REDUCE) )==""\n"
653	R"==(__kernel void gen9_reduce_mean( )==""\n"
654	R"==(__global float reduce_temp, __global float mean) { )==""\n"
655	R"==(__local float local_sum[16 * REDUCE_IC_SUB_GROUPS]; )==""\n"
656	R"==(gen9_reduce_common(reduce_temp, local_sum, mean); )==""\n"
657	R"==(} )==""\n"
658	R"==(#if NHWC_OPTIMIZED )==""\n"
659	R"==(NAMED_KERNEL_ATTR(CALC) )==""\n"
660	R"==(__kernel void gen9_calc_variance(__global DATA_T src, __global float mean, )==""\n"
661	R"==(__global float reduce_temp, volatile __global atomic_float variance) { )==""\n"
662	R"==(const int mb = GWS_GET_STAT_MB(); )==""\n"
663	R"==(const int c = GWS_GET_STAT_IC(); )==""\n"
664	R"==(const int sp_block_idx = GWS_GET_STAT_SP(); )==""\n"
665	R"==(const int ic_block_offset = (c / 16) * IC_BLOCK; )==""\n"
666	R"==(const int group_c_offset )==""\n"
667	R"==(= REDUCE_STAT_NBLOCKS * ic_block_offset + sp_block_idx * 16; )==""\n"
668	R"==(reduce_temp += REDUCE_STAT_NBLOCKS * IC16; )==""\n"
669	R"==(mean += ic_block_offset; )==""\n"
670	R"==(const int src_off = ic_block_offset + sp_block_idx * STAT_SP_BLOCK * IC; )==""\n"
671	R"==(src += src_off; )==""\n"
672	R"==(float v_mean[IC_BLOCK_SGROUPS]; )==""\n"
673	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS; ++sg) { )==""\n"
674	R"==(v_mean[sg] = as_float(intel_sub_group_block_read( )==""\n"
675	R"==((const __global uint )(&mean[(sg 16)]))); )==""\n"
676	R"==(} )==""\n"
677	R"==(float v_var[IC_BLOCK_SGROUPS] = {0.0f}; )==""\n"
678	R"==(float v0[IC_BLOCK_SGROUPS] = {0.0f}; )==""\n"
679	R"==(#if HAS_STAT_SP_BLOCK_TAIL )==""\n"
680	R"==(for (int sp = 0; sp < min(STAT_SP_BLOCK, SP - sp_block_idx * STAT_SP_BLOCK); )==""\n"
681	R"==(++sp) { )==""\n"
682	R"==(#else )==""\n"
683	R"==(for (int sp = 0; sp < STAT_SP_BLOCK; ++sp) { )==""\n"
684	R"==(#endif )==""\n"
685	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS / VECT_SIZE; ++sg) { )==""\n"
686	R"==(VECT_FLOAT_T s_vect = LOAD_VECT_DATA(&src[sg * 16 * VECT_SIZE]); )==""\n"
687	R"==(for (int vect = 0; vect < VECT_SIZE; ++vect) { )==""\n"
688	R"==(int sg_idx = sg * VECT_SIZE + vect; )==""\n"
689	R"==(#if VECT_SIZE > 1 )==""\n"
690	R"==(v0[sg_idx] = s_vect[vect] - v_mean[sg_idx]; )==""\n"
691	R"==(#else )==""\n"
692	R"==(v0[sg_idx] = s_vect - v_mean[sg_idx]; )==""\n"
693	R"==(#endif )==""\n"
694	R"==(v_var[sg_idx] = fma(v0[sg_idx], v0[sg_idx], v_var[sg_idx]); )==""\n"
695	R"==(} )==""\n"
696	R"==(} )==""\n"
697	R"==(#if HAS_IC_VECT_TAIL )==""\n"
698	R"==(for (int sg = 0; sg < IC_TAIL_SGROUPS; ++sg) { )==""\n"
699	R"==(const int sg_idx = IC_VECT_SGROUPS + sg; )==""\n"
700	R"==(float s_tail = LOAD_DATA_1x16(&src[(IC_VECT_SGROUPS + sg) * 16]); )==""\n"
701	R"==(v0[sg_idx] = s_tail - v_mean[sg_idx]; )==""\n"
702	R"==(v_var[sg_idx] = fma(v0[sg_idx], v0[sg_idx], v_var[sg_idx]); )==""\n"
703	R"==(} )==""\n"
704	R"==(#endif )==""\n"
705	R"==(src += IC; )==""\n"
706	R"==(} )==""\n"
707	R"==(#if FUSED_ATOMICS_REDUCTION )==""\n"
708	R"==(__local float local_sum[CALC_SLM_SIZE]; )==""\n"
709	R"==(gen9_calc_fused_reduction(variance, ic_block_offset, v_var, local_sum); )==""\n"
710	R"==(#else )==""\n"
711	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS; ++sg) { )==""\n"
712	R"==(const int reduce_off = group_c_offset + sg * 16 * REDUCE_STAT_NBLOCKS; )==""\n"
713	R"==(STORE_FLOAT_1x16(&reduce_temp[reduce_off], v_var[sg]); )==""\n"
714	R"==(} )==""\n"
715	R"==(#endif )==""\n"
716	R"==(} )==""\n"
717	R"==(#else )==""\n"
718	R"==(NAMED_KERNEL_ATTR(CALC) )==""\n"
719	R"==(__kernel void gen9_calc_variance(__global DATA_T src, __global float mean, )==""\n"
720	R"==(__global float reduce_temp, __global float variance) { )==""\n"
721	R"==(const int mb = GWS_GET_STAT_MB(); )==""\n"
722	R"==(const int c = GWS_GET_STAT_IC(); )==""\n"
723	R"==(const int sp_block_idx = GWS_GET_STAT_SP(); )==""\n"
724	R"==(const int mb_sp_idx = mb * STAT_SP_NBLOCKS + sp_block_idx; )==""\n"
725	R"==(const int group_c_offset = REDUCE_STAT_NBLOCKS * 16 * (int)(c / 16); )==""\n"
726	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
727	R"==(#if HAS_IC_TAIL )==""\n"
728	R"==(const bool is_last_ic_block = c + 16 > IC; )==""\n"
729	R"==(const bool is_last_sp_block = (sp_block_idx == STAT_SP_NBLOCKS - 1); )==""\n"
730	R"==(#endif )==""\n"
731	R"==(reduce_temp += REDUCE_STAT_NBLOCKS * IC16; )==""\n"
732	R"==(#if USE_NHWC )==""\n"
733	R"==(src += c + sp_block_idx * STAT_SP_BLOCK * IC; )==""\n"
734	R"==(#else )==""\n"
735	R"==(src += (c & 15) + sp_block_idx * STAT_SP_BLOCK * 16 + (c & ~15) * SP )==""\n"
736	R"==(+ mb * SP * IC; )==""\n"
737	R"==(#endif )==""\n"
738	R"==(float8 res0 = 0.0f, res1 = 0.0f; )==""\n"
739	R"==(float v_var = 0.0f; )==""\n"
740	R"==(float v_mean = MAYBE_LAST_IC_LOAD_FLOAT_1x16(mean, c); )==""\n"
741	R"==(#if HAS_STAT_SP_TAIL )==""\n"
742	R"==(if (sp_block_idx == STAT_SP_TAIL) { )==""\n"
743	R"==(int sp = SP - STAT_SP_TAIL * STAT_SP_BLOCK; )==""\n"
744	R"==(while (sp >= 16) { )==""\n"
745	R"==(#if USE_NHWC )==""\n"
746	R"==(float8 s0, s1; )==""\n"
747	R"==(#if IS_IC_EQ_8 )==""\n"
748	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX_HALF(8, s0, src, 0); )==""\n"
749	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX_HALF(8, s1, src, 8); )==""\n"
750	R"==(float8 t0 = intel_sub_group_shuffle_down(s0, s0, 8); )==""\n"
751	R"==(float8 t1 = intel_sub_group_shuffle_down(s1, s1, 8); )==""\n"
752	R"==(for (int k = 0; k < 7; k += 2) { )==""\n"
753	R"==(s0[k + 1] = t0[k]; )==""\n"
754	R"==(s1[k + 1] = t1[k]; )==""\n"
755	R"==(} )==""\n"
756	R"==(#elif HAS_IC_TAIL )==""\n"
757	R"==(const bool is_last_sp = sp == 16; )==""\n"
758	R"==(if (is_last_sp && is_last_ic_block) { )==""\n"
759	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(7, s0, src, 0); )==""\n"
760	R"==(s0[7] = simd_id < 8 ? CONVERT_FLOAT_T(src[7 * IC + simd_id]) )==""\n"
761	R"==(: 0.0f; )==""\n"
762	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(7, s1, src, 8); )==""\n"
763	R"==(s1[7] = simd_id < 8 ? CONVERT_FLOAT_T(src[15 * IC + simd_id]) )==""\n"
764	R"==(: 0.0f; )==""\n"
765	R"==(} else { )==""\n"
766	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s0, src, 0); )==""\n"
767	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s1, src, 8); )==""\n"
768	R"==(} )==""\n"
769	R"==(#else )==""\n"
770	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s0, src, 0); )==""\n"
771	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s1, src, 8); )==""\n"
772	R"==(#endif )==""\n"
773	R"==(#else )==""\n"
774	R"==(float8 s0 = LOAD_DATA_8x16(&src[0]); )==""\n"
775	R"==(float8 s1 = LOAD_DATA_8x16(&src[8 * 16]); )==""\n"
776	R"==(#endif )==""\n"
777	R"==(float8 v0 = s0 - v_mean; )==""\n"
778	R"==(float8 v1 = s1 - v_mean; )==""\n"
779	R"==(res0 = fma(v0, v0, res0); )==""\n"
780	R"==(res1 = fma(v1, v1, res1); )==""\n"
781	R"==(src += 16 * IC_BLOCK_STRIDE; )==""\n"
782	R"==(sp -= 16; )==""\n"
783	R"==(} )==""\n"
784	R"==(while (sp >= 1) { )==""\n"
785	R"==(#if HAS_IC_TAIL )==""\n"
786	R"==(float s0; )==""\n"
787	R"==(if (sp == 1 && is_last_ic_block) )==""\n"
788	R"==(s0 = simd_id < 8 ? CONVERT_FLOAT_T(src[simd_id]) : 0.0f; )==""\n"
789	R"==(else )==""\n"
790	R"==(s0 = LOAD_DATA_1x16(&src[0]); )==""\n"
791	R"==(#else )==""\n"
792	R"==(float s0 = LOAD_DATA_1x16(&src[0]); )==""\n"
793	R"==(#endif )==""\n"
794	R"==(float v0 = s0 - v_mean; )==""\n"
795	R"==(v_var = fma(v0, v0, v_var); )==""\n"
796	R"==(src += IC_BLOCK_STRIDE; )==""\n"
797	R"==(--sp; )==""\n"
798	R"==(} )==""\n"
799	R"==(} else )==""\n"
800	R"==(#endif )==""\n"
801	R"==({ )==""\n"
802	R"==(for (int sp = 0; sp < STAT_SP_BLOCK / 16; ++sp) { )==""\n"
803	R"==(#if USE_NHWC )==""\n"
804	R"==(float8 s0, s1; )==""\n"
805	R"==(#if IS_IC_EQ_8 )==""\n"
806	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX_HALF(8, s0, src, 0); )==""\n"
807	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX_HALF(8, s1, src, 8); )==""\n"
808	R"==(float8 t0 = intel_sub_group_shuffle_down(s0, s0, 8); )==""\n"
809	R"==(float8 t1 = intel_sub_group_shuffle_down(s1, s1, 8); )==""\n"
810	R"==(for (int k = 0; k < 7; k += 2) { )==""\n"
811	R"==(s0[k + 1] = t0[k]; )==""\n"
812	R"==(s1[k + 1] = t1[k]; )==""\n"
813	R"==(} )==""\n"
814	R"==(#elif HAS_IC_TAIL )==""\n"
815	R"==(const bool is_last_sp = sp == STAT_SP_BLOCK / 16 - 1; )==""\n"
816	R"==(if (is_last_sp && is_last_ic_block && is_last_sp_block) { )==""\n"
817	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(7, s0, src, 0); )==""\n"
818	R"==(s0[7] = simd_id < 8 ? CONVERT_FLOAT_T(src[7 * IC + simd_id]) )==""\n"
819	R"==(: 0.0f; )==""\n"
820	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(7, s1, src, 8); )==""\n"
821	R"==(s1[7] = simd_id < 8 ? CONVERT_FLOAT_T(src[15 * IC + simd_id]) )==""\n"
822	R"==(: 0.0f; )==""\n"
823	R"==(} else { )==""\n"
824	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s0, src, 0); )==""\n"
825	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s1, src, 8); )==""\n"
826	R"==(} )==""\n"
827	R"==(#else )==""\n"
828	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s0, src, 0); )==""\n"
829	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, s1, src, 8); )==""\n"
830	R"==(#endif )==""\n"
831	R"==(#else )==""\n"
832	R"==(float8 s0 = LOAD_DATA_8x16(&src[0]); )==""\n"
833	R"==(float8 s1 = LOAD_DATA_8x16(&src[8 * 16]); )==""\n"
834	R"==(#endif )==""\n"
835	R"==(float8 v0 = s0 - v_mean; )==""\n"
836	R"==(float8 v1 = s1 - v_mean; )==""\n"
837	R"==(res0 = fma(v0, v0, res0); )==""\n"
838	R"==(res1 = fma(v1, v1, res1); )==""\n"
839	R"==(src += 16 * IC_BLOCK_STRIDE; )==""\n"
840	R"==(} )==""\n"
841	R"==(} )==""\n"
842	R"==(for (int i = 0; i < 8; i++) { )==""\n"
843	R"==(v_var += res0[i] + res1[i]; )==""\n"
844	R"==(} )==""\n"
845	R"==(#if FUSED_ATOMICS_REDUCTION )==""\n"
846	R"==(__local float local_sum[CALC_SLM_SIZE]; )==""\n"
847	R"==(gen9_calc_fused_reduction(variance, c, &v_var, local_sum); )==""\n"
848	R"==(#else )==""\n"
849	R"==(STORE_FLOAT_1x16(&reduce_temp[group_c_offset + mb_sp_idx * 16], v_var); )==""\n"
850	R"==(#endif )==""\n"
851	R"==(} )==""\n"
852	R"==(#endif )==""\n"
853	R"==(NAMED_KERNEL_ATTR(REDUCE) )==""\n"
854	R"==(__kernel void gen9_reduce_variance( )==""\n"
855	R"==(__global float reduce_temp, __global float variance) { )==""\n"
856	R"==(__local float local_sum[16 * REDUCE_IC_SUB_GROUPS]; )==""\n"
857	R"==(gen9_reduce_common( )==""\n"
858	R"==(reduce_temp + REDUCE_STAT_NBLOCKS * IC16, local_sum, variance); )==""\n"
859	R"==(} )==""\n"
860	R"==(#if NHWC_OPTIMIZED )==""\n"
861	R"==(KERNEL_ATTR )==""\n"
862	R"==(__kernel void gen9_bnorm_fwd(__global DATA_T src, __global float mean, )==""\n"
863	R"==(__global float variance, __global DATA_T dst, )==""\n"
864	R"==(__global float scaleshift, __global float shift, __global char *ws, )==""\n"
865	R"==(float eps, __global DATA_T *src_add) { )==""\n"
866	R"==(const int n = GWS_GET_MB(); )==""\n"
867	R"==(const int c = GWS_GET_IC(); )==""\n"
868	R"==(const int sp = GWS_GET_SP() * STAT_SP_BLOCK; )==""\n"
869	R"==(const int ic_block_offset = (c / 16) * IC_BLOCK; )==""\n"
870	R"==(mean += ic_block_offset; )==""\n"
871	R"==(variance += ic_block_offset; )==""\n"
872	R"==(shift += ic_block_offset; )==""\n"
873	R"==(scaleshift += ic_block_offset; )==""\n"
874	R"==(const uint d_off = sp * IC + ic_block_offset; )==""\n"
875	R"==(src += d_off; )==""\n"
876	R"==(#if FUSE_BN_ADD_RELU )==""\n"
877	R"==(src_add += d_off; )==""\n"
878	R"==(#endif )==""\n"
879	R"==(dst += d_off; )==""\n"
880	R"==(#if FUSE_BN_RELU && IS_TRAINING )==""\n"
881	R"==(ws += d_off; )==""\n"
882	R"==(#endif )==""\n"
883	R"==(VECT_FLOAT_T sm[IC_BLOCK_SGROUPS / VECT_SIZE], )==""\n"
884	R"==(sv[IC_BLOCK_SGROUPS / VECT_SIZE], )==""\n"
885	R"==(v_mean[IC_BLOCK_SGROUPS / VECT_SIZE], )==""\n"
886	R"==(v_variance[IC_BLOCK_SGROUPS / VECT_SIZE], )==""\n"
887	R"==(sqrt_variance[IC_BLOCK_SGROUPS / VECT_SIZE]; )==""\n"
888	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS / VECT_SIZE; ++sg) { )==""\n"
889	R"==(const int sg_idx = sg * 16 * VECT_SIZE; )==""\n"
890	R"==(#if USE_SCALE == 1 )==""\n"
891	R"==(sm[sg] = LOAD_VECT_FLOAT(&scaleshift[sg_idx]); )==""\n"
892	R"==(#else )==""\n"
893	R"==(sm[sg] = (VECT_FLOAT_T)1.0f; )==""\n"
894	R"==(#endif )==""\n"
895	R"==(#if USE_SHIFT == 1 )==""\n"
896	R"==(sv[sg] = LOAD_VECT_FLOAT(&shift[sg_idx]); )==""\n"
897	R"==(#else )==""\n"
898	R"==(sv[sg] = (VECT_FLOAT_T)0.0f; )==""\n"
899	R"==(#endif )==""\n"
900	R"==(v_mean[sg] = LOAD_VECT_FLOAT(&mean[sg_idx]); )==""\n"
901	R"==(v_variance[sg] = LOAD_VECT_FLOAT(&variance[sg_idx]); )==""\n"
902	R"==(sqrt_variance[sg] = sm[sg] / sqrt(v_variance[sg] + (VECT_FLOAT_T)eps); )==""\n"
903	R"==(} )==""\n"
904	R"==(#if HAS_IC_VECT_TAIL )==""\n"
905	R"==(float sm_tail[IC_TAIL_SGROUPS], sv_tail[IC_TAIL_SGROUPS], )==""\n"
906	R"==(v_mean_tail[IC_TAIL_SGROUPS], v_variance_tail[IC_TAIL_SGROUPS], )==""\n"
907	R"==(sqrt_variance_tail[IC_TAIL_SGROUPS]; )==""\n"
908	R"==(for (int sg = 0; sg < IC_TAIL_SGROUPS; ++sg) { )==""\n"
909	R"==(const int sg_idx = (IC_VECT_SGROUPS + sg) * 16; )==""\n"
910	R"==(#if USE_SCALE == 1 )==""\n"
911	R"==(sm_tail[sg] = LOAD_FLOAT_1x16(&scaleshift[sg_idx]); )==""\n"
912	R"==(#else )==""\n"
913	R"==(sm_tail[sg] = 1.0f; )==""\n"
914	R"==(#endif )==""\n"
915	R"==(#if USE_SHIFT == 1 )==""\n"
916	R"==(sv_tail[sg] = LOAD_FLOAT_1x16(&shift[sg_idx]); )==""\n"
917	R"==(#else )==""\n"
918	R"==(sv_tail[sg] = 0.0f; )==""\n"
919	R"==(#endif )==""\n"
920	R"==(v_mean_tail[sg] = LOAD_FLOAT_1x16(&mean[sg_idx]); )==""\n"
921	R"==(v_variance_tail[sg] = LOAD_FLOAT_1x16(&variance[sg_idx]); )==""\n"
922	R"==(sqrt_variance_tail[sg] = sm_tail[sg] / sqrt(v_variance_tail[sg] + eps); )==""\n"
923	R"==(} )==""\n"
924	R"==(#endif )==""\n"
925	R"==(#if HAS_STAT_SP_BLOCK_TAIL )==""\n"
926	R"==(for (int sp_idx = 0; sp_idx < min(STAT_SP_BLOCK, SP - sp); ++sp_idx) { )==""\n"
927	R"==(#else )==""\n"
928	R"==(for (int sp_idx = 0; sp_idx < STAT_SP_BLOCK; ++sp_idx) { )==""\n"
929	R"==(#endif )==""\n"
930	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS / VECT_SIZE; ++sg) { )==""\n"
931	R"==(const int sg_idx = sg * 16 * VECT_SIZE; )==""\n"
932	R"==(VECT_FLOAT_T s_vect = LOAD_VECT_DATA(&src[sg_idx]); )==""\n"
933	R"==(VECT_FLOAT_T d_vect )==""\n"
934	R"==(= fma(s_vect - v_mean[sg], sqrt_variance[sg], sv[sg]); )==""\n"
935	R"==(#if FUSE_BN_RELU )==""\n"
936	R"==(#if FUSE_BN_ADD_RELU )==""\n"
937	R"==(VECT_FLOAT_T s_add_vect = LOAD_VECT_DATA(&src_add[sg_idx]); )==""\n"
938	R"==(d_vect += s_add_vect; )==""\n"
939	R"==(#endif )==""\n"
940	R"==(VECT_INT_T ws_vect = isgreater(d_vect, (VECT_FLOAT_T)0.0f); )==""\n"
941	R"==(d_vect = select((VECT_FLOAT_T)0.0f, d_vect, ws_vect); )==""\n"
942	R"==(#if IS_TRAINING )==""\n"
943	R"==(STORE_VECT_CHAR(&ws[sg_idx], ws_vect); )==""\n"
944	R"==(#endif )==""\n"
945	R"==(#endif )==""\n"
946	R"==(#if WITH_RELU )==""\n"
947	R"==(d_vect = max(d_vect, (VECT_FLOAT_T)0.0f); )==""\n"
948	R"==(#endif )==""\n"
949	R"==(STORE_VECT_DATA(&dst[sg_idx], d_vect); )==""\n"
950	R"==(} )==""\n"
951	R"==(#if HAS_IC_VECT_TAIL )==""\n"
952	R"==(for (int sg = 0; sg < IC_TAIL_SGROUPS; ++sg) { )==""\n"
953	R"==(const int sg_idx = (IC_VECT_SGROUPS + sg) * 16; )==""\n"
954	R"==(float s_tail = LOAD_DATA_1x16(&src[sg_idx]); )==""\n"
955	R"==(float d_tail = fma(s_tail - v_mean_tail[sg], sqrt_variance_tail[sg], )==""\n"
956	R"==(sv_tail[sg]); )==""\n"
957	R"==(#if FUSE_BN_RELU )==""\n"
958	R"==(#if FUSE_BN_ADD_RELU )==""\n"
959	R"==(float s_add_tail = LOAD_DATA_1x16(&src_add[sg_idx]); )==""\n"
960	R"==(d_tail += s_add_tail; )==""\n"
961	R"==(#endif )==""\n"
962	R"==(int ws_tail = isgreater(d_tail, 0.0f); )==""\n"
963	R"==(d_tail = select(0.0f, d_tail, ws_tail); )==""\n"
964	R"==(#if IS_TRAINING )==""\n"
965	R"==(STORE_CHAR_1x16(&ws[sg_idx], convert_char(ws_tail)); )==""\n"
966	R"==(#endif )==""\n"
967	R"==(#endif )==""\n"
968	R"==(#if WITH_RELU )==""\n"
969	R"==(d_tail = max(d_tail, 0.0f); )==""\n"
970	R"==(#endif )==""\n"
971	R"==(STORE_DATA_1x16(&dst[sg_idx], d_tail); )==""\n"
972	R"==(} )==""\n"
973	R"==(#endif )==""\n"
974	R"==(src += IC; )==""\n"
975	R"==(#if FUSE_BN_ADD_RELU )==""\n"
976	R"==(src_add += IC; )==""\n"
977	R"==(#endif )==""\n"
978	R"==(dst += IC; )==""\n"
979	R"==(#if FUSE_BN_RELU && IS_TRAINING )==""\n"
980	R"==(ws += IC; )==""\n"
981	R"==(#endif )==""\n"
982	R"==(} )==""\n"
983	R"==(} )==""\n"
984	R"==(#else )==""\n"
985	R"==(inline float8 read_src_block(__global DATA_T *src, int c, int sp) { )==""\n"
986	R"==(float8 blockS0 = 0.0f; )==""\n"
987	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
988	R"==(#if HAS_IC_TAIL )==""\n"
989	R"==(const bool is_last_ic_block = c + 16 > IC; )==""\n"
990	R"==(const bool is_last_sp_block = sp >= SP - VECT_SIZE; )==""\n"
991	R"==(#endif )==""\n"
992	R"==(#if HAS_SP_TAIL )==""\n"
993	R"==(if (sp == SP_TAIL) { )==""\n"
994	R"==(for (int k = 0; k < SP - SP_TAIL; ++k) )==""\n"
995	R"==(#if HAS_IC_TAIL )==""\n"
996	R"==(if (k == SP - SP_TAIL - 1 && is_last_ic_block) )==""\n"
997	R"==(blockS0[k] = simd_id < 8 )==""\n"
998	R"==(? CONVERT_FLOAT_T(src[k * IC_BLOCK_STRIDE + simd_id]) )==""\n"
999	R"==(: 0.0f; )==""\n"
1000	R"==(else )==""\n"
1001	R"==(#endif )==""\n"
1002	R"==(blockS0[k] = LOAD_DATA_1x16(&src[k * IC_BLOCK_STRIDE]); )==""\n"
1003	R"==(} else )==""\n"
1004	R"==(#endif )==""\n"
1005	R"==({ )==""\n"
1006	R"==(#if USE_NHWC )==""\n"
1007	R"==(#if IS_IC_EQ_8 )==""\n"
1008	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX_HALF(8, blockS0, src, 0); )==""\n"
1009	R"==(float8 t0 = intel_sub_group_shuffle_down(blockS0, blockS0, 8); )==""\n"
1010	R"==(for (int k = 0; k < 7; k += 2) )==""\n"
1011	R"==(blockS0[k + 1] = t0[k]; )==""\n"
1012	R"==(#elif HAS_IC_TAIL )==""\n"
1013	R"==(if (is_last_ic_block && is_last_sp_block) { )==""\n"
1014	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(7, blockS0, src, 0); )==""\n"
1015	R"==(blockS0[7] = simd_id < 8 ? CONVERT_FLOAT_T(src[7 * IC + simd_id]) )==""\n"
1016	R"==(: 0.0f; )==""\n"
1017	R"==(} else { )==""\n"
1018	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, blockS0, src, 0); )==""\n"
1019	R"==(} )==""\n"
1020	R"==(#else )==""\n"
1021	R"==(LOAD_DATA_Nx16_USING_LOOP_IDX(8, blockS0, src, 0); )==""\n"
1022	R"==(#endif )==""\n"
1023	R"==(#else )==""\n"
1024	R"==(blockS0 = LOAD_DATA_8x16(&src[0]); )==""\n"
1025	R"==(#endif )==""\n"
1026	R"==(} )==""\n"
1027	R"==(return blockS0; )==""\n"
1028	R"==(} )==""\n"
1029	R"==(KERNEL_ATTR )==""\n"
1030	R"==(__kernel void gen9_bnorm_fwd(__global DATA_T src, __global float mean, )==""\n"
1031	R"==(__global float variance, __global DATA_T dst, )==""\n"
1032	R"==(__global float scaleshift, __global float shift, __global char *ws, )==""\n"
1033	R"==(float eps, __global DATA_T *src_add) { )==""\n"
1034	R"==(const int n = GWS_GET_MB(); )==""\n"
1035	R"==(const int c = GWS_GET_IC(); )==""\n"
1036	R"==(const int sp = GWS_GET_SP() * VECT_SIZE; )==""\n"
1037	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
1038	R"==(#if HAS_IC_TAIL )==""\n"
1039	R"==(const bool is_last_ic_block = c + 16 > IC; )==""\n"
1040	R"==(const bool is_last_sp_block = sp >= SP - VECT_SIZE; )==""\n"
1041	R"==(#endif )==""\n"
1042	R"==(#if USE_NHWC )==""\n"
1043	R"==(const uint d_off = sp * IC + c; )==""\n"
1044	R"==(#else )==""\n"
1045	R"==(const uint d_off = (c & 15) + sp * 16 + (c & ~15) * SP + n * SP * IC; )==""\n"
1046	R"==(#endif )==""\n"
1047	R"==(src += d_off; )==""\n"
1048	R"==(dst += d_off; )==""\n"
1049	R"==(float8 blockS0 = read_src_block(src, c, sp); )==""\n"
1050	R"==(#if FUSE_BN_ADD_RELU )==""\n"
1051	R"==(src_add += d_off; )==""\n"
1052	R"==(float8 block_S0_Add = read_src_block(src_add, c, sp); )==""\n"
1053	R"==(#endif )==""\n"
1054	R"==(float8 blockD0; )==""\n"
1055	R"==(#if USE_SCALE == 1 )==""\n"
1056	R"==(float sm = MAYBE_LAST_IC_LOAD_FLOAT_1x16(scaleshift, c); )==""\n"
1057	R"==(#else )==""\n"
1058	R"==(float sm = 1.0f; )==""\n"
1059	R"==(#endif )==""\n"
1060	R"==(#if USE_SHIFT == 1 )==""\n"
1061	R"==(float sv = MAYBE_LAST_IC_LOAD_FLOAT_1x16(shift, c); )==""\n"
1062	R"==(#else )==""\n"
1063	R"==(float sv = 0.0f; )==""\n"
1064	R"==(#endif )==""\n"
1065	R"==(float v_mean, v_variance; )==""\n"
1066	R"==(#if HAS_IC_TAIL )==""\n"
1067	R"==(if (is_last_ic_block) { )==""\n"
1068	R"==(v_mean = simd_id < 8 ? mean[c + simd_id] : 0.0f; )==""\n"
1069	R"==(v_variance = simd_id < 8 ? variance[c + simd_id] : 0.0f; )==""\n"
1070	R"==(} else )==""\n"
1071	R"==(#endif )==""\n"
1072	R"==({ )==""\n"
1073	R"==(v_mean = LOAD_FLOAT_1x16(&mean[c]); )==""\n"
1074	R"==(v_variance = LOAD_FLOAT_1x16(&variance[c]); )==""\n"
1075	R"==(} )==""\n"
1076	R"==(float sqrt_variance = sm / sqrt(v_variance + eps); )==""\n"
1077	R"==(blockD0 = fma(blockS0 - (float8)v_mean, (float8)sqrt_variance, (float8)sv); )==""\n"
1078	R"==(#if FUSE_BN_RELU )==""\n"
1079	R"==(#if FUSE_BN_ADD_RELU )==""\n"
1080	R"==(blockD0 += block_S0_Add; )==""\n"
1081	R"==(#endif )==""\n"
1082	R"==(int8 blockWS0 = isgreater(blockD0, (float8)0.0f); )==""\n"
1083	R"==(blockD0 = select((float8)0.0f, blockD0, blockWS0); )==""\n"
1084	R"==(#if IS_TRAINING )==""\n"
1085	R"==(ws += d_off; )==""\n"
1086	R"==(#if HAS_SP_TAIL )==""\n"
1087	R"==(if (sp == SP_TAIL) { )==""\n"
1088	R"==(for (int k = 0; k < SP - SP_TAIL; ++k) { )==""\n"
1089	R"==(STORE_CHAR_1x16( )==""\n"
1090	R"==(&ws[k * IC_BLOCK_STRIDE], convert_char(blockWS0[k])); )==""\n"
1091	R"==(} )==""\n"
1092	R"==(} else )==""\n"
1093	R"==(#endif )==""\n"
1094	R"==({ )==""\n"
1095	R"==(#if USE_NHWC )==""\n"
1096	R"==(for (int k = 0; k < 8; ++k) )==""\n"
1097	R"==(STORE_CHAR_1x16( )==""\n"
1098	R"==(&ws[k * IC_BLOCK_STRIDE], convert_char(blockWS0[k])); )==""\n"
1099	R"==(#else )==""\n"
1100	R"==(STORE_CHAR_8x16(&ws[0], convert_char8(blockWS0)); )==""\n"
1101	R"==(#endif )==""\n"
1102	R"==(} )==""\n"
1103	R"==(#endif )==""\n"
1104	R"==(#endif )==""\n"
1105	R"==(#if WITH_RELU )==""\n"
1106	R"==(blockD0 = max(blockD0, (VECT_FLOAT_T)0.0f); )==""\n"
1107	R"==(#endif )==""\n"
1108	R"==(#if HAS_SP_TAIL )==""\n"
1109	R"==(if (sp == SP_TAIL) { )==""\n"
1110	R"==(for (int k = 0; k < SP - SP_TAIL; ++k) { )==""\n"
1111	R"==(#if HAS_IC_TAIL )==""\n"
1112	R"==(if (is_last_ic_block) { )==""\n"
1113	R"==(if (simd_id < 8) )==""\n"
1114	R"==(dst[k * IC_BLOCK_STRIDE + simd_id] )==""\n"
1115	R"==(= CONVERT_DATA_T(blockD0[k]); )==""\n"
1116	R"==(} else )==""\n"
1117	R"==(#endif )==""\n"
1118	R"==(STORE_DATA_1x16(&dst[k * IC_BLOCK_STRIDE], blockD0[k]); )==""\n"
1119	R"==(} )==""\n"
1120	R"==(} else )==""\n"
1121	R"==(#endif )==""\n"
1122	R"==({ )==""\n"
1123	R"==(#if USE_NHWC )==""\n"
1124	R"==(for (int k = 0; k < 8; ++k) )==""\n"
1125	R"==(#if HAS_IC_TAIL )==""\n"
1126	R"==(if (is_last_ic_block) { )==""\n"
1127	R"==(if (simd_id < 8) )==""\n"
1128	R"==(dst[k * IC_BLOCK_STRIDE + simd_id] )==""\n"
1129	R"==(= CONVERT_DATA_T(blockD0[k]); )==""\n"
1130	R"==(} else )==""\n"
1131	R"==(#endif )==""\n"
1132	R"==(STORE_DATA_1x16(&dst[k * IC_BLOCK_STRIDE], blockD0[k]); )==""\n"
1133	R"==(#else )==""\n"
1134	R"==(STORE_DATA_8x16(&dst[0], blockD0); )==""\n"
1135	R"==(#endif )==""\n"
1136	R"==(} )==""\n"
1137	R"==(} )==""\n"
1138	R"==(#endif )==""\n"
1139	R"==(#endif )==""\n"
1140	R"==(#if IS_BWD == 1 )==""\n"
1141	R"==(#define LOAD_DATA_Nx16_USING_LOOP(n, dest, src) \ )==""\n"
1142	R"==({ \ )==""\n"
1143	R"==(for (int k = 0; k < n; ++k) { \ )==""\n"
1144	R"==(dest[k] = LOAD_DATA_1x16(&src[k * IC_BLOCK_STRIDE]); \ )==""\n"
1145	R"==(} \ )==""\n"
1146	R"==(} )==""\n"
1147	R"==(#define LOAD_UINT_Nx16_USING_LOOP(n, dest, src) \ )==""\n"
1148	R"==({ \ )==""\n"
1149	R"==(for (int k = 0; k < n; ++k) { \ )==""\n"
1150	R"==(dest[k] = LOAD_UINT_1x16(&src[k * IC_BLOCK_STRIDE]); \ )==""\n"
1151	R"==(} \ )==""\n"
1152	R"==(} )==""\n"
1153	R"==(#define LOAD_CHAR_Nx16_USING_LOOP(n, dest, src) \ )==""\n"
1154	R"==({ \ )==""\n"
1155	R"==(for (int k = 0; k < n; ++k) { \ )==""\n"
1156	R"==(dest[k] = LOAD_CHAR_1x16(&src[k * IC_BLOCK_STRIDE]); \ )==""\n"
1157	R"==(} \ )==""\n"
1158	R"==(} )==""\n"
1159	R"==(#define LOAD_DATA_8x16_USING_LAYOUT(dest, src) \ )==""\n"
1160	R"==({ \ )==""\n"
1161	R"==(if (USE_NHWC) { \ )==""\n"
1162	R"==(LOAD_DATA_Nx16_USING_LOOP(8, dest, src); \ )==""\n"
1163	R"==(} else { \ )==""\n"
1164	R"==(dest = LOAD_DATA_8x16(src); \ )==""\n"
1165	R"==(} \ )==""\n"
1166	R"==(} )==""\n"
1167	R"==(#define LOAD_UINT_8x16_USING_LAYOUT(dest, src) \ )==""\n"
1168	R"==({ \ )==""\n"
1169	R"==(if (USE_NHWC) { \ )==""\n"
1170	R"==(LOAD_UINT_Nx16_USING_LOOP(8, dest, src); \ )==""\n"
1171	R"==(} else { \ )==""\n"
1172	R"==(dest = LOAD_UINT_8x16(src); \ )==""\n"
1173	R"==(} \ )==""\n"
1174	R"==(} )==""\n"
1175	R"==(#define LOAD_CHAR_8x16_USING_LAYOUT(dest, src) \ )==""\n"
1176	R"==({ \ )==""\n"
1177	R"==(if (USE_NHWC) { \ )==""\n"
1178	R"==(LOAD_CHAR_Nx16_USING_LOOP(8, dest, src); \ )==""\n"
1179	R"==(} else { \ )==""\n"
1180	R"==(dest = LOAD_CHAR_8x16(src); \ )==""\n"
1181	R"==(} \ )==""\n"
1182	R"==(} )==""\n"
1183	R"==(#define LOAD_DATA_Nx16_USING_LOOP_HALF(n, dest, src) \ )==""\n"
1184	R"==({ \ )==""\n"
1185	R"==(for (int k = 0; k < n; k += 2) { \ )==""\n"
1186	R"==(dest[k] = LOAD_DATA_1x16(&src[k * IC_BLOCK_STRIDE]); \ )==""\n"
1187	R"==(} \ )==""\n"
1188	R"==(} )==""\n"
1189	R"==(#if FUSED_ATOMICS_REDUCTION )==""\n"
1190	R"==(#if NHWC_OPTIMIZED )==""\n"
1191	R"==(#if VECT_SIZE > 1 )==""\n"
1192	R"==(#define GET_SCALAR_VAL(v, idx) v[idx / VECT_SIZE][idx % VECT_SIZE] )==""\n"
1193	R"==(#else )==""\n"
1194	R"==(#define GET_SCALAR_VAL(v, idx) v[idx] )==""\n"
1195	R"==(#endif )==""\n"
1196	R"==(#else )==""\n"
1197	R"==(#define GET_SCALAR_VAL(v, idx) v[idx] )==""\n"
1198	R"==(#endif )==""\n"
1199	R"==(void gen9_calc_fused_reduction(volatile __global atomic_float *diff_scale, )==""\n"
1200	R"==(volatile __global atomic_float *diff_shift, int dst_offset, )==""\n"
1201	R"==(#if NHWC_OPTIMIZED )==""\n"
1202	R"==(VECT_FLOAT_T diff_gamma, VECT_FLOAT_T diff_beta, )==""\n"
1203	R"==(#else )==""\n"
1204	R"==(float diff_gamma, float diff_beta, )==""\n"
1205	R"==(#endif )==""\n"
1206	R"==(float diff_gamma_tail, float diff_beta_tail, )==""\n"
1207	R"==(__local float local_gamma, __local float local_beta) { )==""\n"
1208	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
1209	R"==(const int group_size = GWS_LWS1_CALC * GWS_LWS2_CALC; )==""\n"
1210	R"==(const int sg_group_id = get_local_id(0) / 16; )==""\n"
1211	R"==(const int local_id = get_local_id(1); )==""\n"
1212	R"==(for (int sg = 0; sg < REDUCE_NUM_SGROUPS; ++sg) { )==""\n"
1213	R"==(const int slm_offset = CALC_SLM_LINE_SIZE * local_id )==""\n"
1214	R"==(+ REDUCE_NUM_SGROUPS * 16 * sg_group_id + sg * 16 + simd_id; )==""\n"
1215	R"==(#if HAS_IC_VECT_TAIL && NHWC_OPTIMIZED )==""\n"
1216	R"==(if (sg >= IC_VECT_SGROUPS) { )==""\n"
1217	R"==(local_gamma[slm_offset] = diff_gamma_tail[sg - IC_VECT_SGROUPS]; )==""\n"
1218	R"==(local_beta[slm_offset] = diff_beta_tail[sg - IC_VECT_SGROUPS]; )==""\n"
1219	R"==(} else )==""\n"
1220	R"==(#endif )==""\n"
1221	R"==({ )==""\n"
1222	R"==(local_gamma[slm_offset] = GET_SCALAR_VAL(diff_gamma, sg); )==""\n"
1223	R"==(local_beta[slm_offset] = GET_SCALAR_VAL(diff_beta, sg); )==""\n"
1224	R"==(} )==""\n"
1225	R"==(} )==""\n"
1226	R"==(barrier(CLK_LOCAL_MEM_FENCE); )==""\n"
1227	R"==(if (local_id == 0) { )==""\n"
1228	R"==(for (int sg = 0; sg < REDUCE_NUM_SGROUPS; ++sg) { )==""\n"
1229	R"==(float d_gamma = 0.f; )==""\n"
1230	R"==(float d_beta = 0.f; )==""\n"
1231	R"==(for (int gr_id = 0; gr_id < group_size; ++gr_id) { )==""\n"
1232	R"==(const int off_local = CALC_SLM_LINE_SIZE * gr_id )==""\n"
1233	R"==(+ REDUCE_NUM_SGROUPS * 16 * sg_group_id + sg * 16 )==""\n"
1234	R"==(+ simd_id; )==""\n"
1235	R"==(d_gamma += local_gamma[off_local]; )==""\n"
1236	R"==(d_beta += local_beta[off_local]; )==""\n"
1237	R"==(} )==""\n"
1238	R"==(const int offset = dst_offset + sg * 16 + simd_id; )==""\n"
1239	R"==(#if HAS_IC_TAIL )==""\n"
1240	R"==(if (offset < IC) )==""\n"
1241	R"==(#endif )==""\n"
1242	R"==({ )==""\n"
1243	R"==(atomic_add_global(&diff_scale[offset], d_gamma); )==""\n"
1244	R"==(#if DIFF_SHIFT == 1 )==""\n"
1245	R"==(atomic_add_global(&diff_shift[offset], d_beta); )==""\n"
1246	R"==(#else )==""\n"
1247	R"==(atomic_add_global( )==""\n"
1248	R"==(&diff_shift[IC + IC * REDUCE_STAT_NBLOCKS + offset], )==""\n"
1249	R"==(d_beta); )==""\n"
1250	R"==(#endif )==""\n"
1251	R"==(} )==""\n"
1252	R"==(} )==""\n"
1253	R"==(} )==""\n"
1254	R"==(return; )==""\n"
1255	R"==(} )==""\n"
1256	R"==(#endif )==""\n"
1257	R"==(#if NHWC_OPTIMIZED )==""\n"
1258	R"==(NAMED_KERNEL_ATTR(CALC) )==""\n"
1259	R"==(__kernel void gen9_calculate_stats(__global DATA_T src, __global float mean, )==""\n"
1260	R"==(__global DATA_T diff_dst, __global char ws, )==""\n"
1261	R"==(__global float temp_reduce, volatile __global atomic_float diff_scale, )==""\n"
1262	R"==(volatile __global atomic_float *diff_shift) { )==""\n"
1263	R"==(const int mb = GWS_GET_STAT_MB(); )==""\n"
1264	R"==(const int c = GWS_GET_STAT_IC(); )==""\n"
1265	R"==(const int sp_block_idx = GWS_GET_STAT_SP(); )==""\n"
1266	R"==(const int ic_block_offset = (c / 16) * IC_BLOCK; )==""\n"
1267	R"==(const int offset = ic_block_offset + sp_block_idx * STAT_SP_BLOCK * IC; )==""\n"
1268	R"==(mean += ic_block_offset; )==""\n"
1269	R"==(src += offset; )==""\n"
1270	R"==(diff_dst += offset; )==""\n"
1271	R"==(ws += offset; )==""\n"
1272	R"==(float v_mean[IC_BLOCK_SGROUPS]; )==""\n"
1273	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS; ++sg) { )==""\n"
1274	R"==(v_mean[sg] = as_float(intel_sub_group_block_read( )==""\n"
1275	R"==((const __global uint )(&mean[(sg 16)]))); )==""\n"
1276	R"==(} )==""\n"
1277	R"==(VECT_FLOAT_T diff_gamma[IC_BLOCK_SGROUPS / VECT_SIZE] = {0.0f}; )==""\n"
1278	R"==(VECT_FLOAT_T diff_beta[IC_BLOCK_SGROUPS / VECT_SIZE] = {0.0f}; )==""\n"
1279	R"==(#if HAS_IC_VECT_TAIL )==""\n"
1280	R"==(float diff_gamma_tail[IC_TAIL_SGROUPS] = {0.0f}; )==""\n"
1281	R"==(float diff_beta_tail[IC_TAIL_SGROUPS] = {0.0f}; )==""\n"
1282	R"==(#else )==""\n"
1283	R"==(float *diff_gamma_tail = NULL; )==""\n"
1284	R"==(float *diff_beta_tail = NULL; )==""\n"
1285	R"==(#endif )==""\n"
1286	R"==(#if HAS_STAT_SP_BLOCK_TAIL )==""\n"
1287	R"==(for (int sp = 0; sp < min(STAT_SP_BLOCK, SP - sp_block_idx * STAT_SP_BLOCK); )==""\n"
1288	R"==(++sp) { )==""\n"
1289	R"==(#else )==""\n"
1290	R"==(for (int sp = 0; sp < STAT_SP_BLOCK; ++sp) { )==""\n"
1291	R"==(#endif )==""\n"
1292	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS / VECT_SIZE; ++sg) { )==""\n"
1293	R"==(const int sg_idx = sg * 16 * VECT_SIZE; )==""\n"
1294	R"==(#if FUSE_BN_RELU )==""\n"
1295	R"==(VECT_CHAR_T ws_vect = LOAD_VECT_CHAR(&ws[sg_idx]); )==""\n"
1296	R"==(#endif )==""\n"
1297	R"==(VECT_FLOAT_T src_vect = LOAD_VECT_DATA(&src[sg_idx]); )==""\n"
1298	R"==(VECT_FLOAT_T dd_vect = LOAD_VECT_DATA(&diff_dst[sg_idx]); )==""\n"
1299	R"==(VECT_FLOAT_T v0; )==""\n"
1300	R"==(for (int vect = 0; vect < VECT_SIZE; ++vect) { )==""\n"
1301	R"==(int sg_idx = sg * VECT_SIZE + vect; )==""\n"
1302	R"==(#if VECT_SIZE > 1 )==""\n"
1303	R"==(v0[vect] = src_vect[vect] - v_mean[sg_idx]; )==""\n"
1304	R"==(#else )==""\n"
1305	R"==(v0 = src_vect - v_mean[sg_idx]; )==""\n"
1306	R"==(#endif )==""\n"
1307	R"==(} )==""\n"
1308	R"==(#if FUSE_BN_RELU )==""\n"
1309	R"==(dd_vect = select( )==""\n"
1310	R"==((VECT_FLOAT_T)0.0f, dd_vect, CONVERT_VECT_INT_T(ws_vect)); )==""\n"
1311	R"==(#endif )==""\n"
1312	R"==(diff_gamma[sg] = fma(v0, dd_vect, diff_gamma[sg]); )==""\n"
1313	R"==(diff_beta[sg] += dd_vect; )==""\n"
1314	R"==(} )==""\n"
1315	R"==(#if HAS_IC_VECT_TAIL )==""\n"
1316	R"==(for (int sg = 0; sg < IC_TAIL_SGROUPS; ++sg) { )==""\n"
1317	R"==(const int sg_idx = IC_VECT_SGROUPS + sg; )==""\n"
1318	R"==(#if FUSE_BN_RELU )==""\n"
1319	R"==(char ws_tail = LOAD_CHAR_1x16(&ws[sg_idx * 16]); )==""\n"
1320	R"==(#endif )==""\n"
1321	R"==(float src_tail = LOAD_DATA_1x16(&src[sg_idx * 16]); )==""\n"
1322	R"==(float dd_tail = LOAD_DATA_1x16(&diff_dst[sg_idx * 16]); )==""\n"
1323	R"==(float v0 = src_tail - v_mean[sg_idx]; )==""\n"
1324	R"==(#if FUSE_BN_RELU )==""\n"
1325	R"==(dd_tail = select(0.0f, dd_tail, convert_int(ws_tail)); )==""\n"
1326	R"==(#endif )==""\n"
1327	R"==(diff_gamma_tail[sg] = fma(v0, dd_tail, diff_gamma_tail[sg]); )==""\n"
1328	R"==(diff_beta_tail[sg] += dd_tail; )==""\n"
1329	R"==(} )==""\n"
1330	R"==(#endif )==""\n"
1331	R"==(src += IC; )==""\n"
1332	R"==(diff_dst += IC; )==""\n"
1333	R"==(#if FUSE_BN_RELU )==""\n"
1334	R"==(ws += IC; )==""\n"
1335	R"==(#endif )==""\n"
1336	R"==(} )==""\n"
1337	R"==(#if FUSED_ATOMICS_REDUCTION )==""\n"
1338	R"==(__local float local_gamma[2 * CALC_SLM_SIZE]; )==""\n"
1339	R"==(__local float *local_beta = local_gamma + CALC_SLM_SIZE; )==""\n"
1340	R"==(gen9_calc_fused_reduction(diff_scale, diff_shift, ic_block_offset, )==""\n"
1341	R"==(diff_gamma, diff_beta, diff_gamma_tail, diff_beta_tail, local_gamma, )==""\n"
1342	R"==(local_beta); )==""\n"
1343	R"==(#else )==""\n"
1344	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS; ++sg) { )==""\n"
1345	R"==(const int reduce_off = sp_block_idx * 16 )==""\n"
1346	R"==(+ REDUCE_STAT_NBLOCKS * 16 )==""\n"
1347	R"==(* (sg + (int)(c / 16) * (IC_BLOCK / 16)); )==""\n"
1348	R"==(const int diff_gamma_offset = IC16 + reduce_off; )==""\n"
1349	R"==(const int diff_beta_offset )==""\n"
1350	R"==(= 2 * IC16 + REDUCE_STAT_NBLOCKS * IC16 + reduce_off; )==""\n"
1351	R"==(#if HAS_IC_VECT_TAIL )==""\n"
1352	R"==(if (sg >= IC_VECT_SGROUPS) { )==""\n"
1353	R"==(STORE_FLOAT_1x16(&temp_reduce[diff_gamma_offset], )==""\n"
1354	R"==(diff_gamma_tail[sg - IC_VECT_SGROUPS]); )==""\n"
1355	R"==(STORE_FLOAT_1x16(&temp_reduce[diff_beta_offset], )==""\n"
1356	R"==(diff_beta_tail[sg - IC_VECT_SGROUPS]); )==""\n"
1357	R"==(} else )==""\n"
1358	R"==(#endif )==""\n"
1359	R"==({ )==""\n"
1360	R"==(#if VECT_SIZE > 1 )==""\n"
1361	R"==(STORE_FLOAT_1x16(&temp_reduce[diff_gamma_offset], )==""\n"
1362	R"==(diff_gamma[sg / VECT_SIZE][sg % VECT_SIZE]); )==""\n"
1363	R"==(STORE_FLOAT_1x16(&temp_reduce[diff_beta_offset], )==""\n"
1364	R"==(diff_beta[sg / VECT_SIZE][sg % VECT_SIZE]); )==""\n"
1365	R"==(#else )==""\n"
1366	R"==(STORE_FLOAT_1x16(&temp_reduce[diff_gamma_offset], diff_gamma[sg]); )==""\n"
1367	R"==(STORE_FLOAT_1x16(&temp_reduce[diff_beta_offset], diff_beta[sg]); )==""\n"
1368	R"==(#endif )==""\n"
1369	R"==(} )==""\n"
1370	R"==(} )==""\n"
1371	R"==(#endif )==""\n"
1372	R"==(} )==""\n"
1373	R"==(#else )==""\n"
1374	R"==(NAMED_KERNEL_ATTR(CALC) )==""\n"
1375	R"==(__kernel void gen9_calculate_stats(__global DATA_T src, __global float mean, )==""\n"
1376	R"==(__global DATA_T diff_dst, __global char ws, )==""\n"
1377	R"==(__global float temp_reduce, volatile __global atomic_float diff_scale, )==""\n"
1378	R"==(volatile __global atomic_float *diff_shift) { )==""\n"
1379	R"==(const int mb = GWS_GET_STAT_MB(); )==""\n"
1380	R"==(const int c = GWS_GET_STAT_IC(); )==""\n"
1381	R"==(const int sp_block_idx = GWS_GET_STAT_SP(); )==""\n"
1382	R"==(const int mb_sp_idx = mb * STAT_SP_NBLOCKS + sp_block_idx; )==""\n"
1383	R"==(const int group_c_offset = REDUCE_STAT_NBLOCKS * 16 * (int)(c / 16); )==""\n"
1384	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
1385	R"==(#if HAS_IC_TAIL )==""\n"
1386	R"==(const bool is_last_ic_block = c + 16 > IC; )==""\n"
1387	R"==(const bool is_last_sp_block = (sp_block_idx == STAT_SP_NBLOCKS - 1); )==""\n"
1388	R"==(#endif )==""\n"
1389	R"==(temp_reduce += group_c_offset; )==""\n"
1390	R"==(#if USE_NHWC )==""\n"
1391	R"==(const int offset = c + sp_block_idx * STAT_SP_BLOCK * IC; )==""\n"
1392	R"==(#else )==""\n"
1393	R"==(const int offset = (c & 15) + sp_block_idx * STAT_SP_BLOCK * 16 )==""\n"
1394	R"==(+ (c & ~15) * SP + mb * SP * IC; )==""\n"
1395	R"==(#endif )==""\n"
1396	R"==(src += offset; )==""\n"
1397	R"==(diff_dst += offset; )==""\n"
1398	R"==(ws += offset; )==""\n"
1399	R"==(float v_mean = MAYBE_LAST_IC_LOAD_FLOAT_1x16(mean, c); )==""\n"
1400	R"==(float8 diff_gamma = 0.0f; )==""\n"
1401	R"==(float8 diff_beta = 0.0f; )==""\n"
1402	R"==(#if HAS_STAT_SP_TAIL )==""\n"
1403	R"==(int sp; )==""\n"
1404	R"==(if (sp_block_idx == STAT_SP_TAIL) { )==""\n"
1405	R"==(sp = SP - STAT_SP_TAIL * STAT_SP_BLOCK; )==""\n"
1406	R"==(} else { )==""\n"
1407	R"==(sp = STAT_SP_BLOCK; )==""\n"
1408	R"==(} )==""\n"
1409	R"==(#else )==""\n"
1410	R"==(int sp = STAT_SP_BLOCK; )==""\n"
1411	R"==(#endif )==""\n"
1412	R"==(const int C_PARALLEL_FACTOR = 8; )==""\n"
1413	R"==(for (; sp > C_PARALLEL_FACTOR - 1; sp -= C_PARALLEL_FACTOR) { )==""\n"
1414	R"==(float8 src_data; )==""\n"
1415	R"==(float8 dd_data; )==""\n"
1416	R"==(#if FUSE_BN_RELU == 1 )==""\n"
1417	R"==(char8 ws_data; )==""\n"
1418	R"==(LOAD_CHAR_8x16_USING_LAYOUT(ws_data, ws); )==""\n"
1419	R"==(#endif )==""\n"
1420	R"==(#if IS_IC_EQ_8 )==""\n"
1421	R"==(LOAD_DATA_Nx16_USING_LOOP_HALF(8, src_data, src); )==""\n"
1422	R"==(LOAD_DATA_Nx16_USING_LOOP_HALF(8, dd_data, diff_dst); )==""\n"
1423	R"==(float8 t_src = intel_sub_group_shuffle_down(src_data, src_data, 8); )==""\n"
1424	R"==(float8 t_dd = intel_sub_group_shuffle_down(dd_data, dd_data, 8); )==""\n"
1425	R"==(for (int k = 0; k < 7; k += 2) { )==""\n"
1426	R"==(dd_data[k + 1] = t_dd[k]; )==""\n"
1427	R"==(src_data[k + 1] = t_src[k]; )==""\n"
1428	R"==(} )==""\n"
1429	R"==(#elif HAS_IC_TAIL )==""\n"
1430	R"==(const bool is_last_sp = sp - C_PARALLEL_FACTOR <= C_PARALLEL_FACTOR - 1; )==""\n"
1431	R"==(if (is_last_sp && is_last_ic_block && is_last_sp_block) { )==""\n"
1432	R"==(LOAD_DATA_Nx16_USING_LOOP(7, src_data, src); )==""\n"
1433	R"==(LOAD_DATA_Nx16_USING_LOOP(7, dd_data, diff_dst); )==""\n"
1434	R"==(dd_data[7] = simd_id < 8 )==""\n"
1435	R"==(? CONVERT_FLOAT_T(diff_dst[7 * IC_BLOCK_STRIDE + simd_id]) )==""\n"
1436	R"==(: 0.0f; )==""\n"
1437	R"==(src_data[7] = simd_id < 8 )==""\n"
1438	R"==(? CONVERT_FLOAT_T(src[7 * IC_BLOCK_STRIDE + simd_id]) )==""\n"
1439	R"==(: 0.0f; )==""\n"
1440	R"==(} else { )==""\n"
1441	R"==(LOAD_DATA_Nx16_USING_LOOP(8, src_data, src); )==""\n"
1442	R"==(LOAD_DATA_Nx16_USING_LOOP(8, dd_data, diff_dst); )==""\n"
1443	R"==(} )==""\n"
1444	R"==(#else )==""\n"
1445	R"==(LOAD_DATA_8x16_USING_LAYOUT(src_data, src); )==""\n"
1446	R"==(LOAD_DATA_8x16_USING_LAYOUT(dd_data, diff_dst); )==""\n"
1447	R"==(#endif )==""\n"
1448	R"==(src += C_PARALLEL_FACTOR * IC_BLOCK_STRIDE; )==""\n"
1449	R"==(diff_dst += C_PARALLEL_FACTOR * IC_BLOCK_STRIDE; )==""\n"
1450	R"==(#if FUSE_BN_RELU == 1 )==""\n"
1451	R"==(ws += C_PARALLEL_FACTOR * IC_BLOCK_STRIDE; )==""\n"
1452	R"==(#endif )==""\n"
1453	R"==(#if FUSE_BN_RELU == 1 )==""\n"
1454	R"==(const float8 C_ZERO = 0.0; )==""\n"
1455	R"==(dd_data = select(C_ZERO, dd_data, convert_int8(ws_data)); )==""\n"
1456	R"==(#endif )==""\n"
1457	R"==(const float8 v0 = src_data - v_mean; )==""\n"
1458	R"==(diff_gamma = fma(v0, dd_data, diff_gamma); )==""\n"
1459	R"==(diff_beta += dd_data; )==""\n"
1460	R"==(} )==""\n"
1461	R"==(#if HAS_STAT_SP_TAIL )==""\n"
1462	R"==(if (sp_block_idx == STAT_SP_TAIL) { )==""\n"
1463	R"==(sp = (SP - STAT_SP_TAIL * STAT_SP_BLOCK) % C_PARALLEL_FACTOR; )==""\n"
1464	R"==(while (sp-- >= 1) { )==""\n"
1465	R"==(#if FUSE_BN_RELU == 1 )==""\n"
1466	R"==(const char ws_data = LOAD_CHAR_1x16(&ws[0]); )==""\n"
1467	R"==(#else )==""\n"
1468	R"==(const char ws_data = 1; )==""\n"
1469	R"==(#endif )==""\n"
1470	R"==(#if HAS_IC_TAIL )==""\n"
1471	R"==(float src_data, dd_data; )==""\n"
1472	R"==(if (sp == 0 && is_last_ic_block) { )==""\n"
1473	R"==(src_data = simd_id < 8 ? CONVERT_FLOAT_T(src[simd_id]) : 0.0f; )==""\n"
1474	R"==(dd_data = simd_id < 8 ? CONVERT_FLOAT_T(diff_dst[simd_id]) )==""\n"
1475	R"==(: 0.0f; )==""\n"
1476	R"==(} else { )==""\n"
1477	R"==(src_data = LOAD_DATA_1x16(&src[0]); )==""\n"
1478	R"==(dd_data = LOAD_DATA_1x16(&diff_dst[0]); )==""\n"
1479	R"==(} )==""\n"
1480	R"==(#else )==""\n"
1481	R"==(const float src_data = LOAD_DATA_1x16(&src[0]); )==""\n"
1482	R"==(const float dd_data = LOAD_DATA_1x16(&diff_dst[0]); )==""\n"
1483	R"==(#endif )==""\n"
1484	R"==(src += IC_BLOCK_STRIDE; )==""\n"
1485	R"==(diff_dst += IC_BLOCK_STRIDE; )==""\n"
1486	R"==(#if FUSE_BN_RELU == 1 )==""\n"
1487	R"==(ws += IC_BLOCK_STRIDE; )==""\n"
1488	R"==(#endif )==""\n"
1489	R"==(if (ws_data != 0) { )==""\n"
1490	R"==(const float v0 = src_data - v_mean; )==""\n"
1491	R"==(const float diff_gamma_tmp = fma(v0, dd_data, diff_gamma[0]); )==""\n"
1492	R"==(diff_gamma[0] = diff_gamma_tmp; )==""\n"
1493	R"==(diff_beta[0] += dd_data; )==""\n"
1494	R"==(} )==""\n"
1495	R"==(} )==""\n"
1496	R"==(} )==""\n"
1497	R"==(#endif )==""\n"
1498	R"==(for (int i = 1; i < 8; i++) { )==""\n"
1499	R"==(diff_gamma[0] += diff_gamma[i]; )==""\n"
1500	R"==(diff_beta[0] += diff_beta[i]; )==""\n"
1501	R"==(} )==""\n"
1502	R"==(#if FUSED_ATOMICS_REDUCTION )==""\n"
1503	R"==(__local float local_gamma[2 * CALC_SLM_SIZE]; )==""\n"
1504	R"==(__local float *local_beta = local_gamma + CALC_SLM_SIZE; )==""\n"
1505	R"==(gen9_calc_fused_reduction(diff_scale, diff_shift, c, &diff_gamma, )==""\n"
1506	R"==(&diff_beta, NULL, NULL, local_gamma, local_beta); )==""\n"
1507	R"==(#else )==""\n"
1508	R"==(STORE_FLOAT_1x16(&temp_reduce[IC16 + mb_sp_idx * 16], diff_gamma[0]); )==""\n"
1509	R"==(STORE_FLOAT_1x16(&temp_reduce[2 * IC16 + REDUCE_STAT_NBLOCKS * IC16 )==""\n"
1510	R"==(+ mb_sp_idx * 16], )==""\n"
1511	R"==(diff_beta[0]); )==""\n"
1512	R"==(#endif )==""\n"
1513	R"==(} )==""\n"
1514	R"==(#endif )==""\n"
1515	R"==(NAMED_KERNEL_ATTR(REDUCE) )==""\n"
1516	R"==(__kernel void gen9_reduce_stats(__global float *temp_reduce, )==""\n"
1517	R"==(__global float diff_scale, __global float diff_shift, )==""\n"
1518	R"==(__global float *variance, float eps) { )==""\n"
1519	R"==(__local float local_gamma[16 * REDUCE_IC_SUB_GROUPS]; )==""\n"
1520	R"==(__local float local_beta[16 * REDUCE_IC_SUB_GROUPS]; )==""\n"
1521	R"==(const int ic_sub_group = get_global_id(0) / 16; )==""\n"
1522	R"==(const int group_c = get_global_id(1); )==""\n"
1523	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
1524	R"==(const int c = group_c * 16 + simd_id; )==""\n"
1525	R"==(float diff_gamma = 0.0f; )==""\n"
1526	R"==(float diff_beta = 0.0f; )==""\n"
1527	R"==(temp_reduce += IC16 + REDUCE_STAT_NBLOCKS * 16 * group_c )==""\n"
1528	R"==(+ REDUCE_STAT_NBLOCKS / REDUCE_IC_SUB_GROUPS * 16 * ic_sub_group )==""\n"
1529	R"==(+ simd_id; )==""\n"
1530	R"==(for (int i = 0; i < REDUCE_STAT_NBLOCKS / REDUCE_IC_SUB_GROUPS; i++) { )==""\n"
1531	R"==(diff_gamma += temp_reduce[i * 16]; )==""\n"
1532	R"==(} )==""\n"
1533	R"==(temp_reduce += IC16 + IC16 * REDUCE_STAT_NBLOCKS; )==""\n"
1534	R"==(for (int i = 0; i < REDUCE_STAT_NBLOCKS / REDUCE_IC_SUB_GROUPS; i++) { )==""\n"
1535	R"==(diff_beta += temp_reduce[i * 16]; )==""\n"
1536	R"==(} )==""\n"
1537	R"==(if (ic_sub_group > 0) { )==""\n"
1538	R"==(local_gamma[ic_sub_group * 16 + simd_id] = diff_gamma; )==""\n"
1539	R"==(local_beta[ic_sub_group * 16 + simd_id] = diff_beta; )==""\n"
1540	R"==(} )==""\n"
1541	R"==(barrier(CLK_LOCAL_MEM_FENCE); )==""\n"
1542	R"==(if (ic_sub_group == 0) { )==""\n"
1543	R"==(for (int i = 1; i < REDUCE_IC_SUB_GROUPS; i++) { )==""\n"
1544	R"==(diff_gamma += local_gamma[i * 16 + simd_id]; )==""\n"
1545	R"==(diff_beta += local_beta[i * 16 + simd_id]; )==""\n"
1546	R"==(} )==""\n"
1547	R"==(float sqrt_variance = 1.0f / sqrt(variance[c] + eps); )==""\n"
1548	R"==(#if HAS_IC_TAIL )==""\n"
1549	R"==(const bool is_last_ic_block = group_c * 16 + 16 > IC; )==""\n"
1550	R"==(if (!is_last_ic_block \|\| (is_last_ic_block && simd_id < 8)) )==""\n"
1551	R"==(#endif )==""\n"
1552	R"==({ )==""\n"
1553	R"==(diff_scale[c] = diff_gamma * sqrt_variance; )==""\n"
1554	R"==(#if DIFF_SHIFT == 1 )==""\n"
1555	R"==(diff_shift[c] = diff_beta; )==""\n"
1556	R"==(#else )==""\n"
1557	R"==(diff_shift[IC + IC * REDUCE_STAT_NBLOCKS + c] = diff_beta; )==""\n"
1558	R"==(#endif )==""\n"
1559	R"==(} )==""\n"
1560	R"==(} )==""\n"
1561	R"==(} )==""\n"
1562	R"==(NAMED_KERNEL_ATTR(AUX) )==""\n"
1563	R"==(__kernel void gen9_fused_reduce_final( )==""\n"
1564	R"==(__global float diff_scale, __global float variance, float eps) { )==""\n"
1565	R"==(const int c = GWS_GET_IC_AUX(); )==""\n"
1566	R"==(diff_scale[c] *= 1.0f / sqrt(variance[c] + eps); )==""\n"
1567	R"==(return; )==""\n"
1568	R"==(} )==""\n"
1569	R"==(#if NHWC_OPTIMIZED )==""\n"
1570	R"==(KERNEL_ATTR )==""\n"
1571	R"==(__kernel void gen9_bnorm_bwd(__global DATA_T src, __global float mean, )==""\n"
1572	R"==(__global float variance, __global DATA_T diff_dst, )==""\n"
1573	R"==(__global float scaleshift, __global char ws, )==""\n"
1574	R"==(__global DATA_T diff_src, __global float diff_scale, )==""\n"
1575	R"==(__global float diff_shift, float eps, __global DATA_T diff_src_add) { )==""\n"
1576	R"==(const int c = GWS_GET_IC(); )==""\n"
1577	R"==(const int ic_block_offset = (c / 16) * IC_BLOCK; )==""\n"
1578	R"==(variance += ic_block_offset; )==""\n"
1579	R"==(mean += ic_block_offset; )==""\n"
1580	R"==(diff_scale += ic_block_offset; )==""\n"
1581	R"==(diff_shift += ic_block_offset; )==""\n"
1582	R"==(scaleshift += ic_block_offset; )==""\n"
1583	R"==(VECT_FLOAT_T v_variance[IC_BLOCK_SGROUPS / VECT_SIZE], )==""\n"
1584	R"==(v_mean[IC_BLOCK_SGROUPS / VECT_SIZE], )==""\n"
1585	R"==(diff_gamma[IC_BLOCK_SGROUPS / VECT_SIZE], )==""\n"
1586	R"==(diff_beta[IC_BLOCK_SGROUPS / VECT_SIZE], )==""\n"
1587	R"==(sqrt_variance[IC_BLOCK_SGROUPS / VECT_SIZE], )==""\n"
1588	R"==(gamma[IC_BLOCK_SGROUPS / VECT_SIZE]; )==""\n"
1589	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS / VECT_SIZE; ++sg) { )==""\n"
1590	R"==(const int sg_idx = sg * 16 * VECT_SIZE; )==""\n"
1591	R"==(v_variance[sg] = LOAD_VECT_FLOAT(&variance[sg_idx]); )==""\n"
1592	R"==(#if CALCULATE_DIFF_STATS == 1 )==""\n"
1593	R"==(v_mean[sg] = LOAD_VECT_FLOAT(&mean[sg_idx]); )==""\n"
1594	R"==(diff_gamma[sg] = LOAD_VECT_FLOAT(&diff_scale[sg_idx]); )==""\n"
1595	R"==(#if DIFF_SHIFT == 1 )==""\n"
1596	R"==(diff_beta[sg] = LOAD_VECT_FLOAT(&diff_shift[sg_idx]); )==""\n"
1597	R"==(#else )==""\n"
1598	R"==(diff_beta[sg] = LOAD_VECT_FLOAT( )==""\n"
1599	R"==(&diff_shift[IC + REDUCE_STAT_NBLOCKS * IC + sg_idx]); )==""\n"
1600	R"==(#endif )==""\n"
1601	R"==(#endif )==""\n"
1602	R"==(#if USE_SCALE == 1 )==""\n"
1603	R"==(gamma[sg] = LOAD_VECT_FLOAT(&scaleshift[sg_idx]); )==""\n"
1604	R"==(#else )==""\n"
1605	R"==(gamma[sg] = (VECT_FLOAT_T)1.0f; )==""\n"
1606	R"==(#endif )==""\n"
1607	R"==(sqrt_variance[sg] )==""\n"
1608	R"==(= (VECT_FLOAT_T)1.0f / sqrt(v_variance[sg] + (VECT_FLOAT_T)eps); )==""\n"
1609	R"==(} )==""\n"
1610	R"==(#if HAS_IC_VECT_TAIL )==""\n"
1611	R"==(float v_variance_tail[IC_TAIL_SGROUPS], v_mean_tail[IC_TAIL_SGROUPS], )==""\n"
1612	R"==(diff_gamma_tail[IC_TAIL_SGROUPS], diff_beta_tail[IC_TAIL_SGROUPS], )==""\n"
1613	R"==(sqrt_variance_tail[IC_TAIL_SGROUPS], gamma_tail[IC_TAIL_SGROUPS]; )==""\n"
1614	R"==(for (int sg = 0; sg < IC_TAIL_SGROUPS; ++sg) { )==""\n"
1615	R"==(const int sg_idx = (IC_VECT_SGROUPS + sg) * 16; )==""\n"
1616	R"==(v_variance_tail[sg] = LOAD_FLOAT_1x16(&variance[sg_idx]); )==""\n"
1617	R"==(#if CALCULATE_DIFF_STATS == 1 )==""\n"
1618	R"==(v_mean_tail[sg] = LOAD_FLOAT_1x16(&mean[sg_idx]); )==""\n"
1619	R"==(diff_gamma_tail[sg] = LOAD_FLOAT_1x16(&diff_scale[sg_idx]); )==""\n"
1620	R"==(#if DIFF_SHIFT == 1 )==""\n"
1621	R"==(diff_beta_tail[sg] = LOAD_FLOAT_1x16(&diff_shift[sg_idx]); )==""\n"
1622	R"==(#else )==""\n"
1623	R"==(diff_beta_tail[sg] = LOAD_FLOAT_1x16( )==""\n"
1624	R"==(&diff_shift[IC + REDUCE_STAT_NBLOCKS * IC + sg_idx]); )==""\n"
1625	R"==(#endif )==""\n"
1626	R"==(#endif )==""\n"
1627	R"==(#if USE_SCALE == 1 )==""\n"
1628	R"==(gamma_tail[sg] = LOAD_FLOAT_1x16(&scaleshift[sg_idx]); )==""\n"
1629	R"==(#else )==""\n"
1630	R"==(gamma_tail[sg] = 1.0f; )==""\n"
1631	R"==(#endif )==""\n"
1632	R"==(sqrt_variance_tail[sg] = 1.0f / sqrt(v_variance_tail[sg] + eps); )==""\n"
1633	R"==(} )==""\n"
1634	R"==(#endif )==""\n"
1635	R"==(const int sp_block_idx = GWS_GET_SP(); )==""\n"
1636	R"==(const int offset = ic_block_offset + sp_block_idx * STAT_SP_BLOCK * IC; )==""\n"
1637	R"==(src += offset; )==""\n"
1638	R"==(diff_dst += offset; )==""\n"
1639	R"==(ws += offset; )==""\n"
1640	R"==(diff_src += offset; )==""\n"
1641	R"==(#if FUSE_BN_ADD_RELU )==""\n"
1642	R"==(diff_src_add += offset; )==""\n"
1643	R"==(#endif )==""\n"
1644	R"==(#if HAS_STAT_SP_BLOCK_TAIL )==""\n"
1645	R"==(for (int sp = 0; sp < min(STAT_SP_BLOCK, SP - sp_block_idx * STAT_SP_BLOCK); )==""\n"
1646	R"==(++sp) { )==""\n"
1647	R"==(#else )==""\n"
1648	R"==(for (int sp = 0; sp < STAT_SP_BLOCK; ++sp) { )==""\n"
1649	R"==(#endif )==""\n"
1650	R"==(for (int sg = 0; sg < IC_BLOCK_SGROUPS / VECT_SIZE; ++sg) { )==""\n"
1651	R"==(const int sg_idx = sg * 16 * VECT_SIZE; )==""\n"
1652	R"==(VECT_FLOAT_T src_vect = LOAD_VECT_DATA(&src[sg_idx]); )==""\n"
1653	R"==(VECT_FLOAT_T dd_vect = LOAD_VECT_DATA(&diff_dst[sg_idx]); )==""\n"
1654	R"==(#if FUSE_BN_RELU )==""\n"
1655	R"==(VECT_CHAR_T ws_vect = LOAD_VECT_CHAR(&ws[sg_idx]); )==""\n"
1656	R"==(dd_vect = select( )==""\n"
1657	R"==((VECT_FLOAT_T)0.0f, dd_vect, CONVERT_VECT_INT_T(ws_vect)); )==""\n"
1658	R"==(#if FUSE_BN_ADD_RELU )==""\n"
1659	R"==(STORE_VECT_DATA(&diff_src_add[sg_idx], dd_vect); )==""\n"
1660	R"==(#endif )==""\n"
1661	R"==(#endif )==""\n"
1662	R"==(#if CALCULATE_DIFF_STATS == 1 )==""\n"
1663	R"==(dd_vect -= (diff_beta[sg] )==""\n"
1664	R"==(+ (src_vect - v_mean[sg]) * diff_gamma[sg] )==""\n"
1665	R"==(* sqrt_variance[sg]) )==""\n"
1666	R"==(/ (MB * ID * IH * IW); )==""\n"
1667	R"==(#endif )==""\n"
1668	R"==(dd_vect = gamma[sg] sqrt_variance[sg]; )==""\n"
1669	R"==(STORE_VECT_DATA(&diff_src[sg_idx], dd_vect); )==""\n"
1670	R"==(} )==""\n"
1671	R"==(#if HAS_IC_VECT_TAIL )==""\n"
1672	R"==(for (int sg = 0; sg < IC_TAIL_SGROUPS; ++sg) { )==""\n"
1673	R"==(const int sg_idx = (IC_VECT_SGROUPS + sg) * 16; )==""\n"
1674	R"==(float src_tail = LOAD_DATA_1x16(&src[sg_idx]); )==""\n"
1675	R"==(float dd_tail = LOAD_DATA_1x16(&diff_dst[sg_idx]); )==""\n"
1676	R"==(#if FUSE_BN_RELU )==""\n"
1677	R"==(char ws_tail = LOAD_CHAR_1x16(&ws[sg_idx]); )==""\n"
1678	R"==(dd_tail = select(0.0f, dd_tail, convert_int(ws_tail)); )==""\n"
1679	R"==(#if FUSE_BN_ADD_RELU )==""\n"
1680	R"==(STORE_DATA_1x16(&diff_src_add[sg_idx], dd_tail); )==""\n"
1681	R"==(#endif )==""\n"
1682	R"==(#endif )==""\n"
1683	R"==(#if CALCULATE_DIFF_STATS == 1 )==""\n"
1684	R"==(dd_tail -= (diff_beta_tail[sg] )==""\n"
1685	R"==(+ (src_tail - v_mean_tail[sg]) )==""\n"
1686	R"==(* diff_gamma_tail[sg] )==""\n"
1687	R"==(* sqrt_variance_tail[sg]) )==""\n"
1688	R"==(/ (MB * ID * IH * IW); )==""\n"
1689	R"==(#endif )==""\n"
1690	R"==(dd_tail = gamma_tail[sg] sqrt_variance_tail[sg]; )==""\n"
1691	R"==(STORE_DATA_1x16(&diff_src[sg_idx], dd_tail); )==""\n"
1692	R"==(} )==""\n"
1693	R"==(#endif )==""\n"
1694	R"==(src += IC; )==""\n"
1695	R"==(diff_dst += IC; )==""\n"
1696	R"==(diff_src += IC; )==""\n"
1697	R"==(#if FUSE_BN_RELU )==""\n"
1698	R"==(#if FUSE_BN_ADD_RELU )==""\n"
1699	R"==(diff_src_add += IC; )==""\n"
1700	R"==(#endif )==""\n"
1701	R"==(ws += IC; )==""\n"
1702	R"==(#endif )==""\n"
1703	R"==(} )==""\n"
1704	R"==(} )==""\n"
1705	R"==(#else )==""\n"
1706	R"==(inline void write_8x16_block(__global DATA_T *ptr, int c, float8 val) { )==""\n"
1707	R"==(#if USE_NHWC )==""\n"
1708	R"==(#if HAS_IC_TAIL )==""\n"
1709	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
1710	R"==(const bool is_last_ic_block = c + 16 > IC; )==""\n"
1711	R"==(if (is_last_ic_block) { )==""\n"
1712	R"==(if (simd_id < 8) { )==""\n"
1713	R"==(for (int k = 0; k < 8; ++k) )==""\n"
1714	R"==(ptr[k * IC_BLOCK_STRIDE + simd_id] = CONVERT_DATA_T(val[k]); )==""\n"
1715	R"==(} )==""\n"
1716	R"==(} else )==""\n"
1717	R"==(#endif )==""\n"
1718	R"==(for (int k = 0; k < 8; ++k) )==""\n"
1719	R"==(STORE_DATA_1x16(&ptr[k * IC_BLOCK_STRIDE], val[k]); )==""\n"
1720	R"==(#else )==""\n"
1721	R"==(STORE_DATA_8x16(&ptr[0], val); )==""\n"
1722	R"==(#endif )==""\n"
1723	R"==(} )==""\n"
1724	R"==(inline void write_1x16_block(__global DATA_T *ptr, int c, float val) { )==""\n"
1725	R"==(#if HAS_IC_TAIL )==""\n"
1726	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
1727	R"==(const bool is_last_ic_block = c + 16 > IC; )==""\n"
1728	R"==(if (!is_last_ic_block) { )==""\n"
1729	R"==(STORE_DATA_1x16(&ptr[0], val); )==""\n"
1730	R"==(} else { )==""\n"
1731	R"==(if (simd_id < 8) { ptr[simd_id] = CONVERT_DATA_T(val); } )==""\n"
1732	R"==(} )==""\n"
1733	R"==(#else )==""\n"
1734	R"==(STORE_DATA_1x16(&ptr[0], val); )==""\n"
1735	R"==(#endif )==""\n"
1736	R"==(} )==""\n"
1737	R"==(KERNEL_ATTR )==""\n"
1738	R"==(__kernel void gen9_bnorm_bwd(__global DATA_T src, __global float mean, )==""\n"
1739	R"==(__global float variance, __global DATA_T diff_dst, )==""\n"
1740	R"==(__global float scaleshift, __global char ws, )==""\n"
1741	R"==(__global DATA_T diff_src, __global float diff_scale, )==""\n"
1742	R"==(__global float diff_shift, float eps, __global DATA_T diff_src_add) { )==""\n"
1743	R"==(const int c = GWS_GET_IC(); )==""\n"
1744	R"==(const int simd_id = get_sub_group_local_id(); )==""\n"
1745	R"==(#if HAS_IC_TAIL )==""\n"
1746	R"==(const bool is_last_ic_block = c + 16 > IC; )==""\n"
1747	R"==(#endif )==""\n"
1748	R"==(const float v_variance = MAYBE_LAST_IC_LOAD_FLOAT_1x16(variance, c); )==""\n"
1749	R"==(#if CALCULATE_DIFF_STATS == 1 )==""\n"
1750	R"==(const float v_mean = MAYBE_LAST_IC_LOAD_FLOAT_1x16(mean, c); )==""\n"
1751	R"==(const float diff_gamma = MAYBE_LAST_IC_LOAD_FLOAT_1x16(diff_scale, c); )==""\n"
1752	R"==(#if DIFF_SHIFT == 1 )==""\n"
1753	R"==(const float diff_beta = MAYBE_LAST_IC_LOAD_FLOAT_1x16(diff_shift, c); )==""\n"
1754	R"==(#else )==""\n"
1755	R"==(const float diff_beta = MAYBE_LAST_IC_LOAD_FLOAT_1x16( )==""\n"
1756	R"==(diff_shift, IC + REDUCE_STAT_NBLOCKS * IC + c); )==""\n"
1757	R"==(#endif )==""\n"
1758	R"==(#endif )==""\n"
1759	R"==(#if USE_SCALE == 1 )==""\n"
1760	R"==(const float gamma = MAYBE_LAST_IC_LOAD_FLOAT_1x16(scaleshift, c); )==""\n"
1761	R"==(#else )==""\n"
1762	R"==(const float gamma = 1; )==""\n"
1763	R"==(#endif )==""\n"
1764	R"==(const int sp_block_idx = GWS_GET_SP(); )==""\n"
1765	R"==(#if USE_NHWC )==""\n"
1766	R"==(const int offset = c + sp_block_idx * VECT_SIZE * IC; )==""\n"
1767	R"==(#else )==""\n"
1768	R"==(const int mb = GWS_GET_MB(); )==""\n"
1769	R"==(const int offset = (c & 15) + sp_block_idx * VECT_SIZE * 16 + (c & ~15) * SP )==""\n"
1770	R"==(+ mb * SP * IC; )==""\n"
1771	R"==(#endif )==""\n"
1772	R"==(#if HAS_IC_TAIL )==""\n"
1773	R"==(const bool is_last_sp_block = sp_block_idx == SP / VECT_SIZE - 1; )==""\n"
1774	R"==(#endif )==""\n"
1775	R"==(src += offset; )==""\n"
1776	R"==(diff_dst += offset; )==""\n"
1777	R"==(ws += offset; )==""\n"
1778	R"==(diff_src += offset; )==""\n"
1779	R"==(#if FUSE_BN_ADD_RELU )==""\n"
1780	R"==(diff_src_add += offset; )==""\n"
1781	R"==(#endif )==""\n"
1782	R"==(#if HAS_SP_TAIL )==""\n"
1783	R"==(int sp; )==""\n"
1784	R"==(if (sp_block_idx == SP_TAIL / VECT_SIZE) { )==""\n"
1785	R"==(sp = SP - SP_TAIL; )==""\n"
1786	R"==(} else { )==""\n"
1787	R"==(sp = VECT_SIZE; )==""\n"
1788	R"==(} )==""\n"
1789	R"==(#else )==""\n"
1790	R"==(int sp = VECT_SIZE; )==""\n"
1791	R"==(#endif )==""\n"
1792	R"==(const float sqrt_variance = 1.0f / sqrt(v_variance + eps); )==""\n"
1793	R"==(const int C_PARALLEL_FACTOR = 8; )==""\n"
1794	R"==(for (; sp > C_PARALLEL_FACTOR - 1; sp -= C_PARALLEL_FACTOR) { )==""\n"
1795	R"==(float8 src_data; )==""\n"
1796	R"==(float8 dd_data; )==""\n"
1797	R"==(#if FUSE_BN_RELU == 1 )==""\n"
1798	R"==(char8 ws_data; )==""\n"
1799	R"==(LOAD_CHAR_8x16_USING_LAYOUT(ws_data, ws); )==""\n"
1800	R"==(#endif )==""\n"
1801	R"==(#if IS_IC_EQ_8 )==""\n"
1802	R"==(LOAD_DATA_Nx16_USING_LOOP_HALF(8, src_data, src); )==""\n"
1803	R"==(LOAD_DATA_Nx16_USING_LOOP_HALF(8, dd_data, diff_dst); )==""\n"
1804	R"==(float8 t_dd = intel_sub_group_shuffle_down(dd_data, dd_data, 8); )==""\n"
1805	R"==(float8 t_src = intel_sub_group_shuffle_down(src_data, src_data, 8); )==""\n"
1806	R"==(for (int k = 0; k < 7; k += 2) { )==""\n"
1807	R"==(dd_data[k + 1] = t_dd[k]; )==""\n"
1808	R"==(src_data[k + 1] = t_src[k]; )==""\n"
1809	R"==(} )==""\n"
1810	R"==(#elif HAS_IC_TAIL && !HAS_SP_TAIL )==""\n"
1811	R"==(const bool is_last_sp = sp - C_PARALLEL_FACTOR <= C_PARALLEL_FACTOR - 1; )==""\n"
1812	R"==(if (is_last_sp && is_last_ic_block && is_last_sp_block) { )==""\n"
1813	R"==(LOAD_DATA_Nx16_USING_LOOP(7, src_data, src); )==""\n"
1814	R"==(LOAD_DATA_Nx16_USING_LOOP(7, dd_data, diff_dst); )==""\n"
1815	R"==(dd_data[7] = simd_id < 8 )==""\n"
1816	R"==(? CONVERT_FLOAT_T(diff_dst[7 * IC_BLOCK_STRIDE + simd_id]) )==""\n"
1817	R"==(: 0.0f; )==""\n"
1818	R"==(src_data[7] = simd_id < 8 )==""\n"
1819	R"==(? CONVERT_FLOAT_T(src[7 * IC_BLOCK_STRIDE + simd_id]) )==""\n"
1820	R"==(: 0.0f; )==""\n"
1821	R"==(} else { )==""\n"
1822	R"==(LOAD_DATA_Nx16_USING_LOOP(8, src_data, src); )==""\n"
1823	R"==(LOAD_DATA_Nx16_USING_LOOP(8, dd_data, diff_dst); )==""\n"
1824	R"==(} )==""\n"
1825	R"==(#else )==""\n"
1826	R"==(LOAD_DATA_8x16_USING_LAYOUT(dd_data, diff_dst); )==""\n"
1827	R"==(LOAD_DATA_8x16_USING_LAYOUT(src_data, src); )==""\n"
1828	R"==(#endif )==""\n"
1829	R"==(src += C_PARALLEL_FACTOR * IC_BLOCK_STRIDE; )==""\n"
1830	R"==(diff_dst += C_PARALLEL_FACTOR * IC_BLOCK_STRIDE; )==""\n"
1831	R"==(#if FUSE_BN_RELU == 1 )==""\n"
1832	R"==(ws += C_PARALLEL_FACTOR * IC_BLOCK_STRIDE; )==""\n"
1833	R"==(#endif )==""\n"
1834	R"==(#if FUSE_BN_RELU == 1 )==""\n"
1835	R"==(const float8 C_ZERO = 0.0; )==""\n"
1836	R"==(dd_data = select(C_ZERO, dd_data, convert_int8(ws_data)); )==""\n"
1837	R"==(#if FUSE_BN_ADD_RELU )==""\n"
1838	R"==(write_8x16_block(diff_src_add, c, dd_data); )==""\n"
1839	R"==(#endif )==""\n"
1840	R"==(#endif )==""\n"
1841	R"==(#if CALCULATE_DIFF_STATS == 1 )==""\n"
1842	R"==(dd_data -= (diff_beta )==""\n"
1843	R"==(+ (src_data - v_mean) * diff_gamma * sqrt_variance) )==""\n"
1844	R"==(/ (MB * ID * IH * IW); )==""\n"
1845	R"==(#endif )==""\n"
1846	R"==(dd_data = gamma sqrt_variance; )==""\n"
1847	R"==(write_8x16_block(diff_src, c, dd_data); )==""\n"
1848	R"==(diff_src += C_PARALLEL_FACTOR * IC_BLOCK_STRIDE; )==""\n"
1849	R"==(#if FUSE_BN_ADD_RELU )==""\n"
1850	R"==(diff_src_add += C_PARALLEL_FACTOR * IC_BLOCK_STRIDE; )==""\n"
1851	R"==(#endif )==""\n"
1852	R"==(} )==""\n"
1853	R"==(#if HAS_SP_TAIL )==""\n"
1854	R"==(if (sp_block_idx == SP_TAIL / VECT_SIZE) { )==""\n"
1855	R"==(sp = (SP - SP_TAIL) % C_PARALLEL_FACTOR; )==""\n"
1856	R"==(while (sp-- >= 1) { )==""\n"
1857	R"==(#if FUSE_BN_RELU == 1 )==""\n"
1858	R"==(const char ws_data = LOAD_CHAR_1x16(&ws[0]); )==""\n"
1859	R"==(#endif )==""\n"
1860	R"==(#if HAS_IC_TAIL )==""\n"
1861	R"==(float dd_data; )==""\n"
1862	R"==(if (sp == 0 && is_last_ic_block) )==""\n"
1863	R"==(dd_data = simd_id < 8 ? CONVERT_FLOAT_T(diff_dst[simd_id]) )==""\n"
1864	R"==(: 0.0f; )==""\n"
1865	R"==(else )==""\n"
1866	R"==(dd_data = LOAD_DATA_1x16(&diff_dst[0]); )==""\n"
1867	R"==(#if CALCULATE_DIFF_STATS == 1 )==""\n"
1868	R"==(float src_data; )==""\n"
1869	R"==(if (sp == 0 && is_last_ic_block) )==""\n"
1870	R"==(src_data = simd_id < 8 ? CONVERT_FLOAT_T(src[simd_id]) : 0.0f; )==""\n"
1871	R"==(else )==""\n"
1872	R"==(src_data = LOAD_DATA_1x16(&src[0]); )==""\n"
1873	R"==(#endif )==""\n"
1874	R"==(#else )==""\n"
1875	R"==(float dd_data = LOAD_DATA_1x16(&diff_dst[0]); )==""\n"
1876	R"==(#if CALCULATE_DIFF_STATS == 1 )==""\n"
1877	R"==(const float src_data = LOAD_DATA_1x16(&src[0]); )==""\n"
1878	R"==(#endif )==""\n"
1879	R"==(#endif )==""\n"
1880	R"==(src += IC_BLOCK_STRIDE; )==""\n"
1881	R"==(diff_dst += IC_BLOCK_STRIDE; )==""\n"
1882	R"==(#if FUSE_BN_RELU == 1 )==""\n"
1883	R"==(ws += IC_BLOCK_STRIDE; )==""\n"
1884	R"==(#endif )==""\n"
1885	R"==(#if FUSE_BN_RELU == 1 )==""\n"
1886	R"==(if (ws_data == 0) dd_data = 0; )==""\n"
1887	R"==(#if FUSE_BN_ADD_RELU )==""\n"
1888	R"==(write_1x16_block(diff_src_add, c, dd_data); )==""\n"
1889	R"==(#endif )==""\n"
1890	R"==(#endif )==""\n"
1891	R"==(#if CALCULATE_DIFF_STATS == 1 )==""\n"
1892	R"==(dd_data -= (diff_beta )==""\n"
1893	R"==(+ (src_data - v_mean) * diff_gamma )==""\n"
1894	R"==(* sqrt_variance) )==""\n"
1895	R"==(/ (MB * ID * IH * IW); )==""\n"
1896	R"==(#endif )==""\n"
1897	R"==(dd_data = gamma sqrt_variance; )==""\n"
1898	R"==(write_1x16_block(diff_src, c, dd_data); )==""\n"
1899	R"==(diff_src += IC_BLOCK_STRIDE; )==""\n"
1900	R"==(#if FUSE_BN_ADD_RELU )==""\n"
1901	R"==(diff_src_add += IC_BLOCK_STRIDE; )==""\n"
1902	R"==(#endif )==""\n"
1903	R"==(} )==""\n"
1904	R"==(} )==""\n"
1905	R"==(#endif )==""\n"
1906	R"==(} )==""\n"
1907	R"==(#endif )==""\n"
1908	R"==(#endif )==""\n"
1909	R"==()==";
1910	}
1911	}
1912	}
1913	}

Browse the source code of oneDNN/build/src/gpu/ocl/gen9_bnorm_kernel.cpp