1namespace dnnl {
2namespace impl {
3namespace gpu {
4namespace ocl {
5const char *ref_resampling_kernel = R"==(/******************************************************************************* )==""\n"
6R"==(* Copyright 2019-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 IS_FWD == 1 )==""\n"
23R"==(KERNEL_ATTR )==""\n"
24R"==(__kernel void ref_resampling_fwd( )==""\n"
25R"==(__global const DATA_T *src, __global DST_DATA_T *dst POST_OP_ARGS) { )==""\n"
26R"==(const uint mb = GWS_GET_MB(); )==""\n"
27R"==(const uint c = GWS_GET_C(); )==""\n"
28R"==(const uint od = GWS_GET_OD(); )==""\n"
29R"==(const uint oh = GWS_GET_OH(); )==""\n"
30R"==(const uint ow = GWS_GET_OW(); )==""\n"
31R"==(const float id = (od + .5f) * ID / OD; )==""\n"
32R"==(const float ih = (oh + .5f) * IH / OH; )==""\n"
33R"==(const float iw = (ow + .5f) * IW / OW; )==""\n"
34R"==(float result; )==""\n"
35R"==(const uint dst_index = DST_OFF(mb, c, od, oh, ow); )==""\n"
36R"==(if (mb >= DST_D0 || c >= DST_D1) { )==""\n"
37R"==(dst[dst_index] = TO_DST(0.0f); )==""\n"
38R"==(return; )==""\n"
39R"==(} )==""\n"
40R"==(#if RESAMPLING_ALG_NEAREST )==""\n"
41R"==(const uint src_index = SRC_OFF(mb, c, (uint)id, (uint)ih, (uint)iw); )==""\n"
42R"==(result = CONVERT_FLOAT_T(src[src_index]); )==""\n"
43R"==(#else )==""\n"
44R"==(const int id0 = max((int)floor(id - .5f), 0); )==""\n"
45R"==(const int id1 = min((int)ceil(id - .5f), ID - 1); )==""\n"
46R"==(const int ih0 = max((int)floor(ih - .5f), 0); )==""\n"
47R"==(const int ih1 = min((int)ceil(ih - .5f), IH - 1); )==""\n"
48R"==(const int iw0 = max((int)floor(iw - .5f), 0); )==""\n"
49R"==(const int iw1 = min((int)ceil(iw - .5f), IW - 1); )==""\n"
50R"==(const float wd[2] = {1.0f - fabs(id - .5f - id0), fabs(id - .5f - id0)}; )==""\n"
51R"==(const float wh[2] = {1.0f - fabs(ih - .5f - ih0), fabs(ih - .5f - ih0)}; )==""\n"
52R"==(const float ww[2] = {1.0f - fabs(iw - .5f - iw0), fabs(iw - .5f - iw0)}; )==""\n"
53R"==(const int ih_arr[2] = {ih0, ih1}; )==""\n"
54R"==(const int iw_arr[2] = {iw0, iw1}; )==""\n"
55R"==(float cd[2][2]; )==""\n"
56R"==(for_(int i = 0; i < 2; i++) )==""\n"
57R"==(for (int j = 0; j < 2; j++) )==""\n"
58R"==(cd[i][j] = CONVERT_FLOAT_T( )==""\n"
59R"==(src[SRC_OFF(mb, c, id0, ih_arr[i], iw_arr[j])]) )==""\n"
60R"==(* wd[0] )==""\n"
61R"==(+ CONVERT_FLOAT_T( )==""\n"
62R"==(src[SRC_OFF(mb, c, id1, ih_arr[i], iw_arr[j])]) )==""\n"
63R"==(* wd[1]; )==""\n"
64R"==(float ch[2]; )==""\n"
65R"==(for (int i = 0; i < 2; i++) )==""\n"
66R"==(ch[i] = cd[0][i] * wh[0] + cd[1][i] * wh[1]; )==""\n"
67R"==(result = ch[0] * ww[0] + ch[1] * ww[1]; )==""\n"
68R"==(#endif )==""\n"
69R"==(float sum_src; )==""\n"
70R"==(#if WITH_SUM )==""\n"
71R"==(sum_src = DST_TO_REF(dst[dst_index]); )==""\n"
72R"==(#endif )==""\n"
73R"==(#if NDIMS == 3 )==""\n"
74R"==(const unsigned po_d2 = ow; )==""\n"
75R"==(const unsigned po_d3 = 0; )==""\n"
76R"==(const unsigned po_d4 = 0; )==""\n"
77R"==(#elif NDIMS == 4 )==""\n"
78R"==(const unsigned po_d2 = oh; )==""\n"
79R"==(const unsigned po_d3 = ow; )==""\n"
80R"==(const unsigned po_d4 = 0; )==""\n"
81R"==(#elif NDIMS == 5 )==""\n"
82R"==(const unsigned po_d2 = od; )==""\n"
83R"==(const unsigned po_d3 = oh; )==""\n"
84R"==(const unsigned po_d4 = ow; )==""\n"
85R"==(#else )==""\n"
86R"==(const unsigned po_d2 = 0; )==""\n"
87R"==(const unsigned po_d3 = 0; )==""\n"
88R"==(const unsigned po_d4 = 0; )==""\n"
89R"==(#endif )==""\n"
90R"==(APPLY_POST_OPS_SERIAL(result, float, sum_src, float, mb, 1, c, 1, po_d2, 1, )==""\n"
91R"==(po_d3, 1, po_d4, 1, 0, 1); )==""\n"
92R"==(dst[dst_index] = TO_DST(result); )==""\n"
93R"==(} )==""\n"
94R"==(#endif )==""\n"
95R"==(#if IS_BWD == 1 )==""\n"
96R"==(float linear(int x, int fo, int fi) { )==""\n"
97R"==(return ((x + .5f) * fo / fi) - .5f; )==""\n"
98R"==(} )==""\n"
99R"==(KERNEL_ATTR )==""\n"
100R"==(__kernel void ref_resampling_bwd( )==""\n"
101R"==(__global DATA_T *diff_src, __global const DST_DATA_T *diff_dst) { )==""\n"
102R"==(#define CEIL(x) max((int)ceil(x), (int)0) )==""\n"
103R"==(#define L(x, fo, fi) linear(x, fo, fi) )==""\n"
104R"==(#define LS(x, fo, fi) CEIL(L(x, fo, fi)) )==""\n"
105R"==(#define RS(x, fo, fi) \ )==""\n"
106R"==(L((int)x - 1, fo, fi) < 0 ? 0 : (int)(L((int)x - 1, fo, fi)) + 1 )==""\n"
107R"==(#define LE(x, fo, fi, lim) min(CEIL(L(x + 1, fo, fi)), (int)lim) )==""\n"
108R"==(#define RE(x, fo, fi, lim) \ )==""\n"
109R"==(min((L(x, fo, fi) < 0 ? 0 : (int)(L(x, fo, fi)) + 1), (int)lim) )==""\n"
110R"==(const uint mb = GWS_GET_MB(); )==""\n"
111R"==(const uint c = GWS_GET_C(); )==""\n"
112R"==(const uint id = GWS_GET_ID(); )==""\n"
113R"==(const uint ih = GWS_GET_IH(); )==""\n"
114R"==(const uint iw = GWS_GET_IW(); )==""\n"
115R"==(const uint src_index = SRC_OFF(mb, c, id, ih, iw); )==""\n"
116R"==(if (mb >= DST_D0 || c >= DST_D1) { )==""\n"
117R"==(diff_src[src_index] = TO_DST(0.f); )==""\n"
118R"==(return; )==""\n"
119R"==(} )==""\n"
120R"==(#if RESAMPLING_ALG_NEAREST )==""\n"
121R"==(int od_start = CEIL(id * FD - .5f); )==""\n"
122R"==(int oh_start = CEIL(ih * FH - .5f); )==""\n"
123R"==(int ow_start = CEIL(iw * FW - .5f); )==""\n"
124R"==(int od_end = CEIL((id + 1.f) * FD - .5f); )==""\n"
125R"==(int oh_end = CEIL((ih + 1.f) * FH - .5f); )==""\n"
126R"==(int ow_end = CEIL((iw + 1.f) * FW - .5f); )==""\n"
127R"==(float src_val = 0; )==""\n"
128R"==(for (int i = od_start; i < od_end; i++) { )==""\n"
129R"==(for (int j = oh_start; j < oh_end; j++) { )==""\n"
130R"==(for (int k = ow_start; k < ow_end; k++) { )==""\n"
131R"==(const int dst_index = DST_OFF(mb, c, i, j, k); )==""\n"
132R"==(src_val += DST_TO_REF(diff_dst[dst_index]); )==""\n"
133R"==(} )==""\n"
134R"==(} )==""\n"
135R"==(} )==""\n"
136R"==(#else )==""\n"
137R"==(int left_sd = id == 0 ? 0 : LS(id, OD, ID); )==""\n"
138R"==(int left_sh = ih == 0 ? 0 : LS(ih, OH, IH); )==""\n"
139R"==(int left_sw = iw == 0 ? 0 : LS(iw, OW, IW); )==""\n"
140R"==(int right_sd = RS(id, OD, ID); )==""\n"
141R"==(int right_sh = RS(ih, OH, IH); )==""\n"
142R"==(int right_sw = RS(iw, OW, IW); )==""\n"
143R"==(int left_ed = LE(id, OD, ID, OD); )==""\n"
144R"==(int left_eh = LE(ih, OH, IH, OH); )==""\n"
145R"==(int left_ew = LE(iw, OW, IW, OW); )==""\n"
146R"==(int right_ed = id == (ID - 1) ? OD : RE(id, OD, ID, OD); )==""\n"
147R"==(int right_eh = ih == (IH - 1) ? OH : RE(ih, OH, IH, OH); )==""\n"
148R"==(int right_ew = iw == (IW - 1) ? OW : RE(iw, OW, IW, OW); )==""\n"
149R"==(int od_start[2] = {left_sd, right_sd}; )==""\n"
150R"==(int oh_start[2] = {left_sh, right_sh}; )==""\n"
151R"==(int ow_start[2] = {left_sw, right_sw}; )==""\n"
152R"==(int od_end[2] = {left_ed, right_ed}; )==""\n"
153R"==(int oh_end[2] = {left_eh, right_eh}; )==""\n"
154R"==(int ow_end[2] = {left_ew, right_ew}; )==""\n"
155R"==(float src_val = 0.0f; )==""\n"
156R"==(for (int c1 = 0; c1 < 2; c1++) { )==""\n"
157R"==(for (int c2 = 0; c2 < 2; c2++) { )==""\n"
158R"==(for (int c3 = 0; c3 < 2; c3++) { )==""\n"
159R"==(for (int i = od_start[c1]; i < od_end[c1]; i++) { )==""\n"
160R"==(for (int j = oh_start[c2]; j < oh_end[c2]; j++) { )==""\n"
161R"==(for (int k = ow_start[c3]; k < ow_end[c3]; k++) { )==""\n"
162R"==(float dst_val = DST_TO_REF( )==""\n"
163R"==(diff_dst[DST_OFF(mb, c, i, j, k)]); )==""\n"
164R"==(float d = L(i, ID, OD); )==""\n"
165R"==(float h = L(j, IH, OH); )==""\n"
166R"==(float w = L(k, IW, OW); )==""\n"
167R"==(float Wid = c1 == 0 ? 1.f - fabs(d - (int)d) )==""\n"
168R"==(: fabs(d - (int)d); )==""\n"
169R"==(float Wih = c2 == 0 ? 1.f - fabs(h - (int)h) )==""\n"
170R"==(: fabs(h - (int)h); )==""\n"
171R"==(float Wiw = c3 == 0 ? 1.f - fabs(w - (int)w) )==""\n"
172R"==(: fabs(w - (int)w); )==""\n"
173R"==(src_val += dst_val * Wid * Wih * Wiw; )==""\n"
174R"==(} )==""\n"
175R"==(} )==""\n"
176R"==(} )==""\n"
177R"==(} )==""\n"
178R"==(} )==""\n"
179R"==(} )==""\n"
180R"==(#endif )==""\n"
181R"==(#if DT_S32 == 1 )==""\n"
182R"==(diff_src[src_index] = CONVERT_DATA_T(src_val); )==""\n"
183R"==(#else )==""\n"
184R"==(diff_src[src_index] = TO_DATA_T(src_val); )==""\n"
185R"==(#endif )==""\n"
186R"==(} )==""\n"
187R"==(#endif )==""\n"
188R"==()==";
189}
190}
191}
192}