1namespace dnnl {
2namespace impl {
3namespace gpu {
4namespace ocl {
5const char *gen9_reduction_kernel = R"==(/******************************************************************************* )==""\n"
6R"==(* Copyright 2021-2022 Intel Corporation )==""\n"
7R"==(* )==""\n"
8R"==(* Licensed under the Apache License, Version 2.0 (the "License"); )==""\n"
9R"==(* you may not use this file except in compliance with the License. )==""\n"
10R"==(* You may obtain a copy of the License at )==""\n"
11R"==(* )==""\n"
12R"==(* http: )==""\n"
13R"==(* )==""\n"
14R"==(* Unless required by applicable law or agreed to in writing, software )==""\n"
15R"==(* distributed under the License is distributed on an "AS IS" BASIS, )==""\n"
16R"==(* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. )==""\n"
17R"==(* See the License for the specific language governing permissions and )==""\n"
18R"==(* limitations under the License. )==""\n"
19R"==(*******************************************************************************/ )==""\n"
20R"==(#include "gpu/ocl/ocl_post_ops.h" )==""\n"
21R"==(#include "gpu/ocl/ocl_types.h" )==""\n"
22R"==(#if defined(IS_MAX) )==""\n"
23R"==(#define INIT_ACC TO_DEF_ACC_DATA_T(DATA_MIN) )==""\n"
24R"==(#elif defined(IS_MIN) )==""\n"
25R"==(#define INIT_ACC TO_DEF_ACC_DATA_T(DATA_MAX) )==""\n"
26R"==(#elif defined(IS_MUL) )==""\n"
27R"==(#define INIT_ACC TO_DEF_ACC_DATA_T(DATA_ONE) )==""\n"
28R"==(#else )==""\n"
29R"==(#define INIT_ACC TO_DEF_ACC_DATA_T(DATA_ZERO) )==""\n"
30R"==(#endif )==""\n"
31R"==(#if defined(IS_MAX) )==""\n"
32R"==(#if defined(SRC_DT_S8) || defined(SRC_DT_U8) )==""\n"
33R"==(#define ACCUMULATE(x, y) max(x, y) )==""\n"
34R"==(#else )==""\n"
35R"==(#define ACCUMULATE(x, y) fmax(x, y) )==""\n"
36R"==(#endif )==""\n"
37R"==(#elif defined(IS_MIN) )==""\n"
38R"==(#if defined(SRC_DT_S8) || defined(SRC_DT_U8) )==""\n"
39R"==(#define ACCUMULATE(x, y) min(x, y) )==""\n"
40R"==(#else )==""\n"
41R"==(#define ACCUMULATE(x, y) fmin(x, y) )==""\n"
42R"==(#endif )==""\n"
43R"==(#elif defined(IS_MEAN) || defined(IS_SUM) )==""\n"
44R"==(#define ACCUMULATE(x, y) (x + y) )==""\n"
45R"==(#elif defined(IS_MUL) )==""\n"
46R"==(#define ACCUMULATE(x, y) (x * y) )==""\n"
47R"==(#else )==""\n"
48R"==(#define ACCUMULATE(x, y) (x + pow(fabs(y), POWER)) )==""\n"
49R"==(#endif )==""\n"
50R"==(#if defined(IS_MAX) || defined(IS_MIN) || defined(IS_MUL) )==""\n"
51R"==(#define ACCUMULATE_AGAIN(x, y) ACCUMULATE(x, y) )==""\n"
52R"==(#else )==""\n"
53R"==(#define ACCUMULATE_AGAIN(x, y) (x + y) )==""\n"
54R"==(#endif )==""\n"
55R"==(#if defined(IS_MEAN) )==""\n"
56R"==(#define FINALIZE_REDUCTION(x) (x / REDUCTION_SIZE) )==""\n"
57R"==(#elif defined(IS_LP_MAX) )==""\n"
58R"==(#define FINALIZE_REDUCTION(x) rootn(fmax(x, EPS), POWER) )==""\n"
59R"==(#elif defined(IS_LP_SUM) )==""\n"
60R"==(#define FINALIZE_REDUCTION(x) rootn(x + EPS, POWER) )==""\n"
61R"==(#elif defined(IS_P_MAX) )==""\n"
62R"==(#define FINALIZE_REDUCTION(x) fmax(x, EPS) )==""\n"
63R"==(#elif defined(IS_P_SUM) )==""\n"
64R"==(#define FINALIZE_REDUCTION(x) (x + EPS) )==""\n"
65R"==(#else )==""\n"
66R"==(#define FINALIZE_REDUCTION(x) (x) )==""\n"
67R"==(#endif )==""\n"
68R"==(#if WITH_SUM )==""\n"
69R"==(#define INIT_SUM(sum_data) DST_TO_REF(sum_data) )==""\n"
70R"==(#else )==""\n"
71R"==(#define INIT_SUM(sum_data) 0.0f )==""\n"
72R"==(#endif )==""\n"
73R"==(#if defined(IS_MAX) )==""\n"
74R"==(#define SUB_GROUP_REDUCE(x, c_block) sub_group_reduce_max(x) )==""\n"
75R"==(#elif defined(IS_MIN) )==""\n"
76R"==(#define SUB_GROUP_REDUCE(x, c_block) sub_group_reduce_min(x) )==""\n"
77R"==(#elif defined(IS_MUL) )==""\n"
78R"==(#define SUB_GROUP_REDUCE(x, c_block) \ )==""\n"
79R"==(({ \ )==""\n"
80R"==(int cid_end \ )==""\n"
81R"==(= (INITIAL_C % SUB_GROUP_SIZE == 0 ? SUB_GROUP_SIZE \ )==""\n"
82R"==(: (INITIAL_C - c_block)); \ )==""\n"
83R"==(DEF_ACC_DATA_T sub_group_acc = 1.0; \ )==""\n"
84R"==(for (int channel_id = 0; channel_id < cid_end; channel_id++) { \ )==""\n"
85R"==(sub_group_acc *= intel_sub_group_shuffle(c_acc, channel_id); \ )==""\n"
86R"==(} \ )==""\n"
87R"==(sub_group_acc; \ )==""\n"
88R"==(}) )==""\n"
89R"==(#else )==""\n"
90R"==(#define SUB_GROUP_REDUCE(x, c_block) sub_group_reduce_add(x) )==""\n"
91R"==(#endif )==""\n"
92R"==(#if INITIAL_C_CHUNKS == 1 )==""\n"
93R"==(#define C_BLOCK_READ BLOCK_READ )==""\n"
94R"==(#define AS_C_BLOCK_DATA_T AS_DATA_T )==""\n"
95R"==(#define CONVERT_C_BLOCK_T TO_DEF_ACC_DATA_T )==""\n"
96R"==(#define C_BLOCK_T DEF_ACC_DATA_T )==""\n"
97R"==(#elif INITIAL_C_CHUNKS == 2 )==""\n"
98R"==(#define C_BLOCK_READ BLOCK_READ2 )==""\n"
99R"==(#define AS_C_BLOCK_DATA_T AS_DATA2_T )==""\n"
100R"==(#define CONVERT_C_BLOCK_T TO_DEF_ACC_DATA2_T )==""\n"
101R"==(#define C_BLOCK_T DEF_ACC_DATA2_T )==""\n"
102R"==(#elif INITIAL_C_CHUNKS == 4 )==""\n"
103R"==(#define C_BLOCK_READ BLOCK_READ4 )==""\n"
104R"==(#define AS_C_BLOCK_DATA_T AS_DATA4_T )==""\n"
105R"==(#define CONVERT_C_BLOCK_T TO_DEF_ACC_DATA4_T )==""\n"
106R"==(#define C_BLOCK_T DEF_ACC_DATA4_T )==""\n"
107R"==(#elif INITIAL_C_CHUNKS == 8 )==""\n"
108R"==(#define C_BLOCK_READ BLOCK_READ8 )==""\n"
109R"==(#define AS_C_BLOCK_DATA_T AS_DATA8_T )==""\n"
110R"==(#define CONVERT_C_BLOCK_T TO_DEF_ACC_DATA8_T )==""\n"
111R"==(#define C_BLOCK_T DEF_ACC_DATA8_T )==""\n"
112R"==(#endif )==""\n"
113R"==(#define ROUND_DOWN(a, b) ((a) - ((a) % (b))) )==""\n"
114R"==(#undef ROUND_UP )==""\n"
115R"==(#define ROUND_UP(a, b) ROUND_DOWN((a + b - 1), (b)) )==""\n"
116R"==(#define INITIAL_SRC_OFFSET(n, c, hwd) \ )==""\n"
117R"==((((n) / N_BLOCK_SIZE) * INITIAL_HWD_DIM * N_BLOCK_SIZE * ROUND_UP(INITIAL_C, C_BLOCK_SIZE) + \ )==""\n"
118R"==(((c) / C_BLOCK_SIZE) * INITIAL_HWD_DIM * N_BLOCK_SIZE * C_BLOCK_SIZE + \ )==""\n"
119R"==((hwd) * N_BLOCK_SIZE * C_BLOCK_SIZE + \ )==""\n"
120R"==(((n) % N_BLOCK_SIZE) * C_BLOCK_SIZE + \ )==""\n"
121R"==(((c) % C_BLOCK_SIZE)) )==""\n"
122R"==(#define INITIAL_DST_OFFSET(n, c, hwd) \ )==""\n"
123R"==(((n / N_BLOCK_SIZE) * FINAL_HWD_DIM * N_BLOCK_SIZE * ROUND_UP(FINAL_C_DIM, C_BLOCK_SIZE) + \ )==""\n"
124R"==(((c) / C_BLOCK_SIZE) * FINAL_HWD_DIM * N_BLOCK_SIZE * C_BLOCK_SIZE + \ )==""\n"
125R"==((hwd) * N_BLOCK_SIZE * C_BLOCK_SIZE + \ )==""\n"
126R"==(((n) % N_BLOCK_SIZE) * C_BLOCK_SIZE + \ )==""\n"
127R"==(((c) % C_BLOCK_SIZE)) )==""\n"
128R"==(#define FINAL_SRC_OFFSET(n, c, hwd) INITIAL_DST_OFFSET(n, c, hwd) )==""\n"
129R"==(#define FINAL_DST_OFFSET(n, c, hwd) \ )==""\n"
130R"==(((n) / N_BLOCK_SIZE) * (FINAL_HWD_DIM / FINAL_HWD_CHUNK_SIZE) * N_BLOCK_SIZE * ROUND_UP(FINAL_C_DIM / FINAL_C_CHUNK_SIZE, C_BLOCK_SIZE) + \ )==""\n"
131R"==(((c) / C_BLOCK_SIZE) * (FINAL_HWD_DIM / FINAL_HWD_CHUNK_SIZE) * N_BLOCK_SIZE * C_BLOCK_SIZE + \ )==""\n"
132R"==((hwd) * N_BLOCK_SIZE * C_BLOCK_SIZE + \ )==""\n"
133R"==(((n) % N_BLOCK_SIZE) * C_BLOCK_SIZE + \ )==""\n"
134R"==(((c) % C_BLOCK_SIZE) )==""\n"
135R"==(#if WITH_POST_OP )==""\n"
136R"==(#define APPLY_POST_OPS(sum_data, data, n_idx, c_idx, hwd_idx) \ )==""\n"
137R"==({ \ )==""\n"
138R"==(float sum_init_val = INIT_SUM(sum_data); \ )==""\n"
139R"==(const int D = hwd_start / (DST_H_DIM * DST_W_DIM); \ )==""\n"
140R"==(const int H = (hwd_start % (DST_H_DIM * DST_W_DIM)) / DST_W_DIM; \ )==""\n"
141R"==(const int W = hwd_start % DST_W_DIM; \ )==""\n"
142R"==(APPLY_POST_OPS_SERIAL(data, float, sum_init_val, float, n_idx, 1, \ )==""\n"
143R"==(c_idx, 1, D, 1, H, 1, W, 1, 0, 1); \ )==""\n"
144R"==(} )==""\n"
145R"==(#else )==""\n"
146R"==(#define APPLY_POST_OPS(sum_data, data, n_idx, c_idx, hwd_idx) \ )==""\n"
147R"==({} )==""\n"
148R"==(#endif )==""\n"
149R"==(#define WRITE_FINAL_RESULT(dst_elem, acc, n_start, c_start, hwd_start) \ )==""\n"
150R"==({ \ )==""\n"
151R"==(float acc_float = FINALIZE_REDUCTION(convert_float(acc)); \ )==""\n"
152R"==(APPLY_POST_OPS(dst_elem, acc_float, n_start, c_start, hwd_start); \ )==""\n"
153R"==(dst_elem = TO_DST(acc_float); \ )==""\n"
154R"==(} )==""\n"
155R"==(#if SKIP_FINAL_PHASE )==""\n"
156R"==(#define WRITE_INITIAL_RESULT WRITE_FINAL_RESULT )==""\n"
157R"==(#define INITIAL_DST_DTYPE DST_DATA_T )==""\n"
158R"==(#else )==""\n"
159R"==(#define WRITE_INITIAL_RESULT(dst_elem, data, n_start, c_start, hwd_start) \ )==""\n"
160R"==({ dst_elem = data; } )==""\n"
161R"==(#define INITIAL_DST_DTYPE DEF_ACC_DATA_T )==""\n"
162R"==(#endif )==""\n"
163R"==(NAMED_KERNEL_ATTR(INITIAL) )==""\n"
164R"==(__kernel void gen9_initial_reduce(__global SRC_DATA_T *src, )==""\n"
165R"==(__global INITIAL_DST_DTYPE *dst )==""\n"
166R"==(#if SKIP_FINAL_PHASE )==""\n"
167R"==(POST_OP_ARGS )==""\n"
168R"==(#endif )==""\n"
169R"==() { )==""\n"
170R"==(const int n_chunk_idx = GWS_GET_INITIAL_N(); )==""\n"
171R"==(const int c = GWS_GET_INITIAL_C(); )==""\n"
172R"==(const int c_block_idx = c / C_BLOCK_SIZE; )==""\n"
173R"==(const int hwd_chunk_idx = GWS_GET_INITIAL_HWD_CHUNK_ID(); )==""\n"
174R"==(const int hwd_start = hwd_chunk_idx * INITIAL_HWD_CHUNK_SIZE; )==""\n"
175R"==(const int current_hwd_chunk = min(INITIAL_HWD_CHUNK_SIZE, )==""\n"
176R"==(INITIAL_HWD_DIM - hwd_chunk_idx * INITIAL_HWD_CHUNK_SIZE); )==""\n"
177R"==(const int aligned_hwd_chunk = ROUND_DOWN(current_hwd_chunk, VECT_DT_N); )==""\n"
178R"==(const int n_start = n_chunk_idx * INITIAL_N_CHUNK_SIZE; )==""\n"
179R"==(const int n_end = min(n_start + INITIAL_N_CHUNK_SIZE, INITIAL_N); )==""\n"
180R"==(#if SKIP_FINAL_PHASE )==""\n"
181R"==(for (int n_idx = n_start; n_idx < n_start + INITIAL_N_CHUNK_SIZE; n_idx++) { )==""\n"
182R"==(for (int c_idx = c; c_idx < c + INITIAL_C_CHUNKS * SUB_GROUP_SIZE; )==""\n"
183R"==(c_idx++) { )==""\n"
184R"==(if (n_idx >= DST_N && n_idx < DST_N_PADDED )==""\n"
185R"==(|| c_idx >= DST_C && c_idx < DST_C_PADDED) { )==""\n"
186R"==(for (int hwd_idx = hwd_start; )==""\n"
187R"==(#if IS_HWD_REDUCED )==""\n"
188R"==(hwd_idx < hwd_start + FINAL_HWD_CHUNK_SIZE; )==""\n"
189R"==(#else )==""\n"
190R"==(hwd_idx < hwd_start + INITIAL_HWD_CHUNK_SIZE; )==""\n"
191R"==(#endif )==""\n"
192R"==(hwd_idx++) { )==""\n"
193R"==(int n = (!IS_C_REDUCED && IS_N_REDUCED && NDIMS == 3 )==""\n"
194R"==(&& DST_N_PADDED == 1) )==""\n"
195R"==(? 0 )==""\n"
196R"==(: n_idx; )==""\n"
197R"==(const int dst_off = FINAL_DST_OFFSET(n, c_idx, hwd_idx); )==""\n"
198R"==(dst[dst_off] = TO_DST(0.0f); )==""\n"
199R"==(} )==""\n"
200R"==(} )==""\n"
201R"==(} )==""\n"
202R"==(} )==""\n"
203R"==(#endif )==""\n"
204R"==(int channel_id = c + get_sub_group_local_id(); )==""\n"
205R"==(if (channel_id >= INITIAL_C || n_start >= INITIAL_N) { return; } )==""\n"
206R"==(VECT_DEF_ACC_DATA_T vector_acc = INIT_ACC; )==""\n"
207R"==(for (int n = n_start; n < n_end; n++) { )==""\n"
208R"==(for (int hwd_id = 0; hwd_id < aligned_hwd_chunk; hwd_id += VECT_DT_N) { )==""\n"
209R"==(for (int c_chunk = 0; c_chunk < INITIAL_C_CHUNKS; c_chunk++) { )==""\n"
210R"==(const int off = INITIAL_SRC_OFFSET(n, c, )==""\n"
211R"==(hwd_start + hwd_id )==""\n"
212R"==(+ c_chunk * VECT_DT_N / INITIAL_C_CHUNKS); )==""\n"
213R"==(VECT_DEF_ACC_DATA_T data )==""\n"
214R"==(= AS_VECT_DEF_ACC_DATA_T(AS_VECT_DATA_T(VECT_BLOCK_READ( )==""\n"
215R"==((const __global BLOCK_DATA_T *)&src[off]))); )==""\n"
216R"==(vector_acc = ACCUMULATE(vector_acc, data); )==""\n"
217R"==(} )==""\n"
218R"==(} )==""\n"
219R"==(for (int hwd_id = aligned_hwd_chunk; hwd_id < current_hwd_chunk; )==""\n"
220R"==(hwd_id++) { )==""\n"
221R"==(const int off = INITIAL_SRC_OFFSET(n, c, hwd_start + hwd_id); )==""\n"
222R"==(C_BLOCK_T data = CONVERT_C_BLOCK_T(AS_C_BLOCK_DATA_T( )==""\n"
223R"==(C_BLOCK_READ((const __global BLOCK_DATA_T *)&src[off]))); )==""\n"
224R"==(#if VECT_DT_N == 1 )==""\n"
225R"==(vector_acc = ACCUMULATE(vector_acc, data); )==""\n"
226R"==(#else )==""\n"
227R"==(#if INITIAL_C_CHUNKS == 1 )==""\n"
228R"==(vector_acc[0] = ACCUMULATE(vector_acc[0], data); )==""\n"
229R"==(#else )==""\n"
230R"==(for (int i = 0; i < INITIAL_C_CHUNKS; i++) { )==""\n"
231R"==(vector_acc[i] = ACCUMULATE(vector_acc[i], data[i]); )==""\n"
232R"==(} )==""\n"
233R"==(#endif )==""\n"
234R"==(#endif )==""\n"
235R"==(} )==""\n"
236R"==(} )==""\n"
237R"==(#if VECT_DT_N == 1 )==""\n"
238R"==(VECT_DEF_ACC_DATA_T acc = vector_acc; )==""\n"
239R"==(#else )==""\n"
240R"==(const int elems_to_accumulate = aligned_hwd_chunk > 0 ? VECT_DT_N : 1; )==""\n"
241R"==(#if INITIAL_C_CHUNKS == 1 )==""\n"
242R"==(DEF_ACC_DATA_T acc = INIT_ACC; )==""\n"
243R"==(for (int i = 0; i < elems_to_accumulate; i++) { )==""\n"
244R"==(acc = ACCUMULATE_AGAIN(acc, vector_acc[i]); )==""\n"
245R"==(} )==""\n"
246R"==(#else )==""\n"
247R"==(C_BLOCK_T acc = INIT_ACC; )==""\n"
248R"==(for (int i = 0; i < elems_to_accumulate; i += INITIAL_C_CHUNKS) { )==""\n"
249R"==(unroll_for(int j = 0; j < INITIAL_C_CHUNKS; j++) { )==""\n"
250R"==(acc[j] = ACCUMULATE_AGAIN(acc[j], vector_acc[i + j]); )==""\n"
251R"==(} )==""\n"
252R"==(} )==""\n"
253R"==(#endif )==""\n"
254R"==(#endif )==""\n"
255R"==(const int local_id = get_sub_group_local_id(); )==""\n"
256R"==(#if IS_C_REDUCED )==""\n"
257R"==(#if INITIAL_C_CHUNKS == 1 )==""\n"
258R"==(DEF_ACC_DATA_T c_acc = acc; )==""\n"
259R"==(#else )==""\n"
260R"==(DEF_ACC_DATA_T c_acc = acc[0]; )==""\n"
261R"==(for (int i = 1; i < INITIAL_C_CHUNKS; i++) { )==""\n"
262R"==(c_acc += acc[i]; )==""\n"
263R"==(} )==""\n"
264R"==(#endif )==""\n"
265R"==(const int dst_off )==""\n"
266R"==(= INITIAL_DST_OFFSET(n_chunk_idx, c_block_idx, hwd_chunk_idx); )==""\n"
267R"==(c_acc = SUB_GROUP_REDUCE(c_acc, c); )==""\n"
268R"==(if (local_id == 0) { )==""\n"
269R"==(WRITE_INITIAL_RESULT(dst[dst_off], c_acc, n_start, c, hwd_start); )==""\n"
270R"==(} )==""\n"
271R"==(#else )==""\n"
272R"==(const int dst_c = c + local_id; )==""\n"
273R"==(#if INITIAL_C_CHUNKS == 1 )==""\n"
274R"==(WRITE_INITIAL_RESULT( )==""\n"
275R"==(dst[INITIAL_DST_OFFSET(n_chunk_idx, dst_c, hwd_chunk_idx)], acc, )==""\n"
276R"==(n_start, dst_c, hwd_start); )==""\n"
277R"==(#else )==""\n"
278R"==(for (int i = 0; i < INITIAL_C_CHUNKS; i++) { )==""\n"
279R"==(int c_off = i * SUB_GROUP_SIZE; )==""\n"
280R"==(WRITE_INITIAL_RESULT(dst[INITIAL_DST_OFFSET(n_chunk_idx, dst_c + c_off, )==""\n"
281R"==(hwd_chunk_idx)], )==""\n"
282R"==(acc[i], n_start, c, hwd_start); )==""\n"
283R"==(} )==""\n"
284R"==(#endif )==""\n"
285R"==(#endif )==""\n"
286R"==(} )==""\n"
287R"==(#if !SKIP_FINAL_PHASE )==""\n"
288R"==(NAMED_KERNEL_ATTR(FINAL) )==""\n"
289R"==(__kernel void gen9_final_reduce( )==""\n"
290R"==(__global DEF_ACC_DATA_T *src, __global DST_DATA_T *dst POST_OP_ARGS) { )==""\n"
291R"==(const int n_start = GWS_GET_FINAL_N() * FINAL_N_CHUNK_SIZE; )==""\n"
292R"==(const int c_start = GWS_GET_FINAL_C() * FINAL_C_CHUNK_SIZE; )==""\n"
293R"==(const int hwd_start = GWS_GET_FINAL_HWD() * FINAL_HWD_CHUNK_SIZE; )==""\n"
294R"==(DEF_ACC_DATA_T acc = INIT_ACC; )==""\n"
295R"==(const int max_n = max(DST_N_PADDED, FINAL_N_DIM); )==""\n"
296R"==(const int max_c = max(DST_C_PADDED, FINAL_C_DIM); )==""\n"
297R"==(const int n_end = min(max_n, n_start + FINAL_N_CHUNK_SIZE); )==""\n"
298R"==(const int c_end = min(max_c, c_start + FINAL_C_CHUNK_SIZE); )==""\n"
299R"==(const int hwd_end = min(FINAL_HWD_DIM, hwd_start + FINAL_HWD_CHUNK_SIZE); )==""\n"
300R"==(for (int n = n_start; n < n_end; n++) { )==""\n"
301R"==(for (int c = c_start; c < c_end; c++) { )==""\n"
302R"==(for (int hwd = hwd_start; hwd < hwd_end; hwd++) { )==""\n"
303R"==(if ((n >= DST_N && n < DST_N_PADDED) )==""\n"
304R"==(|| (c >= DST_C && c < DST_C_PADDED)) { )==""\n"
305R"==(#if NDIMS == 2 && DST_N_PADDED == 1 )==""\n"
306R"==(const int dst_off = FINAL_DST_OFFSET(0, c, hwd); )==""\n"
307R"==(#elif NDIMS >= 3 && DST_N_PADDED == 1 && IS_HWD_REDUCED )==""\n"
308R"==(const int dst_off = FINAL_DST_OFFSET(0, c, 0); )==""\n"
309R"==(#elif IS_HWD_REDUCED )==""\n"
310R"==(const int dst_off = FINAL_DST_OFFSET(n, c, 0); )==""\n"
311R"==(#else )==""\n"
312R"==(const int dst_off = FINAL_DST_OFFSET(n, c, hwd); )==""\n"
313R"==(#endif )==""\n"
314R"==(dst[dst_off] = TO_DST(0.0f); )==""\n"
315R"==(} )==""\n"
316R"==(if (n < FINAL_N_DIM && c < FINAL_C_DIM) { )==""\n"
317R"==(const int off = FINAL_SRC_OFFSET(n, c, hwd); )==""\n"
318R"==(const DEF_ACC_DATA_T data = src[off]; )==""\n"
319R"==(acc = ACCUMULATE_AGAIN(acc, data); )==""\n"
320R"==(} )==""\n"
321R"==(} )==""\n"
322R"==(} )==""\n"
323R"==(} )==""\n"
324R"==(if (n_start < DST_N && c_start < DST_C) { )==""\n"
325R"==(const int off = FINAL_DST_OFFSET(n_start, c_start, hwd_start); )==""\n"
326R"==(WRITE_FINAL_RESULT(dst[off], acc, n_start, c_start, hwd_start); )==""\n"
327R"==(} )==""\n"
328R"==(} )==""\n"
329R"==(#endif )==""\n"
330R"==()==";
331}
332}
333}
334}