From edf26ea4921fd0c162abb962fc4863846e73f53f Mon Sep 17 00:00:00 2001 From: Michalis Spyrou Date: Wed, 21 Nov 2018 14:17:42 +0000 Subject: COMPMID-1451 Fix macros in SpaceToBatch Change-Id: I95fdf5bd85becfe081f6ae587284f3b294681308 --- src/core/CL/cl_kernels/space_to_batch.cl | 32 +++++++++++------------ 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(CLKernelLibrary::get().create_kernel("space_to_batch_static_" + lower_string(string_from_data_layout(input->info()->data_layout())), build_opts.options())); // Configure kernel window -- cgit v1.2.1