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 | |
parent | be2de40cffe375ba521ce8b3575277b81a71d0bf (diff) | |
download | ComputeLibrary-edf26ea4921fd0c162abb962fc4863846e73f53f.tar.gz |
COMPMID-1451 Fix macros in SpaceToBatch
Change-Id: I95fdf5bd85becfe081f6ae587284f3b294681308
Diffstat (limited to 'src/core')
-rw-r--r-- | src/core/CL/cl_kernels/space_to_batch.cl | 32 | ||||
-rw-r--r-- | src/core/CL/kernels/CLSpaceToBatchLayerKernel.cpp | 8 |
2 files changed, 20 insertions, 20 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); diff --git a/src/core/CL/kernels/CLSpaceToBatchLayerKernel.cpp b/src/core/CL/kernels/CLSpaceToBatchLayerKernel.cpp index 9e4010e6c6..d488631ae9 100644 --- a/src/core/CL/kernels/CLSpaceToBatchLayerKernel.cpp +++ b/src/core/CL/kernels/CLSpaceToBatchLayerKernel.cpp @@ -134,10 +134,10 @@ void CLSpaceToBatchLayerKernel::configure(const ICLTensor *input, const int bloc build_opts.add_option("-DBATCH_SIZE=" + support::cpp11::to_string(output->info()->dimension(idx_batch))); build_opts.add_option("-DBLOCK_SHAPE_X=" + support::cpp11::to_string(block_shape_x)); build_opts.add_option("-DBLOCK_SHAPE_Y=" + support::cpp11::to_string(block_shape_y)); - build_opts.add_option("-DPAD_START_X=" + support::cpp11::to_string(padding_left.x())); - build_opts.add_option("-DPAD_END_X=" + support::cpp11::to_string(padding_right.x())); - build_opts.add_option("-DPAD_START_Y=" + support::cpp11::to_string(padding_left.y())); - build_opts.add_option("-DPAD_END_Y=" + support::cpp11::to_string(padding_right.y())); + build_opts.add_option("-DPAD_LEFT_X=" + support::cpp11::to_string(padding_left.x())); + build_opts.add_option("-DPAD_RIGHT_X=" + support::cpp11::to_string(padding_right.x())); + build_opts.add_option("-DPAD_LEFT_Y=" + support::cpp11::to_string(padding_left.y())); + build_opts.add_option("-DPAD_RIGHT_Y=" + support::cpp11::to_string(padding_right.y())); _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel("space_to_batch_static_" + lower_string(string_from_data_layout(input->info()->data_layout())), build_opts.options())); // Configure kernel window |