diff options
Diffstat (limited to 'src/core')
-rw-r--r-- | src/core/CL/cl_kernels/convolution_layer.cl | 15 | ||||
-rw-r--r-- | src/core/CL/kernels/CLCol2ImKernel.cpp | 10 |
2 files changed, 13 insertions, 12 deletions
diff --git a/src/core/CL/cl_kernels/convolution_layer.cl b/src/core/CL/cl_kernels/convolution_layer.cl index cd886b1b0f..162632bce6 100644 --- a/src/core/CL/cl_kernels/convolution_layer.cl +++ b/src/core/CL/cl_kernels/convolution_layer.cl @@ -241,6 +241,7 @@ __kernel void im2col_kernel3x3_padx0_pady0( } #endif //defined(CONVOLVED_WIDTH) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_X) && defined(PAD_Y) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) +#if defined(WIDTH_OUTPUT) /** This kernel performs a reshaping of the output of the convolution layer. * * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float @@ -262,22 +263,22 @@ __kernel void im2col_kernel3x3_padx0_pady0( * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) - * @param[in] width The output tensor width */ __kernel void col2im( TENSOR3D_DECLARATION(src), TENSOR3D_DECLARATION(dst), - uint dst_stride_w, - uint width) + uint dst_stride_w) { Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(dst); - int idx = get_global_id(0) * dst.stride_z + (get_global_id(1) / width) * dst.stride_y + (get_global_id(1) % width) * dst.stride_x + get_global_id(2) * dst_stride_w; - __global uchar *tmp_out_ptr = dst.ptr + idx; - *((__global DATA_TYPE *)tmp_out_ptr) = *((__global DATA_TYPE *)(src.ptr)); -} + // Compute output offset + int idx = get_global_id(0) * dst.stride_z + (get_global_id(1) / WIDTH_OUTPUT) * dst_stride_y + (get_global_id(1) % WIDTH_OUTPUT) * dst_stride_x + get_global_id(2) * dst_stride_w; + // Store value + *((__global DATA_TYPE *)(dst.ptr + idx)) = *((__global DATA_TYPE *)(src.ptr)); +} +#endif // defined(WIDTH_OUTPUT) /** This kernel reshapes the tensor's low three dimensions to single row for GEMM operation * * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float diff --git a/src/core/CL/kernels/CLCol2ImKernel.cpp b/src/core/CL/kernels/CLCol2ImKernel.cpp index c1a1ef2829..c7884e3c15 100644 --- a/src/core/CL/kernels/CLCol2ImKernel.cpp +++ b/src/core/CL/kernels/CLCol2ImKernel.cpp @@ -64,6 +64,7 @@ void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::p // Create kernel std::set<std::string> build_opts = { ("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())) }; + build_opts.emplace("-DWIDTH_OUTPUT=" + support::cpp11::to_string(_convolved_dims.first)); if(is_data_type_fixed_point(input->info()->data_type())) { build_opts.emplace("-DFIXED_POINT_POSITION=" + support::cpp11::to_string(input->info()->fixed_point_position())); @@ -71,11 +72,6 @@ void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::p _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel("col2im", build_opts)); - // Set static kernel arguments - unsigned int idx = 2 * num_arguments_per_3D_tensor(); - _kernel.setArg<cl_uint>(idx++, _output->info()->strides_in_bytes()[3]); - _kernel.setArg<cl_uint>(idx++, _convolved_dims.first); - // Configure window Window win = calculate_max_window(*input->info(), Steps()); @@ -97,6 +93,10 @@ void CLCol2ImKernel::run(const Window &window, cl::CommandQueue &queue) Window collapsed_window = window.collapse_if_possible(ICLKernel::window(), Window::DimZ); Window slice = collapsed_window.first_slice_window_3D(); + // Set static kernel arguments + unsigned int idx = 2 * num_arguments_per_3D_tensor(); + _kernel.setArg<cl_uint>(idx++, _output->info()->strides_in_bytes()[3]); + do { // Set inputs |