26 #if defined(BATCH_SIZE) && defined(DATA_TYPE) && defined(WIDTH_IN) && defined(HEIGHT_IN) 60 __kernel
void space_to_batch_nchw(
72 const int pad_left_x = *((__global
int *)
offset(&pad, 0, 0));
73 const int pad_right_x = *((__global
int *)
offset(&pad, 1, 0));
74 const int pad_left_y = *((__global
int *)
offset(&pad, 0, 1));
75 const int pad_right_y = *((__global
int *)
offset(&pad, 1, 1));
80 const int out_x = get_global_id(0);
81 const int out_y = get_global_id(1);
82 const int z = get_global_id(2);
84 const int pos_x = out_x * block_x + ((batch_id / BATCH_IN) % block_x);
85 const int pos_y = out_y * block_y + ((batch_id / BATCH_IN) / block_x);
87 if(((pos_y >= pad_left_y) && (pos_y < pad_left_y + HEIGHT_IN) && (pos_x >= pad_left_x) && (pos_x < pad_left_x + WIDTH_IN)))
89 const int w = batch_id % BATCH_IN;
90 const int in_x = pos_x - pad_left_x;
91 const int in_y = pos_y - pad_left_y;
129 __kernel
void space_to_batch_nhwc(
141 const int pad_left_x = *((__global
int *)
offset(&pad, 0, 0));
142 const int pad_right_x = *((__global
int *)
offset(&pad, 1, 0));
143 const int pad_left_y = *((__global
int *)
offset(&pad, 0, 1));
144 const int pad_right_y = *((__global
int *)
offset(&pad, 1, 1));
149 const int out_x = get_global_id(1);
150 const int out_y = get_global_id(2);
151 const int z = get_global_id(0);
153 const int pos_x = out_x * block_x + ((batch_id / BATCH_IN) % block_x);
154 const int pos_y = out_y * block_y + ((batch_id / BATCH_IN) / block_x);
156 if(((pos_y >= pad_left_y) && (pos_y < pad_left_y + HEIGHT_IN) && (pos_x >= pad_left_x) && (pos_x < pad_left_x + WIDTH_IN)))
158 const int w = batch_id % BATCH_IN;
159 const int in_x = pos_x - pad_left_x;
160 const int in_y = pos_y - pad_left_y;
165 #endif // defined(BATCH_SIZE) && defined(DATA_TYPE) && defined(WIDTH_IN) && defined(HEIGHT_IN) 167 #if defined(BATCH_SIZE) && defined(DATA_TYPE) && defined(BLOCK_SHAPE_X) && defined(BLOCK_SHAPE_Y) && defined(PAD_LEFT_X) && defined(PAD_RIGHT_X) && defined(PAD_LEFT_Y) && defined(PAD_RIGHT_Y) && defined(WIDTH_IN) && defined(HEIGHT_IN) 197 __kernel
void space_to_batch_static_nchw(
205 int block_x = BLOCK_SHAPE_X;
206 int block_y = BLOCK_SHAPE_Y;
208 const int out_x = get_global_id(0);
209 const int out_y = get_global_id(1);
210 const int z = get_global_id(2);
212 const int pos_x = out_x * block_x + ((batch_id / BATCH_IN) % block_x);
213 const int pos_y = out_y * block_y + ((batch_id / BATCH_IN) / block_x);
215 if(pos_y >= PAD_LEFT_Y && pos_y < PAD_LEFT_Y + HEIGHT_IN && pos_x >= PAD_LEFT_X && pos_x < PAD_LEFT_X + WIDTH_IN)
217 const int w = batch_id % BATCH_IN;
218 const int in_x = pos_x - PAD_LEFT_X;
219 const int in_y = pos_y - PAD_LEFT_Y;
253 __kernel
void space_to_batch_static_nhwc(
261 int block_x = BLOCK_SHAPE_X;
262 int block_y = BLOCK_SHAPE_Y;
264 const int out_x = get_global_id(1);
265 const int out_y = get_global_id(2);
266 const int z = get_global_id(0);
268 const int pos_x = out_x * block_x + ((batch_id / BATCH_IN) % block_x);
269 const int pos_y = out_y * block_y + ((batch_id / BATCH_IN) / block_x);
271 if(pos_y >= PAD_LEFT_Y && pos_y < PAD_LEFT_Y + HEIGHT_IN && pos_x >= PAD_LEFT_X && pos_x < PAD_LEFT_X + WIDTH_IN)
273 const int w = batch_id % BATCH_IN;
274 const int in_x = pos_x - PAD_LEFT_X;
275 const int in_y = pos_y - PAD_LEFT_Y;
280 #endif // defined(BATCH_SIZE) && defined(DATA_TYPE) && defined(BLOCK_SHAPE_X) && defined(BLOCK_SHAPE_Y) && defined(PAD_LEFT_X) && defined(PAD_RIGHT_X) && defined(PAD_LEFT_Y) && defined(PAD_RIGHT_Y) && defined(WIDTH_IN) && defined(HEIGHT_IN) Structure to hold Vector information.
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
#define IMAGE_DECLARATION(name)
Structure to hold 3D tensor information.
Structure to hold 4D tensor information.
#define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size)
#define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name)
#define VECTOR_DECLARATION(name)
__global const uchar * tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)
Get the pointer position of a Tensor4D.
Structure to hold Image information.
#define CONVERT_TO_TENSOR3D_STRUCT(name)
__global const uchar * vector_offset(const Vector *vec, int x)
Get the pointer position of a Vector.
#define TENSOR4D_DECLARATION(name)
__global uchar * ptr
Pointer to the starting postion of the buffer.
#define TENSOR3D_DECLARATION(name)
#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name)