26 #if defined(DATA_TYPE) && defined(BATCH_SIZE) 57 __kernel
void batch_to_space_nchw(
67 const int block_x = *((__global
int *)
vector_offset(&block, 0));
68 const int block_y = *((__global
int *)
vector_offset(&block, 1));
70 const int r = (BATCH_SIZE / (block_x * block_y));
71 const int x = get_global_id(0);
72 const int y = get_global_id(1);
73 const int z = get_global_id(2);
74 const int w = batch_id % r;
76 const int out_x = x * block_x + (batch_id / r) % block_x;
77 const int out_y = y * block_y + (batch_id / r) / block_x;
111 __kernel
void batch_to_space_nhwc(
121 const int block_x = *((__global
int *)
vector_offset(&block, 0));
122 const int block_y = *((__global
int *)
vector_offset(&block, 1));
124 const int r = (BATCH_SIZE / (block_x * block_y));
125 const int x = get_global_id(1);
126 const int y = get_global_id(2);
127 const int z = get_global_id(0);
128 const int w = batch_id % r;
130 const int out_x = x * block_x + (batch_id / r) % block_x;
131 const int out_y = y * block_y + (batch_id / r) / block_x;
135 #endif // defined(DATA_TYPE) && defined(BATCH_SIZE) 137 #if defined(DATA_TYPE) && defined(BATCH_SIZE) && defined(BLOCK_SHAPE_X) && defined(BLOCK_SHAPE_Y) 163 __kernel
void batch_to_space_static_nchw(
171 const int block_x = BLOCK_SHAPE_X;
172 const int block_y = BLOCK_SHAPE_Y;
174 const int r = (BATCH_SIZE / (block_x * block_y));
175 const int x = get_global_id(0);
176 const int y = get_global_id(1);
177 const int z = get_global_id(2);
178 const int w = batch_id % r;
180 const int out_x = x * block_x + (batch_id / r) % block_x;
181 const int out_y = y * block_y + (batch_id / r) / block_x;
210 __kernel
void batch_to_space_static_nhwc(
218 const int block_x = BLOCK_SHAPE_X;
219 const int block_y = BLOCK_SHAPE_Y;
221 const int r = (BATCH_SIZE / (block_x * block_y));
222 const int x = get_global_id(1);
223 const int y = get_global_id(2);
224 const int z = get_global_id(0);
225 const int w = batch_id % r;
227 const int out_x = x * block_x + (batch_id / r) % block_x;
228 const int out_y = y * block_y + (batch_id / r) / block_x;
232 #endif // defined(DATA_TYPE) && defined(BATCH_SIZE) && defined(BLOCK_SHAPE_X) && defined(BLOCK_SHAPE_Y) Structure to hold Vector information.
Structure to hold 3D tensor information.
Structure to hold 4D tensor information.
#define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size)
#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.
#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)