diff options
| -rw-r--r-- | src/core/CL/cl_kernels/direct_convolution9x9.cl | 4 | ||||
| -rw-r--r-- | src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp | 6 |
2 files changed, 8 insertions, 2 deletions
diff --git a/src/core/CL/cl_kernels/direct_convolution9x9.cl b/src/core/CL/cl_kernels/direct_convolution9x9.cl index 22f437f78e..8d0417a31d 100644 --- a/src/core/CL/cl_kernels/direct_convolution9x9.cl +++ b/src/core/CL/cl_kernels/direct_convolution9x9.cl @@ -284,7 +284,9 @@ __kernel void direct_convolution9x9_nhwc( const int id2 = get_global_id(2); __global uchar *weights_addr = (__global uchar *)tensor3D_offset(&weights, 0, 0, 0); - __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0) + ((id2 * STRIDE_Y) - PAD_TOP) * (int)src_stride_z; + __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0) - src_stride_x * id0 + ((id2 * STRIDE_Y) - PAD_TOP) * (int)src_stride_z; + + weights_addr += id0 * weights_stride_w; #if(PAD_TOP == 1) const int coordy = id2 - PAD_TOP; diff --git a/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp b/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp index b878a2121f..dc4c431c5d 100644 --- a/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp +++ b/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp @@ -239,6 +239,10 @@ inline void setup_num_elems(unsigned int &num_elems_read_per_iteration_x, unsign { num_elems_read_per_iteration_x = 4; } + else + { + num_elems_read_per_iteration_x = 1; + } switch(kernel_size) { @@ -346,7 +350,7 @@ std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITen if(data_layout == DataLayout::NHWC) { AccessWindowStatic input_access(input, 0, -conv_pad_left, - num_elems_read_per_iteration_x, + ceil_to_multiple(input->dimension(0), num_elems_read_per_iteration_x), ceil_to_multiple(input->dimension(1) + conv_info.pad_right(), num_elems_read_per_iteration_y)); AccessWindowStatic weights_access(weights, 0, 0, weights->dimension(0), weights->dimension(1)); AccessWindowRectangle output_access(output, 0, 0, num_elems_written_per_iteration_x, num_elems_written_per_iteration_y); |