diff options
| author | Michalis Spyrou <michalis.spyrou@arm.com> | 2018-11-21 14:17:42 +0000 |
|---|---|---|
| committer | Isabella Gottardi <isabella.gottardi@arm.com> | 2018-11-21 17:14:43 +0000 |
| commit | edf26ea4921fd0c162abb962fc4863846e73f53f (patch) | |
| tree | b5e32c9a53c55035c08e4ba0480da2446f9f7937 /src/core/CL/cl_kernels/space_to_batch.cl | |
| parent | be2de40cffe375ba521ce8b3575277b81a71d0bf (diff) | |
| download | ComputeLibrary-edf26ea4921fd0c162abb962fc4863846e73f53f.tar.gz | |
COMPMID-1451 Fix macros in SpaceToBatch
Change-Id: I95fdf5bd85becfe081f6ae587284f3b294681308
Diffstat (limited to 'src/core/CL/cl_kernels/space_to_batch.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/space_to_batch.cl | 32 |
1 files changed, 16 insertions, 16 deletions
diff --git a/src/core/CL/cl_kernels/space_to_batch.cl b/src/core/CL/cl_kernels/space_to_batch.cl index 1d641cc0f4..d42a79d3ff 100644 --- a/src/core/CL/cl_kernels/space_to_batch.cl +++ b/src/core/CL/cl_kernels/space_to_batch.cl @@ -71,10 +71,10 @@ __kernel void space_to_batch_nchw( Vector block = CONVERT_TO_VECTOR_STRUCT_NO_STEP(block_shape); Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output); - const int PAD_LEFT_X = *((__global int *)offset(&pad, 0, 0)); - const int PAD_RIGHT_X = *((__global int *)offset(&pad, 1, 0)); - const int PAD_LEFT_Y = *((__global int *)offset(&pad, 0, 1)); - const int PAD_RIGHT_Y = *((__global int *)offset(&pad, 1, 1)); + const int pad_left_x = *((__global int *)offset(&pad, 0, 0)); + const int pad_right_x = *((__global int *)offset(&pad, 1, 0)); + const int pad_left_y = *((__global int *)offset(&pad, 0, 1)); + const int pad_right_y = *((__global int *)offset(&pad, 1, 1)); int block_x = *((__global int *)vector_offset(&block, 0)); int block_y = *((__global int *)vector_offset(&block, 1)); @@ -83,12 +83,12 @@ __kernel void space_to_batch_nchw( const int out_y = get_global_id(1); const int z = get_global_id(2); - if((out_x >= PAD_LEFT_X && out_x < WIDTH_OUT - PAD_RIGHT_X) && (out_y >= PAD_LEFT_Y && out_y < HEIGHT_OUT - PAD_RIGHT_Y)) + if((out_x >= pad_left_x && out_x < WIDTH_OUT - pad_right_x) && (out_y >= pad_left_y && out_y < HEIGHT_OUT - pad_right_y)) { const int r = (BATCH_SIZE / (block_x * block_y)); const int w = batch_id % r; - const int in_x = (out_x - PAD_LEFT_X) * block_x + (batch_id / r) % block_x; - const int in_y = (out_y - PAD_LEFT_Y) * block_y + (batch_id / r) / block_x; + const int in_x = (out_x - pad_left_x) * block_x + (batch_id / r) % block_x; + const int in_y = (out_y - pad_left_y) * block_y + (batch_id / r) / block_x; *((__global DATA_TYPE *)out.ptr) = *((__global DATA_TYPE *)tensor4D_offset(&in, in_x, in_y, z, w)); } } @@ -139,10 +139,10 @@ __kernel void space_to_batch_nhwc( Vector block = CONVERT_TO_VECTOR_STRUCT_NO_STEP(block_shape); Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output); - const int PAD_LEFT_X = *((__global int *)offset(&pad, 0, 0)); - const int PAD_RIGHT_X = *((__global int *)offset(&pad, 1, 0)); - const int PAD_LEFT_Y = *((__global int *)offset(&pad, 0, 1)); - const int PAD_RIGHT_Y = *((__global int *)offset(&pad, 1, 1)); + const int pad_left_x = *((__global int *)offset(&pad, 0, 0)); + const int pad_right_x = *((__global int *)offset(&pad, 1, 0)); + const int pad_left_y = *((__global int *)offset(&pad, 0, 1)); + const int pad_right_y = *((__global int *)offset(&pad, 1, 1)); int block_x = *((__global int *)vector_offset(&block, 0)); int block_y = *((__global int *)vector_offset(&block, 1)); @@ -151,12 +151,12 @@ __kernel void space_to_batch_nhwc( const int out_y = get_global_id(2); const int z = get_global_id(0); - if((out_x >= PAD_LEFT_X && out_x < WIDTH_OUT - PAD_RIGHT_X) && (out_y >= PAD_LEFT_Y && out_y < HEIGHT_OUT - PAD_RIGHT_Y)) + if((out_x >= pad_left_x && out_x < WIDTH_OUT - pad_right_x) && (out_y >= pad_left_y && out_y < HEIGHT_OUT - pad_right_y)) { const int r = (BATCH_SIZE / (block_x * block_y)); const int w = batch_id % r; - const int in_x = (out_x - PAD_LEFT_X) * block_x + (batch_id / r) % block_x; - const int in_y = (out_y - PAD_LEFT_Y) * block_y + (batch_id / r) / block_x; + const int in_x = (out_x - pad_left_x) * block_x + (batch_id / r) % block_x; + const int in_y = (out_y - pad_left_y) * block_y + (batch_id / r) / block_x; *((__global DATA_TYPE *)out.ptr) = *((__global DATA_TYPE *)tensor4D_offset(&in, z, in_x, in_y, w)); } } @@ -201,7 +201,7 @@ __kernel void space_to_batch_static_nchw( Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output); int block_x = BLOCK_SHAPE_X; - int block_y = *((__global int *)vector_offset(&block, 1)); + int block_y = BLOCK_SHAPE_Y; const int out_x = get_global_id(0); const int out_y = get_global_id(1); @@ -254,7 +254,7 @@ __kernel void space_to_batch_static_nhwc( Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output); int block_x = BLOCK_SHAPE_X; - int block_y = *((__global int *)vector_offset(&block, 1)); + int block_y = BLOCK_SHAPE_Y; const int out_x = get_global_id(1); const int out_y = get_global_id(2); |