diff --git a/src/shogun/mathematics/linalg/internal/implementation/Convolve.h b/src/shogun/mathematics/linalg/internal/implementation/Convolve.h deleted file mode 100644 index 6fce30b8e77..00000000000 --- a/src/shogun/mathematics/linalg/internal/implementation/Convolve.h +++ /dev/null @@ -1,418 +0,0 @@ -/* - * Copyright (c) The Shogun Machine Learning Toolbox - * Written (w) 2014 Khaled Nasr - * All rights reserved. - * - * Redistribution and use in source and binary forms, with or without - * modification, are permitted provided that the following conditions are met: - * - * 1. Redistributions of source code must retain the above copyright notice, this - * list of conditions and the following disclaimer. - * 2. Redistributions in binary form must reproduce the above copyright notice, - * this list of conditions and the following disclaimer in the documentation - * and/or other materials provided with the distribution. - * - * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND - * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED - * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE - * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR - * ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES - * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; - * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND - * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT - * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS - * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - * - * The views and conclusions contained in the software and documentation are those - * of the authors and should not be interpreted as representing official policies, - * either expressed or implied, of the Shogun Development Team. - */ - -#ifndef CONVOLVE_H_ -#define CONVOLVE_H_ - -#include -#include -#include - -#include - -#include - -#ifdef HAVE_VIENNACL -#include -#include -#endif // HAVE_VIENNACL - -namespace shogun -{ - -namespace linalg -{ - -namespace implementation -{ - -/** Generic class which is specialized for different backends to perform - * the convolve operation - */ -template -struct convolve -{ - /** The scalar type */ - typedef typename Matrix::Scalar T; - - /** Computes the 2D convolution of X with W - * - * NOTE: For the ViennaCL backend, the size of W (number of bytes) must not exceed - * [CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE](http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clGetDeviceInfo.html). - * - * @param X Input image - * @param W Filter coefficients. The dimensions of the matrix must be odd-numbered. - * @param Y Output image of the same size as the input image, as the borders - * of the input image are implicitly padded with zeros during the computation - * @param flip If true the filter coefficients are flipped, performing cross-correlation - * instead of convolution - * @param overwrite If true, the values in Y are overwritten with result of the - * computation. Otherwise, the result is added to the existing values in Y. - * @param stride_x Stride in the x (column) direction - * @param stride_y Stride in the y (row) direction - */ - static void compute(Matrix X, Matrix W, Matrix Y, bool flip , - bool overwrite, int32_t stride_x, int32_t stride_y); -}; - - -/** Partial specialization of convolve for the Eigen3 backend */ -template -struct convolve -{ - /** The scalar type */ - typedef typename Matrix::Scalar T; - - /** Eigen3 matrix type */ - typedef Eigen::Matrix MatrixXt; - - /** Eigen3 vector type */ - typedef Eigen::Matrix VectorXt; - - /** Computes the 2D convolution of X with W - * - * @param X Input image - * @param W Filter coefficients. The dimensions of the matrix must be odd-numbered. - * @param Y Output image of the same size as the input image, as the borders - * of the input image are implicitly padded with zeros during the computation - * @param flip If true the filter coefficients are flipped, performing cross-correlation - * instead of convolution - * @param overwrite If true, the values in Y are overwritten with result of the - * computation. Otherwise, the result is added to the existing values in Y. - * @param stride_x Stride in the x (column) direction - * @param stride_y Stride in the y (row) direction - */ - static void compute(SGMatrix X, SGMatrix W, SGMatrix Y, bool flip , - bool overwrite, int32_t stride_x, int32_t stride_y) - { - int32_t width = X.num_cols; - int32_t height = X.num_rows; - - int32_t kx = W.num_cols; - int32_t ky = W.num_rows; - - int32_t rx = (kx-1)/2; - int32_t ry = (ky-1)/2; - - for (int32_t x=0; x=0 && y1>=0 && x1 -struct convolve -{ - /** The scalar type */ - typedef typename Matrix::Scalar T; - - /** Generates the computation kernel for convolution with a stride of 1*/ - template - static viennacl::ocl::kernel& generate_kernel_unity_stride( - int32_t radius_x, int32_t radius_y, bool flip, bool overwrite) - { - std::string kernel_name = - "convolve_unity_stride_" + ocl::get_type_string() + "_" + - std::to_string(radius_x) + "_" + std::to_string(radius_y); - - if (flip) kernel_name.append("_flip"); - if (overwrite) kernel_name.append("_overwrite"); - - if (ocl::kernel_exists(kernel_name)) - return ocl::get_kernel(kernel_name); - - std::string source = ocl::generate_kernel_preamble(kernel_name); - - if (flip) source.append("#define FLIP\n"); - if (overwrite) source.append("#define OVERWRITE\n"); - - source.append("#define RADIUS_X " + std::to_string(radius_x) + "\n"); - source.append("#define RADIUS_Y " + std::to_string(radius_y) + "\n"); - - source.append( - R"( - #define W_WIDTH (2*RADIUS_X+1) - #define W_HEIGHT (2*RADIUS_Y+1) - - #define X_LOCAL_WIDTH (WORK_GROUP_SIZE_2D+2*RADIUS_X) - #define X_LOCAL_HEIGHT (WORK_GROUP_SIZE_2D+2*RADIUS_Y) - - inline DATATYPE readX(read_only __global DATATYPE* X, int x, int y, - int X_width, int X_height, int X_offset) - { - if (x>=0 && y>=0 && x=X_width || y>=X_height) - return; - - DATATYPE sum = 0; - for (int x1=0; x1 - static viennacl::ocl::kernel& generate_kernel_arbitrary_stride( - int32_t radius_x, int32_t radius_y, bool flip, bool overwrite) - { - std::string kernel_name = - "convolve_arbitrary_stride_" + ocl::get_type_string() + "_" + - std::to_string(radius_x) + "_" + std::to_string(radius_y); - - if (flip) kernel_name.append("_flip"); - if (overwrite) kernel_name.append("_overwrite"); - - if (ocl::kernel_exists(kernel_name)) - return ocl::get_kernel(kernel_name); - - std::string source = ocl::generate_kernel_preamble(kernel_name); - - if (flip) source.append("#define FLIP\n"); - if (overwrite) source.append("#define OVERWRITE\n"); - - source.append("#define RADIUS_X " + std::to_string(radius_x) + "\n"); - source.append("#define RADIUS_Y " + std::to_string(radius_y) + "\n"); - - source.append( - R"( - #define W_WIDTH (2*RADIUS_X+1) - #define W_HEIGHT (2*RADIUS_Y+1) - - #define X_LOCAL_WIDTH (WORK_GROUP_SIZE_2D+2*RADIUS_X) - #define X_LOCAL_HEIGHT (WORK_GROUP_SIZE_2D+2*RADIUS_Y) - - __kernel void KERNEL_NAME( - read_only __global DATATYPE* X, int X_width, int X_height, int X_offset, - __constant DATATYPE* W, int W_offset, - __global DATATYPE* Y, int Y_offset, - int stride_x, int stride_y) - { - __local DATATYPE X_local[WORK_GROUP_SIZE_2D][WORK_GROUP_SIZE_2D]; - - int x = get_global_id(0)*stride_x; - int y = get_global_id(1)*stride_y; - - int Y_width = X_width/stride_x; - int Y_height = X_height/stride_y; - - if (get_global_id(0)>=Y_width || get_global_id(1)>=Y_height) - return; - - DATATYPE sum = 0; - for (int x1=0; x1=0 && iny>=0 && inx X, CGPUMatrix W, CGPUMatrix Y, bool flip , - bool overwrite, int32_t stride_x, int32_t stride_y) - { - if (stride_x==1 && stride_y==1) - { - viennacl::ocl::kernel& kernel = generate_kernel_unity_stride( - (W.num_cols-1)/2, (W.num_rows-1)/2, flip, overwrite); - - kernel.global_work_size(0, ocl::align_to_multiple_2d(Y.num_cols)); - kernel.global_work_size(1, ocl::align_to_multiple_2d(Y.num_rows)); - - viennacl::ocl::enqueue(kernel( - X.vcl_matrix(), cl_int(X.num_cols), cl_int(X.num_rows), cl_int(X.offset), - W.vcl_matrix(), cl_int(W.offset), - Y.vcl_matrix(), cl_int(Y.offset))); - } - else - { - viennacl::ocl::kernel& kernel = generate_kernel_arbitrary_stride( - (W.num_cols-1)/2, (W.num_rows-1)/2, flip, overwrite); - - kernel.global_work_size(0, ocl::align_to_multiple_2d(Y.num_cols)); - kernel.global_work_size(1, ocl::align_to_multiple_2d(Y.num_rows)); - - viennacl::ocl::enqueue(kernel( - X.vcl_matrix(), cl_int(X.num_cols), cl_int(X.num_rows), cl_int(X.offset), - W.vcl_matrix(), cl_int(W.offset), - Y.vcl_matrix(), cl_int(Y.offset), - cl_int(stride_x), cl_int(stride_y))); - } - } -}; - -#endif // HAVE_VIENNACL - -} - -} - -} -#endif // CONVOLVE_H_ diff --git a/src/shogun/mathematics/linalg/internal/modules/Core.h b/src/shogun/mathematics/linalg/internal/modules/Core.h index f46528d5c34..62136ab15c0 100644 --- a/src/shogun/mathematics/linalg/internal/modules/Core.h +++ b/src/shogun/mathematics/linalg/internal/modules/Core.h @@ -35,7 +35,6 @@ #include #include #include -#include namespace shogun { @@ -139,28 +138,6 @@ void elementwise_square(Matrix m, ResultMatrix result) implementation::elementwise_square::compute(m,result); } -/** Computes the 2D convolution of X with W - * - * NOTE: For the ViennaCL backend, the size of W (number of bytes) must not exceed - * [CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE](http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clGetDeviceInfo.html). - * - * @param X Input image - * @param W Filter coefficients. The dimensions of the matrix must be odd-numbered. - * @param Y Output image of the same size as the input image, as the borders - * of the input image are implicitly padded with zeros during the computation - * @param flip If true the filter coefficients are flipped, performing cross-correlation - * instead of convolution - * @param overwrite If true, the values in Y are overwritten with result of the - * computation. Otherwise, the result is added to the existing values in Y. - * @param stride_x Stride in the x (column) direction - * @param stride_y Stride in the y (row) direction - */ -template ::backend,class Matrix> -void convolve(Matrix X, Matrix W, Matrix Y, bool flip = false, - bool overwrite=true, int32_t stride_x=1, int32_t stride_y=1) -{ - implementation::convolve::compute(X, W, Y, flip, overwrite, stride_x, stride_y); -} #endif // HAVE_LINALG_LIB } diff --git a/tests/unit/mathematics/linalg/Convolve_unittest.cc b/tests/unit/mathematics/linalg/Convolve_unittest.cc deleted file mode 100644 index 5e49c3610c4..00000000000 --- a/tests/unit/mathematics/linalg/Convolve_unittest.cc +++ /dev/null @@ -1,237 +0,0 @@ -/* - * Copyright (c) The Shogun Machine Learning Toolbox - * Written (w) 2014 Khaled Nasr - * All rights reserved. - * - * Redistribution and use in source and binary forms, with or without - * modification, are permitted provided that the following conditions are met: - * - * 1. Redistributions of source code must retain the above copyright notice, this - * list of conditions and the following disclaimer. - * 2. Redistributions in binary form must reproduce the above copyright notice, - * this list of conditions and the following disclaimer in the documentation - * and/or other materials provided with the distribution. - * - * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND - * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED - * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE - * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR - * ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES - * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; - * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND - * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT - * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS - * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - * - * The views and conclusions contained in the software and documentation are those - * of the authors and should not be interpreted as representing official policies, - * either expressed or implied, of the Shogun Development Team. - */ - -#include - -#ifdef HAVE_LINALG_LIB -#include -#include -#include -#include - -#include - -#ifdef HAVE_VIENNACL -#include -#include -#endif - -using namespace shogun; - -TEST(Convolve, eigen3_backend) -{ - const int32_t w = 6; - const int32_t h = 5; - const int32_t sx = 1; - const int32_t sy = 1; - - SGMatrix X(h,w); - SGMatrix W(3,3); - SGMatrix Y(h/sy,w/sx); - - for (int32_t i=0; i(X, W, Y, false, true, sx, sy); - - // generated with scipy.signal.convolve2d - float64_t ref[] = { - 2.00000, 6.50000, 10.25000, 14.00000, 14.00000, 13.50000, - 30.00000, 39.00000, 48.00000, 42.00000, 39.75000, 75.00000, - 84.00000, 93.00000, 75.75000, 66.00000, 120.00000, 129.00000, - 138.00000, 109.50000, 92.25000, 165.00000, 174.00000, 183.00000, - 143.25000, 111.00000, 187.25000, 195.50000, 203.75000, 152.00000}; - - for (int32_t i=0; i X(h,w); - SGMatrix W(3,3); - SGMatrix Y(h/sy,w/sx); - - for (int32_t i=0; i(X, W, Y, true, true, sx, sy); - - // generated with scipy.signal.convolve2d - float64_t ref[] = { - 22.00000, 35.50000, 43.75000, 52.00000, 34.00000, 52.50000, - 78.00000, 87.00000, 96.00000, 60.00000, 86.25000, 123.00000, - 132.00000, 141.00000, 86.25000, 120.00000, 168.00000, 177.00000, - 186.00000, 112.50000, 153.75000, 213.00000, 222.00000, 231.00000, - 138.75000, 73.00000, 94.75000, 98.50000, 102.25000, 56.00000}; - - for (int32_t i=0; i X(h,w); - SGMatrix W(3,3); - SGMatrix Y(h/sy,w/sx); - - for (int32_t i=0; i(X, W, Y, false, true, sx, sy); - - // generated with scipy.signal.convolve2d - float64_t ref[] = { - 3.75000, 19.25000, 30.50000, 41.75000, 92.25000, 187.50000, - 214.50000, 241.50000, 218.25000, 403.50000, 430.50000, 457.50000, - 344.25000, 619.50000, 646.50000, 673.50000, 470.25000, 835.50000, - 862.50000, 889.50000}; - - for (int32_t i=0; i X(h,w); - CGPUMatrix W(3,3); - CGPUMatrix Y(h/sy,w/sx); - - for (int32_t i=0; i(X, W, Y, false, true, sx, sy); - - // generated with scipy.signal.convolve2d - float64_t ref[] = { - 2.00000, 6.50000, 10.25000, 14.00000, 14.00000, 13.50000, - 30.00000, 39.00000, 48.00000, 42.00000, 39.75000, 75.00000, - 84.00000, 93.00000, 75.75000, 66.00000, 120.00000, 129.00000, - 138.00000, 109.50000, 92.25000, 165.00000, 174.00000, 183.00000, - 143.25000, 111.00000, 187.25000, 195.50000, 203.75000, 152.00000}; - - for (int32_t i=0; i X(h,w); - CGPUMatrix W(3,3); - CGPUMatrix Y(h/sy,w/sx); - - for (int32_t i=0; i(X, W, Y, true, true, sx, sy); - - // generated with scipy.signal.convolve2d - float64_t ref[] = { - 22.00000, 35.50000, 43.75000, 52.00000, 34.00000, 52.50000, - 78.00000, 87.00000, 96.00000, 60.00000, 86.25000, 123.00000, - 132.00000, 141.00000, 86.25000, 120.00000, 168.00000, 177.00000, - 186.00000, 112.50000, 153.75000, 213.00000, 222.00000, 231.00000, - 138.75000, 73.00000, 94.75000, 98.50000, 102.25000, 56.00000}; - - for (int32_t i=0; i X(h,w); - CGPUMatrix W(3,3); - CGPUMatrix Y(h/sy,w/sx); - - for (int32_t i=0; i(X, W, Y, false, true, sx, sy); - - // generated with scipy.signal.convolve2d - float64_t ref[] = { - 3.75000, 19.25000, 30.50000, 41.75000, 92.25000, 187.50000, - 214.50000, 241.50000, 218.25000, 403.50000, 430.50000, 457.50000, - 344.25000, 619.50000, 646.50000, 673.50000, 470.25000, 835.50000, - 862.50000, 889.50000}; - - for (int32_t i=0; i