1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
|
/******************************************************************************/
/**************************** Macros ******************************************/
/******************************************************************************/
// Error handling
#ifndef IMAGE_DIMS
#error "IMAGE_DIMS must be defined"
#endif
#ifndef FILTER_DIMS
#error "FILTER_DIMS must be defined"
#endif
#if FILTER_DIMS > IMAGE_DIMS
#error "Filter cannot have more dimensions than image"
#endif
// Boundary handling modes
#define CONV_MODE_REFLECT 0 // CLK_ADDRESS_MIRRORED_REPEAT : cba|abcd|dcb
#define CONV_MODE_NEAREST 1 // CLK_ADDRESS_CLAMP_TO_EDGE : aaa|abcd|ddd
#define CONV_MODE_WRAP 2 // CLK_ADDRESS_REPEAT : bcd|abcd|abc
#define CONV_MODE_CONSTANT 3 // CLK_ADDRESS_CLAMP : 000|abcd|000
#ifndef USED_CONV_MODE
#define USED_CONV_MODE CONV_MODE_NEAREST
#endif
#if USED_CONV_MODE == CONV_MODE_REFLECT
#define CLK_BOUNDARY CLK_ADDRESS_MIRRORED_REPEAT
#define CLK_COORDS CLK_NORMALIZED_COORDS_TRUE
#define USE_NORM_COORDS
#elif USED_CONV_MODE == CONV_MODE_NEAREST
#define CLK_BOUNDARY CLK_ADDRESS_CLAMP_TO_EDGE
#define CLK_COORDS CLK_NORMALIZED_COORDS_FALSE
#elif USED_CONV_MODE == CONV_MODE_WRAP
#define CLK_BOUNDARY CLK_ADDRESS_REPEAT
#define CLK_COORDS CLK_NORMALIZED_COORDS_TRUE
#define USE_NORM_COORDS
#elif USED_CONV_MODE == CONV_MODE_CONSTANT
#define CLK_BOUNDARY CLK_ADDRESS_CLAMP
#define CLK_COORDS CLK_NORMALIZED_COORDS_FALSE
#else
#error "Unknown convolution mode"
#endif
// Convolution index for filter
#define FILTER_INDEX(j) (Lx - 1 - j)
// Filter access patterns
#define READ_FILTER_1D(j) read_imagef(filter, (int2) (FILTER_INDEX(j), 0)).x;
#define READ_FILTER_2D(jx, jy) read_imagef(filter, (int2) (FILTER_INDEX(jx), FILTER_INDEX(jy))).x;
#define READ_FILTER_3D(jx, jy, jz) read_imagef(filter, (int4) (FILTER_INDEX(jx), FILTER_INDEX(jy), FILTER_INDEX(jz), 0)).x;
// Convolution index for image
#ifdef USE_NORM_COORDS
#define IMAGE_INDEX_X (gidx*1.0f +0.5f - c + jx)/Nx
#define IMAGE_INDEX_Y (gidy*1.0f +0.5f - c + jy)/Ny
#define IMAGE_INDEX_Z (gidz*1.0f +0.5f - c + jz)/Nz
#define RET_TYPE_1 float
#define RET_TYPE_2 float2
#define RET_TYPE_4 float4
#define C_ZERO 0.5f
#define GIDX (gidx*1.0f + 0.5f)/Nx
#define GIDY (gidy*1.0f + 0.5f)/Ny
#define GIDZ (gidz*1.0f + 0.5f)/Nz
#else
#define IMAGE_INDEX_X (gidx - c + jx)
#define IMAGE_INDEX_Y (gidy - c + jy)
#define IMAGE_INDEX_Z (gidz - c + jz)
#define RET_TYPE_1 int
#define RET_TYPE_2 int2
#define RET_TYPE_4 int4
#define C_ZERO 0
#define GIDX gidx
#define GIDY gidy
#define GIDZ gidz
#endif
static const sampler_t sampler = CLK_COORDS | CLK_BOUNDARY | CLK_FILTER_NEAREST;
// Image access patterns
#define READ_IMAGE_1D read_imagef(input, sampler, (RET_TYPE_2) (IMAGE_INDEX_X, C_ZERO)).x
#define READ_IMAGE_2D_X read_imagef(input, sampler, (RET_TYPE_2) (IMAGE_INDEX_X , GIDY)).x
#define READ_IMAGE_2D_Y read_imagef(input, sampler, (RET_TYPE_2) (GIDX, IMAGE_INDEX_Y)).x
#define READ_IMAGE_2D_XY read_imagef(input, sampler, (RET_TYPE_2) (IMAGE_INDEX_X, IMAGE_INDEX_Y)).x
#define READ_IMAGE_3D_X read_imagef(input, sampler, (RET_TYPE_4) (IMAGE_INDEX_X, GIDY, GIDZ, C_ZERO)).x
#define READ_IMAGE_3D_Y read_imagef(input, sampler, (RET_TYPE_4) (GIDX, IMAGE_INDEX_Y, GIDZ, C_ZERO)).x
#define READ_IMAGE_3D_Z read_imagef(input, sampler, (RET_TYPE_4) (GIDX, GIDY, IMAGE_INDEX_Z, C_ZERO)).x
#define READ_IMAGE_3D_XY read_imagef(input, sampler, (RET_TYPE_4) (IMAGE_INDEX_X, IMAGE_INDEX_Y, GIDZ, C_ZERO)).x
#define READ_IMAGE_3D_XZ read_imagef(input, sampler, (RET_TYPE_4) (IMAGE_INDEX_X, GIDY, IMAGE_INDEX_Z, C_ZERO)).x
#define READ_IMAGE_3D_YZ read_imagef(input, sampler, (RET_TYPE_4) (GIDX, IMAGE_INDEX_Y, IMAGE_INDEX_Z, C_ZERO)).x
#define READ_IMAGE_3D_XYZ read_imagef(input, sampler, (RET_TYPE_4) (IMAGE_INDEX_X, IMAGE_INDEX_Y, IMAGE_INDEX_Z, C_ZERO)).x
// NOTE: pyopencl and OpenCL < 1.2 do not support image1d_t
#if FILTER_DIMS == 1
#define FILTER_TYPE image2d_t
#define READ_FILTER_VAL(j) READ_FILTER_1D(j)
#elif FILTER_DIMS == 2
#define FILTER_TYPE image2d_t
#define READ_FILTER_VAL(jx, jy) READ_FILTER_2D(jx, jy)
#elif FILTER_DIMS == 3
#define FILTER_TYPE image3d_t
#define READ_FILTER_VAL(jx, jy, jz) READ_FILTER_3D(jx, jy, jz)
#endif
#if IMAGE_DIMS == 1
#define IMAGE_TYPE image2d_t
#define READ_IMAGE_X READ_IMAGE_1D
#elif IMAGE_DIMS == 2
#define IMAGE_TYPE image2d_t
#define READ_IMAGE_X READ_IMAGE_2D_X
#define READ_IMAGE_Y READ_IMAGE_2D_Y
#define READ_IMAGE_XY READ_IMAGE_2D_XY
#elif IMAGE_DIMS == 3
#define IMAGE_TYPE image3d_t
#define READ_IMAGE_X READ_IMAGE_3D_X
#define READ_IMAGE_Y READ_IMAGE_3D_Y
#define READ_IMAGE_Z READ_IMAGE_3D_Z
#define READ_IMAGE_XY READ_IMAGE_3D_XY
#define READ_IMAGE_XZ READ_IMAGE_3D_XZ
#define READ_IMAGE_YZ READ_IMAGE_3D_YZ
#define READ_IMAGE_XYZ READ_IMAGE_3D_XYZ
#endif
// Get the center index of the filter,
// and the "half-Left" and "half-Right" lengths.
// In the case of an even-sized filter, the center is shifted to the left.
#define GET_CENTER_HL(hlen){\
if (hlen & 1) {\
c = hlen/2;\
hL = c;\
hR = c;\
}\
else {\
c = hlen/2 - 1;\
hL = c;\
hR = c+1;\
}\
}\
/******************************************************************************/
/**************************** 1D Convolution **********************************/
/******************************************************************************/
#if FILTER_DIMS == 1
// Convolution with 1D kernel along axis "X" (fast dimension)
// Works for batched 1D on 2D and batched 2D on 3D, along axis "X".
__kernel void convol_1D_X_tex(
read_only IMAGE_TYPE input,
__global float * output,
read_only FILTER_TYPE filter,
int Lx, // filter size
int Nx, // input/output number of columns
int Ny, // input/output number of rows
int Nz // input/output depth
)
{
uint gidx = get_global_id(0);
uint gidy = get_global_id(1);
uint gidz = get_global_id(2);
if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return;
int c, hL, hR;
GET_CENTER_HL(Lx);
float sum = 0.0f;
for (int jx = 0; jx <= hR+hL; jx++) {
sum += READ_IMAGE_X * READ_FILTER_VAL(jx);
}
output[(gidz*Ny + gidy)*Nx + gidx] = sum;
}
#if IMAGE_DIMS >= 2
// Convolution with 1D kernel along axis "Y"
// Works for batched 1D on 2D and batched 2D on 3D, along axis "Y".
__kernel void convol_1D_Y_tex(
read_only IMAGE_TYPE input,
__global float * output,
read_only FILTER_TYPE filter,
int Lx, // filter size
int Nx, // input/output number of columns
int Ny, // input/output number of rows
int Nz // input/output depth
)
{
uint gidx = get_global_id(0);
uint gidy = get_global_id(1);
uint gidz = get_global_id(2);
if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return;
int c, hL, hR;
GET_CENTER_HL(Lx);
float sum = 0.0f;
for (int jy = 0; jy <= hR+hL; jy++) {
sum += READ_IMAGE_Y * READ_FILTER_VAL(jy);
}
output[(gidz*Ny + gidy)*Nx + gidx] = sum;
}
#endif
#if IMAGE_DIMS == 3
// Convolution with 1D kernel along axis "Z"
// Works for batched 1D on 2D and batched 2D on 3D, along axis "Z".
__kernel void convol_1D_Z_tex(
read_only IMAGE_TYPE input,
__global float * output,
read_only FILTER_TYPE filter,
int Lx, // filter size
int Nx, // input/output number of columns
int Ny, // input/output number of rows
int Nz // input/output depth
)
{
uint gidx = get_global_id(0);
uint gidy = get_global_id(1);
uint gidz = get_global_id(2);
if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return;
int c, hL, hR;
GET_CENTER_HL(Lx);
float sum = 0.0f;
for (int jz = 0; jz <= hR+hL; jz++) {
sum += READ_IMAGE_Z * READ_FILTER_VAL(jz);
}
output[(gidz*Ny + gidy)*Nx + gidx] = sum;
}
#endif
#endif
/******************************************************************************/
/**************************** 2D Convolution **********************************/
/******************************************************************************/
#if IMAGE_DIMS >= 2 && FILTER_DIMS == 2
// Convolution with 2D kernel
// Works for batched 2D on 3D.
__kernel void convol_2D_XY_tex(
read_only IMAGE_TYPE input,
__global float * output,
read_only FILTER_TYPE filter,
int Lx, // filter number of columns,
int Ly, // filter number of rows,
int Nx, // input/output number of columns
int Ny, // input/output number of rows
int Nz // input/output depth
)
{
uint gidx = get_global_id(0);
uint gidy = get_global_id(1);
uint gidz = get_global_id(2);
if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return;
int c, hL, hR;
GET_CENTER_HL(Lx);
float sum = 0.0f;
for (int jy = 0; jy <= hR+hL; jy++) {
for (int jx = 0; jx <= hR+hL; jx++) {
sum += READ_IMAGE_XY * READ_FILTER_VAL(jx, jy);
}
}
output[(gidz*Ny + gidy)*Nx + gidx] = sum;
}
#endif
#if IMAGE_DIMS == 3 && FILTER_DIMS == 2
// Convolution with 2D kernel
// Works for batched 2D on 3D.
__kernel void convol_2D_XZ_tex(
read_only IMAGE_TYPE input,
__global float * output,
read_only FILTER_TYPE filter,
int Lx, // filter number of columns,
int Lz, // filter number of rows,
int Nx, // input/output number of columns
int Ny, // input/output number of rows
int Nz // input/output depth
)
{
uint gidx = get_global_id(0);
uint gidy = get_global_id(1);
uint gidz = get_global_id(2);
if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return;
int c, hL, hR;
GET_CENTER_HL(Lx);
float sum = 0.0f;
for (int jz = 0; jz <= hR+hL; jz++) {
for (int jx = 0; jx <= hR+hL; jx++) {
sum += READ_IMAGE_XZ * READ_FILTER_VAL(jx, jz);
}
}
output[(gidz*Ny + gidy)*Nx + gidx] = sum;
}
// Convolution with 2D kernel
// Works for batched 2D on 3D.
__kernel void convol_2D_YZ_tex(
read_only IMAGE_TYPE input,
__global float * output,
read_only FILTER_TYPE filter,
int Lx, // filter number of columns,
int Lz, // filter number of rows,
int Nx, // input/output number of columns
int Ny, // input/output number of rows
int Nz // input/output depth
)
{
uint gidx = get_global_id(0);
uint gidy = get_global_id(1);
uint gidz = get_global_id(2);
if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return;
int c, hL, hR;
GET_CENTER_HL(Lx);
float sum = 0.0f;
for (int jz = 0; jz <= hR+hL; jz++) {
for (int jy = 0; jy <= hR+hL; jy++) {
sum += READ_IMAGE_YZ * READ_FILTER_VAL(jy, jz);
}
}
output[(gidz*Ny + gidy)*Nx + gidx] = sum;
}
#endif
/******************************************************************************/
/**************************** 3D Convolution **********************************/
/******************************************************************************/
#if IMAGE_DIMS == 3 && FILTER_DIMS == 3
// Convolution with 3D kernel
__kernel void convol_3D_XYZ_tex(
read_only IMAGE_TYPE input,
__global float * output,
read_only FILTER_TYPE filter,
int Lx, // filter number of columns,
int Ly, // filter number of rows,
int Lz, // filter number of rows,
int Nx, // input/output number of columns
int Ny, // input/output number of rows
int Nz // input/output depth
)
{
uint gidx = get_global_id(0);
uint gidy = get_global_id(1);
uint gidz = get_global_id(2);
if ((gidx >= Nx) || (gidy >= Ny) || (gidz >= Nz)) return;
int c, hL, hR;
GET_CENTER_HL(Lx);
float sum = 0.0f;
for (int jz = 0; jz <= hR+hL; jz++) {
for (int jy = 0; jy <= hR+hL; jy++) {
for (int jx = 0; jx <= hR+hL; jx++) {
sum += READ_IMAGE_XYZ * READ_FILTER_VAL(jx, jy, jz);
}
}
}
output[(gidz*Ny + gidy)*Nx + gidx] = sum;
}
#endif
|
