From 1c79aecc1bc70095d8e508d40438b86bab758bbf Mon Sep 17 00:00:00 2001 From: lambday Date: Fri, 20 Feb 2015 19:36:06 +0530 Subject: [PATCH] fix issue #2729 --- src/shogun/kernel/GaussianARDKernel.cpp | 4 +- src/shogun/kernel/GaussianARDKernel.h | 8 +- src/shogun/kernel/LinearARDKernel.cpp | 8 +- src/shogun/kernel/LinearARDKernel.h | 8 +- .../linalg/internal/implementation/Add.h | 80 ++++++++------- .../linalg/internal/implementation/Max.h | 97 ++++++++++--------- .../internal/implementation/VectorSum.h | 16 +-- .../linalg/internal/modules/Core.h | 43 ++++---- .../linalg/internal/modules/Redux.h | 43 ++++---- 9 files changed, 166 insertions(+), 141 deletions(-) diff --git a/src/shogun/kernel/GaussianARDKernel.cpp b/src/shogun/kernel/GaussianARDKernel.cpp index 41f4f21e40a..1e71fbb91ab 100644 --- a/src/shogun/kernel/GaussianARDKernel.cpp +++ b/src/shogun/kernel/GaussianARDKernel.cpp @@ -29,7 +29,7 @@ void CGaussianARDKernel::init() SG_ADD(&m_width, "width", "Kernel width", MS_AVAILABLE, GRADIENT_AVAILABLE); } -#ifdef HAVE_CXX11 +#ifdef HAVE_LINALG_LIB #include CGaussianARDKernel::CGaussianARDKernel(int32_t size, float64_t width) : CLinearARDKernel(size) @@ -122,4 +122,4 @@ float64_t CGaussianARDKernel::distance(int32_t idx_a, int32_t idx_b) float64_t result=compute_helper(avec, avec); return result/m_width; } -#endif /* HAVE_CXX11 */ +#endif /* HAVE_LINALG_LIB */ diff --git a/src/shogun/kernel/GaussianARDKernel.h b/src/shogun/kernel/GaussianARDKernel.h index a7b96b84ffb..a0c6cc00cc6 100644 --- a/src/shogun/kernel/GaussianARDKernel.h +++ b/src/shogun/kernel/GaussianARDKernel.h @@ -33,7 +33,7 @@ namespace shogun * * where \f$\tau\f$ is the kernel width. * - * There are three variants based on \f$||\cdot||\f$. + * There are three variants based on \f$||\cdot||\f$. * The default case is * \f$\sum_{i=1}^{p}{{[\lambda \times ({\bf x_i}-{\bf y_i})] }^2}\f$ * where \f$\lambda\f$ is a scalar and \f$p\f$ is # of features @@ -50,7 +50,7 @@ namespace shogun * The last case is * \f$({\bf x}-{\bf y})^T \Lambda^T \Lambda ({\bf x}-{\bf y})\f$ * where \f$\Lambda\f$ is a \f$d\f$-by-\f$p\f$ matrix, - * \f$p\f$ is # of features and \f$d\f$ can be \f$ d \ge p\f$ or \f$ d \le p\f$ + * \f$p\f$ is # of features and \f$d\f$ can be \f$ d \ge p\f$ or \f$ d \le p\f$ * To use this case, please call set_matrix_weights(\f$\Lambda\f$) * where \f$\Lambda\f$ is a \f$d\f$-by-\f$p\f$ matrix * @@ -91,7 +91,7 @@ class CGaussianARDKernel: public CLinearARDKernel /** kernel width */ float64_t m_width; -#ifdef HAVE_CXX11 +#ifdef HAVE_LINALG_LIB public: /** constructor * @@ -160,7 +160,7 @@ class CGaussianARDKernel: public CLinearARDKernel */ virtual float64_t distance(int32_t idx_a, int32_t idx_b); -#endif /* HAVE_CXX11 */ +#endif /* HAVE_LINALG_LIB */ }; } #endif /* _GAUSSIANARDKERNEL_H_ */ diff --git a/src/shogun/kernel/LinearARDKernel.cpp b/src/shogun/kernel/LinearARDKernel.cpp index 6f77d386ef7..ba4736f8fbb 100644 --- a/src/shogun/kernel/LinearARDKernel.cpp +++ b/src/shogun/kernel/LinearARDKernel.cpp @@ -32,7 +32,7 @@ void CLinearARDKernel::init() GRADIENT_AVAILABLE); } -#ifdef HAVE_CXX11 +#ifdef HAVE_LINALG_LIB #include CLinearARDKernel::CLinearARDKernel(int32_t size) : CDotKernel(size) @@ -78,7 +78,7 @@ SGMatrix CLinearARDKernel::compute_right_product(SGVectorr right=SGMatrix(right_vec.vector,right_vec.vlen,1,false); scalar_weight*=m_weights[0]; } - else + else { right=SGMatrix (m_weights.num_rows,1); @@ -110,7 +110,7 @@ float64_t CLinearARDKernel::compute_helper(SGVector avec, SGVector(avec.vector,1,avec.vlen,false); scalar_weight=m_weights[0]; } - else + else { left=SGMatrix(1,m_weights.num_rows); @@ -288,4 +288,4 @@ void CLinearARDKernel::set_matrix_weights(SGMatrix weights) { set_weights(weights); } -#endif //HAVE_CXX11 +#endif //HAVE_LINALG_LIB diff --git a/src/shogun/kernel/LinearARDKernel.h b/src/shogun/kernel/LinearARDKernel.h index 1c0000e1aab..8c3a9f3cce9 100644 --- a/src/shogun/kernel/LinearARDKernel.h +++ b/src/shogun/kernel/LinearARDKernel.h @@ -36,7 +36,7 @@ enum EARDKernelType * k({\bf x},{\bf y})= \frac{||{\bf x}-{\bf y}||} * \f] * - * There are three variants based on \f$||\cdot||\f$. + * There are three variants based on \f$||\cdot||\f$. * The default case is * \f$\sum_{i=1}^{p}{{[\lambda \times ({\bf x_i}-{\bf y_i})] }^2}\f$ * where \f$\lambda\f$ is a scalar and \f$p\f$ is # of features @@ -53,7 +53,7 @@ enum EARDKernelType * The last case is * \f$({\bf x}-{\bf y})^T \Lambda^T \Lambda ({\bf x}-{\bf y})\f$ * where \f$\Lambda\f$ is a \f$d\f$-by-\f$p\f$ matrix, - * \f$p\f$ is # of features and \f$d\f$ can be \f$ d \ge p\f$ or \f$ d \le p\f$ + * \f$p\f$ is # of features and \f$d\f$ can be \f$ d \ge p\f$ or \f$ d \le p\f$ * To use this case, please call set_matrix_weights(\f$\Lambda\f$), * where \f$\Lambda\f$ is a \f$d\f$-by-\f$p\f$ matrix * @@ -105,7 +105,7 @@ class CLinearARDKernel: public CDotKernel /** type of ARD kernel */ EARDKernelType m_ARD_type; -#ifdef HAVE_CXX11 +#ifdef HAVE_LINALG_LIB public: /** constructor * @@ -219,7 +219,7 @@ class CLinearARDKernel: public CDotKernel */ virtual SGMatrix compute_right_product(SGVectorright_vec, float64_t & scalar_weight); -#endif /* HAVE_CXX11 */ +#endif /* HAVE_LINALG_LIB */ }; } #endif /* _LINEARARDKERNEL_H_ */ diff --git a/src/shogun/mathematics/linalg/internal/implementation/Add.h b/src/shogun/mathematics/linalg/internal/implementation/Add.h index 92f05d1bae1..002e52070c9 100644 --- a/src/shogun/mathematics/linalg/internal/implementation/Add.h +++ b/src/shogun/mathematics/linalg/internal/implementation/Add.h @@ -64,7 +64,7 @@ struct add { /** Scalar type */ typedef typename Matrix::Scalar T; - + /** * Performs the operation C = alpha*A + beta*B. Works for both matrices and vectors * @param A first matrix @@ -77,11 +77,12 @@ struct add }; /** - * @brief Specialization of add for the Native backend + * @brief Partial specialization of add for the Native backend */ -template <> template +template struct add { + /** Scalar type */ typedef typename Matrix::Scalar T; /** @@ -101,7 +102,7 @@ struct add "Matrices should have same number of columns!\n"); compute(A.matrix, B.matrix, C.matrix, alpha, beta, A.num_rows*A.num_cols); } - + /** * Performs the operation C = alpha*A + beta*B. * @param A first vector @@ -126,27 +127,31 @@ struct add * @param beta constant to be multiplied by the second vector * @param len length of the vectors/matrices */ - static void compute(T* A, T* B, T* C, - T alpha, T beta, index_t len) + static void compute(T* A, T* B, T* C, T alpha, T beta, index_t len) { for (int32_t i=0; i template +template struct add { + /** Scalar type */ typedef typename Matrix::Scalar T; + + /** Eigen3 matrix type */ typedef Eigen::Matrix MatrixXt; + + /** Eigen3 vector type */ typedef Eigen::Matrix VectorXt; - - /** + + /** * Performs the operation C = alpha*A + beta*B using Eigen3 * @param A first matrix * @param B second matrix @@ -154,17 +159,17 @@ struct add * @param alpha constant to be multiplied by the first matrix * @param beta constant to be multiplied by the second matrix */ - static void compute(SGMatrix A, SGMatrix B, SGMatrix C, + static void compute(SGMatrix A, SGMatrix B, SGMatrix C, T alpha, T beta) { - Eigen::Map A_eig = A; - Eigen::Map B_eig = B; - Eigen::Map C_eig = C; - - C_eig = alpha*A_eig + beta*B_eig; + Eigen::Map A_eig=A; + Eigen::Map B_eig=B; + Eigen::Map C_eig=C; + + C_eig=alpha*A_eig+beta*B_eig; } - - /** + + /** * Performs the operation C = alpha*A + beta*B using Eigen3 * @param A first vector * @param B second vector @@ -172,29 +177,30 @@ struct add * @param alpha constant to be multiplied by the first vector * @param beta constant to be multiplied by the second vector */ - static void compute(SGVector A, SGVector B, SGVector C, + static void compute(SGVector A, SGVector B, SGVector C, T alpha, T beta) { - Eigen::Map A_eig = A; - Eigen::Map B_eig = B; - Eigen::Map C_eig = C; - - C_eig = alpha*A_eig + beta*B_eig; + Eigen::Map A_eig=A; + Eigen::Map B_eig=B; + Eigen::Map C_eig=C; + + C_eig=alpha*A_eig+beta*B_eig; } }; #endif // HAVE_EIGEN3 #ifdef HAVE_VIENNACL -/** - * @brief Specialization of add for the ViennaCL backend +/** + * @brief Partial specialization of add for the ViennaCL backend */ -template <> template +template struct add { + /** Scalar type */ typedef typename Matrix::Scalar T; - - /** + + /** * Performs the operation C = alpha*A + beta*B using Viennacl * @param A first matrix * @param B second matrix @@ -202,13 +208,13 @@ struct add * @param alpha constant to be multiplied by the first matrix * @param beta constant to be multiplied by the second matrix */ - static void compute(CGPUMatrix A, CGPUMatrix B, CGPUMatrix C, + static void compute(CGPUMatrix A, CGPUMatrix B, CGPUMatrix C, T alpha, T beta) { - C.vcl_matrix() = alpha*A.vcl_matrix() + beta*B.vcl_matrix(); + C.vcl_matrix()=alpha*A.vcl_matrix()+beta*B.vcl_matrix(); } - - /** + + /** * Performs the operation C = alpha*A + beta*B using Viennacl * @param A first vector * @param B second vector @@ -216,10 +222,10 @@ struct add * @param alpha constant to be multiplied by the first vector * @param beta constant to be multiplied by the second vector */ - static void compute(CGPUVector A, CGPUVector B, CGPUVector C, + static void compute(CGPUVector A, CGPUVector B, CGPUVector C, T alpha, T beta) { - C.vcl_vector() = alpha*A.vcl_vector() + beta*B.vcl_vector(); + C.vcl_vector()=alpha*A.vcl_vector()+beta*B.vcl_vector(); } }; diff --git a/src/shogun/mathematics/linalg/internal/implementation/Max.h b/src/shogun/mathematics/linalg/internal/implementation/Max.h index 6a028aeca5a..247d5ed91a5 100644 --- a/src/shogun/mathematics/linalg/internal/implementation/Max.h +++ b/src/shogun/mathematics/linalg/internal/implementation/Max.h @@ -57,8 +57,8 @@ namespace linalg namespace implementation { -/** - * @brief Generic class which is specialized for different backends to perform +/** + * @brief Generic class which is specialized for different backends to perform * the max operation */ template @@ -66,8 +66,8 @@ struct max { /** Scalar type */ typedef typename Matrix::Scalar T; - - /** + + /** * Returns the largest element in a matrix or a vector. * @param m input matrix or vector * @return largest value in the input matrix or vector @@ -78,9 +78,10 @@ struct max /** * @brief Specialization of add for the Native backend */ -template <> template +template struct max { + /** Scalar type */ typedef typename Matrix::Scalar T; /** @@ -93,7 +94,7 @@ struct max REQUIRE(mat.num_cols*mat.num_rows > 0, "Matrix can not be empty!\n"); return compute(mat.matrix, mat.num_cols*mat.num_rows); } - + /** * Returns the largest element in a vector. * @param vec input vector @@ -122,13 +123,18 @@ struct max /** * @brief Specialization of max for the Eigen3 backend */ -template <> template +template struct max { + /** Scalar type */ typedef typename Matrix::Scalar T; + + /** Eigen3 matrix type */ typedef Eigen::Matrix MatrixXt; + + /** Eigen3 vector type */ typedef Eigen::Matrix VectorXt; - + /** * Returns the largest element in a matrix * @param mat input matrix @@ -137,10 +143,10 @@ struct max static T compute(SGMatrix mat) { Eigen::Map m = mat; - + return m.maxCoeff(); } - + /** * Returns the largest element in a vector * @param vec input vector @@ -149,7 +155,7 @@ struct max static T compute(SGVector vec) { Eigen::Map v = vec; - + return v.maxCoeff(); } }; @@ -160,65 +166,66 @@ struct max /** * @brief Specialization of max for the ViennaCL backend */ -template <> template +template struct max { + /** Scalar type */ typedef typename Matrix::Scalar T; - + /** Generates the computation kernel */ template static viennacl::ocl::kernel& generate_kernel() { std::string kernel_name = "max_" + ocl::get_type_string(); - + if (ocl::kernel_exists(kernel_name)) return ocl::get_kernel(kernel_name); - + std::string source = ocl::generate_kernel_preamble(kernel_name); - + source.append( R"( __kernel void KERNEL_NAME( - __global DATATYPE* vec, int size, int offset, + __global DATATYPE* vec, int size, int offset, __global DATATYPE* result) { __local DATATYPE buffer[WORK_GROUP_SIZE_1D]; - + int local_id = get_local_id(0); - + DATATYPE thread_max = -INFINITY; for (int i=local_id; i 0; j = j>>1) - { - barrier(CLK_LOCAL_MEM_FENCE); - if (local_id < j) - buffer[local_id] = max(buffer[local_id], buffer[local_id + j]); - } - + + for (int j = WORK_GROUP_SIZE_1D/2; j > 0; j = j>>1) + { + barrier(CLK_LOCAL_MEM_FENCE); + if (local_id < j) + buffer[local_id] = max(buffer[local_id], buffer[local_id + j]); + } + barrier(CLK_LOCAL_MEM_FENCE); - + if (get_global_id(0)==0) *result = buffer[0]; } )" ); - + viennacl::ocl::kernel& kernel = ocl::compile_kernel(kernel_name, source); - + kernel.local_work_size(0, OCL_WORK_GROUP_SIZE_1D); kernel.global_work_size(0, OCL_WORK_GROUP_SIZE_1D); - + return kernel; } - - /** + + /** * Returns the largest element in a matrix * @param mat input matrix * @return largest value in the matrix @@ -226,16 +233,16 @@ struct max static T compute(CGPUMatrix mat) { viennacl::ocl::kernel& kernel = generate_kernel(); - + CGPUVector result(1); - - viennacl::ocl::enqueue(kernel(mat.vcl_matrix(), - cl_int(mat.num_rows*mat.num_cols), cl_int(mat.offset), + + viennacl::ocl::enqueue(kernel(mat.vcl_matrix(), + cl_int(mat.num_rows*mat.num_cols), cl_int(mat.offset), result.vcl_vector())); - + return result[0]; } - + /** * Returns the largest element in a vector * @param vec input vector @@ -244,13 +251,13 @@ struct max static T compute(CGPUVector vec) { viennacl::ocl::kernel& kernel = generate_kernel(); - + CGPUVector result(1); - - viennacl::ocl::enqueue(kernel(vec.vcl_vector(), - cl_int(vec.vlen), cl_int(vec.offset), + + viennacl::ocl::enqueue(kernel(vec.vcl_vector(), + cl_int(vec.vlen), cl_int(vec.offset), result.vcl_vector())); - + return result[0]; } }; diff --git a/src/shogun/mathematics/linalg/internal/implementation/VectorSum.h b/src/shogun/mathematics/linalg/internal/implementation/VectorSum.h index 8ff6e0fce43..6e5d6ca2550 100644 --- a/src/shogun/mathematics/linalg/internal/implementation/VectorSum.h +++ b/src/shogun/mathematics/linalg/internal/implementation/VectorSum.h @@ -79,15 +79,15 @@ struct vector_sum /** * @brief Specialization of generic vector_sum for the Native backend */ -template<> template +template struct vector_sum { /** Scalar type */ typedef typename Vector::Scalar T; - + /** * Method that computes the sum of SGVectors - * + * * @param vec a vector whose sum has to be computed * @return the vector sum \f$\sum_i a_i\f$ */ @@ -99,7 +99,7 @@ struct vector_sum /** * Method that computes the sum of SGVectors - * + * * @param vec a vector whose sum has to be computed * @param len the length of the vector * @return the vector sum \f$\sum_i a_i\f$ @@ -109,14 +109,15 @@ struct vector_sum return std::accumulate(vec,vec+len,0); } }; - + #ifdef HAVE_EIGEN3 /** * @brief Specialization of generic vector_sum for the Eigen3 backend */ -template <> template +template struct vector_sum { + /** Scalar type */ typedef typename Vector::Scalar T; /** @@ -139,9 +140,10 @@ struct vector_sum /** * @brief Specialization of generic vector_sum for the ViennaCL backend */ -template <> template +template struct vector_sum { + /** Scalar type */ typedef typename Vector::Scalar T; /** diff --git a/src/shogun/mathematics/linalg/internal/modules/Core.h b/src/shogun/mathematics/linalg/internal/modules/Core.h index 14ed47447e3..2eb1571cc7e 100644 --- a/src/shogun/mathematics/linalg/internal/modules/Core.h +++ b/src/shogun/mathematics/linalg/internal/modules/Core.h @@ -32,8 +32,6 @@ #ifndef CORE_H_ #define CORE_H_ -#ifdef HAVE_LINALG_LIB - #include #include #include @@ -47,6 +45,23 @@ namespace shogun namespace linalg { +/** Performs the operation \f$C = \alpha A + \beta B\f$. + * Works for both matrices and vectors. + * + * @param A First matrix + * @param B Second matrix + * @param C Result of the operation + * @param alpha scaling parameter for first matrix + * @param beta scaling parameter for second matrix + */ +template ::backend,class Matrix> +void add(Matrix A, Matrix B, Matrix C, typename Matrix::Scalar alpha=1.0, + typename Matrix::Scalar beta=1.0) +{ + implementation::add::compute(A, B, C, alpha, beta); +} + +#ifdef HAVE_LINALG_LIB /** Performs matrix multiplication * * @param A First matrix @@ -64,14 +79,6 @@ void matrix_product(Matrix A, Matrix B, Matrix C, implementation::matrix_product::compute(A, B, C, transpose_A, transpose_B, overwrite); } -/** Performs the operation C = alpha*A + beta*B. Works for both matrices and vectors */ -template ::backend,class Matrix> -void add(Matrix A, Matrix B, Matrix C, - typename Matrix::Scalar alpha=1.0, typename Matrix::Scalar beta=1.0) -{ - implementation::add::compute(A, B, C, alpha, beta); -} - /** Performs the operation C = alpha*A - beta*B. Works for both matrices and vectors */ template ::backend,class Matrix> void subtract(Matrix A, Matrix B, Matrix C, @@ -122,17 +129,17 @@ void elementwise_square(Matrix m, ResultMatrix 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). - * + * + * 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 + * @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 + * @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 + * @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 @@ -143,9 +150,9 @@ void convolve(Matrix X, Matrix W, Matrix Y, bool flip = false, { implementation::convolve::compute(X, W, Y, flip, overwrite, stride_x, stride_y); } +#endif // HAVE_LINALG_LIB } } -#endif // HAVE_LINALG_LIB #endif // CORE_H_ diff --git a/src/shogun/mathematics/linalg/internal/modules/Redux.h b/src/shogun/mathematics/linalg/internal/modules/Redux.h index 63aba04eeae..b15902092d3 100644 --- a/src/shogun/mathematics/linalg/internal/modules/Redux.h +++ b/src/shogun/mathematics/linalg/internal/modules/Redux.h @@ -58,6 +58,29 @@ typename Vector::Scalar dot(Vector a, Vector b) return implementation::dot::compute(a, b); } +/** Returns the largest element in a matrix or vector + * @param m the matrix or the vector + * @return the value of the largest element + */ +template ::backend, class Matrix> +typename Matrix::Scalar max(Matrix m) +{ + return implementation::max::compute(m); +} + +/** + * Wrapper method for internal implementation of vector sum of values that works + * with generic dense vectors + + * @param a vector whose sum has to be computed + * @return the vector sum \f$\sum_i a_i\f$ + */ +template ::backend, class Vector> +typename Vector::Scalar vector_sum(Vector a) +{ + return implementation::vector_sum::compute(a); +} + #ifdef HAVE_LINALG_LIB /** @@ -161,26 +184,6 @@ void rowwise_sum(Matrix m, Vector result, bool no_diag=false) implementation::rowwise_sum::compute(m, result, no_diag); } -/** - * Wrapper method for internal implementation of vector sum of values that works - * with generic dense vectors - - * @param a vector whose sum has to be computed - * @return the vector sum \f$\sum_i a_i\f$ - */ -template ::backend, class Vector> -typename Vector::Scalar vector_sum(Vector a) -{ - return implementation::vector_sum::compute(a); -} - -/** Returns the largest element in a matrix or vector */ -template ::backend, class Matrix> -typename Matrix::Scalar max(Matrix m) -{ - return implementation::max::compute(m); -} - #endif // HAVE_LINALG_LIB }