62 namespace implementation
70 template <enum Backend,
class Matrix>
74 typedef typename Matrix::Scalar
T;
83 static T
compute(Matrix m,
bool no_diag);
100 template <enum Backend,
class Matrix>
104 typedef typename Matrix::Scalar
T;
113 static T
compute(Matrix m,
bool no_diag);
130 template <enum Backend,
class Matrix>
134 typedef typename Matrix::Scalar
T;
179 template <enum Backend,
class Matrix>
183 typedef typename Matrix::Scalar
T;
228 template <
class Matrix>
229 struct sum<Backend::EIGEN3,Matrix>
232 typedef typename Matrix::Scalar
T;
252 sum-=m.diagonal().sum();
268 Eigen::Block< Eigen::Map<MatrixXt> > b_eigen = map.block(
269 b.m_row_begin, b.m_col_begin,
270 b.m_row_size, b.m_col_size);
276 sum-=b_eigen.diagonal().sum();
286 template <
class Matrix>
290 typedef typename Matrix::Scalar
T;
306 REQUIRE(m.rows()==m.cols(),
"Matrix is not square!\n");
310 const MatrixXt& m_upper=m.template triangularView<Eigen::StrictlyUpper>();
318 sum+=m.diagonal().sum();
334 const MatrixXt& m=map.template
block(b.m_row_begin, b.m_col_begin,
335 b.m_row_size, b.m_col_size);
337 REQUIRE(m.rows()==m.cols(),
"Matrix is not square!\n");
341 const MatrixXt& m_upper=m.template triangularView<Eigen::StrictlyUpper>();
349 sum+=m.diagonal().sum();
359 template <
class Matrix>
363 typedef typename Matrix::Scalar
T;
415 r = m.colwise().sum();
420 index_t len_major_diag=m.rows() < m.cols() ? m.rows() : m.cols();
421 for (
index_t i=0; i<len_major_diag; ++i)
439 Eigen::Block< Eigen::Map<MatrixXt> > m = map.block(
440 b.m_row_begin, b.m_col_begin,
441 b.m_row_size, b.m_col_size);
443 r = m.colwise().sum();
448 index_t len_major_diag=m.rows() < m.cols() ? m.rows() : m.cols();
449 for (
index_t i=0; i<len_major_diag; ++i)
459 template <
class Matrix>
463 typedef typename Matrix::Scalar
T;
515 r = m.rowwise().sum();
520 index_t len_major_diag=m.rows() < m.cols() ? m.rows() : m.cols();
521 for (
index_t i=0; i<len_major_diag; ++i)
539 Eigen::Block< Eigen::Map<MatrixXt> > m = map.block(
540 b.m_row_begin, b.m_col_begin,
541 b.m_row_size, b.m_col_size);
543 r = m.rowwise().sum();
548 index_t len_major_diag=m.rows() < m.cols() ? m.rows() : m.cols();
549 for (
index_t i=0; i<len_major_diag; ++i)
555 #endif // HAVE_EIGEN3
562 template <
class Matrix>
563 struct sum<Backend::VIENNACL,Matrix>
566 typedef typename Matrix::Scalar
T;
570 static viennacl::ocl::kernel& generate_kernel(
bool no_diag)
572 std::string kernel_name =
"sum_" + ocl::get_type_string<T>();
573 if (no_diag) kernel_name.append(
"_no_diag");
575 if (ocl::kernel_exists(kernel_name))
576 return ocl::get_kernel(kernel_name);
578 std::string source = ocl::generate_kernel_preamble<T>(kernel_name);
579 if (no_diag) source.append(
"#define NO_DIAG\n");
583 __kernel void KERNEL_NAME(
584 __global DATATYPE* mat, int nrows, int ncols, int offset,
585 __global DATATYPE* result)
587 __local DATATYPE buffer[WORK_GROUP_SIZE_1D];
588 int size = nrows*ncols;
590 int local_id = get_local_id(0);
592 DATATYPE thread_sum = 0;
593 for (int i=local_id; i<size; i+=WORK_GROUP_SIZE_1D)
596 if (!(i/nrows == i%nrows))
598 thread_sum += mat[i+offset];
601 buffer[local_id] = thread_sum;
603 for (int j = WORK_GROUP_SIZE_1D/2; j > 0; j = j>>1)
605 barrier(CLK_LOCAL_MEM_FENCE);
607 buffer[local_id] += buffer[local_id + j];
610 barrier(CLK_LOCAL_MEM_FENCE);
612 if (get_global_id(0)==0)
618 viennacl::ocl::kernel& kernel = ocl::compile_kernel(kernel_name, source);
620 kernel.local_work_size(0, OCL_WORK_GROUP_SIZE_1D);
621 kernel.global_work_size(0, OCL_WORK_GROUP_SIZE_1D);
633 static T
compute(CGPUMatrix<T> mat,
bool no_diag)
635 viennacl::ocl::kernel& kernel = generate_kernel<T>(no_diag);
637 CGPUVector<T> result(1);
639 viennacl::ocl::enqueue(kernel(mat.vcl_matrix(),
640 cl_int(mat.num_rows), cl_int(mat.num_cols), cl_int(mat.offset),
641 result.vcl_vector()));
653 static T
compute(Block<CGPUMatrix<T> > b,
bool no_diag)
655 SG_SERROR(
"The operation sum() on a matrix block is currently not supported\n");
664 template <
class Matrix>
665 struct sum_symmetric<Backend::VIENNACL,Matrix>
668 typedef typename Matrix::Scalar
T;
677 static T
compute(CGPUMatrix<T> mat,
bool no_diag)
679 return sum<Backend::VIENNACL, CGPUMatrix<T> >
::compute(mat, no_diag);
689 static T
compute(Block<CGPUMatrix<T> > b,
bool no_diag)
691 SG_SERROR(
"The operation sum_symmetric() on a matrix block is currently not supported\n");
700 template <
class Matrix>
701 struct colwise_sum<Backend::VIENNACL,Matrix>
704 typedef typename Matrix::Scalar
T;
707 typedef CGPUVector<T> ReturnType;
711 static viennacl::ocl::kernel& generate_kernel(
bool no_diag)
713 std::string kernel_name =
"colwise_sum_" + ocl::get_type_string<T>();
714 if (no_diag) kernel_name.append(
"_no_diag");
716 if (ocl::kernel_exists(kernel_name))
717 return ocl::get_kernel(kernel_name);
719 std::string source = ocl::generate_kernel_preamble<T>(kernel_name);
720 if (no_diag) source.append(
"#define NO_DIAG\n");
724 __kernel void KERNEL_NAME(
725 __global DATATYPE* mat, int nrows, int ncols, int offset,
726 __global DATATYPE* result, int result_offset)
728 int j = get_global_id(0);
734 for (int i=0; i<nrows; i++)
739 sum += mat[offset+i+j*nrows];
742 result[j+result_offset] = sum;
747 viennacl::ocl::kernel& kernel = ocl::compile_kernel(kernel_name, source);
749 kernel.local_work_size(0, OCL_WORK_GROUP_SIZE_1D);
761 static CGPUVector<T>
compute(CGPUMatrix<T> m,
bool no_diag)
763 CGPUVector<T> result(m.num_cols);
776 static CGPUVector<T>
compute(Block<CGPUMatrix<T> > b,
bool no_diag)
778 SG_SERROR(
"The operation colwise_sum() on a matrix block is currently not supported\n");
779 return CGPUVector<T>();
789 static void compute(CGPUMatrix<T> mat, CGPUVector<T> result,
bool no_diag)
791 viennacl::ocl::kernel& kernel = generate_kernel<T>(no_diag);
792 kernel.global_work_size(0, ocl::align_to_multiple_1d(mat.num_cols));
794 viennacl::ocl::enqueue(kernel(mat.vcl_matrix(),
795 cl_int(mat.num_rows), cl_int(mat.num_cols), cl_int(mat.offset),
796 result.vcl_vector(), cl_int(result.offset)));
807 static void compute(Block<CGPUMatrix<T> > b, CGPUVector<T> result,
bool no_diag)
809 SG_SERROR(
"The operation colwise_sum() on a matrix block is currently not supported\n");
817 template <
class Matrix>
818 struct rowwise_sum<Backend::VIENNACL,Matrix>
821 typedef typename Matrix::Scalar
T;
824 typedef CGPUVector<T> ReturnType;
828 static viennacl::ocl::kernel& generate_kernel(
bool no_diag)
830 std::string kernel_name =
"rowwise_sum_" + ocl::get_type_string<T>();
831 if (no_diag) kernel_name.append(
"_no_diag");
833 if (ocl::kernel_exists(kernel_name))
834 return ocl::get_kernel(kernel_name);
836 std::string source = ocl::generate_kernel_preamble<T>(kernel_name);
837 if (no_diag) source.append(
"#define NO_DIAG\n");
841 __kernel void KERNEL_NAME(
842 __global DATATYPE* mat, int nrows, int ncols, int offset,
843 __global DATATYPE* result, int result_offset)
845 int i = get_global_id(0);
851 for (int j=0; j<ncols; j++)
856 sum += mat[offset+i+j*nrows];
859 result[i+result_offset] = sum;
864 viennacl::ocl::kernel& kernel = ocl::compile_kernel(kernel_name, source);
866 kernel.local_work_size(0, OCL_WORK_GROUP_SIZE_1D);
878 static CGPUVector<T>
compute(CGPUMatrix<T> m,
bool no_diag)
880 CGPUVector<T> result(m.num_rows);
893 static CGPUVector<T>
compute(Block<CGPUMatrix<T> > b,
bool no_diag)
895 SG_SERROR(
"The operation rowwise_sum() on a matrix block is currently not supported\n");
896 return CGPUVector<T>();
906 static void compute(CGPUMatrix<T> mat, CGPUVector<T> result,
bool no_diag)
908 viennacl::ocl::kernel& kernel = generate_kernel<T>(no_diag);
909 kernel.global_work_size(0, ocl::align_to_multiple_1d(mat.num_rows));
911 viennacl::ocl::enqueue(kernel(mat.vcl_matrix(),
912 cl_int(mat.num_rows), cl_int(mat.num_cols), cl_int(mat.offset),
913 result.vcl_vector(), cl_int(result.offset)));
924 static void compute(Block<CGPUMatrix<T> > b, CGPUVector<T> result,
bool no_diag)
926 SG_SERROR(
"The operation rowwise_sum() on a matrix block is currently not supported\n");
930 #endif // HAVE_VIENNACL
937 #endif // SUM_IMPL_H_
Generic class colwise_sum which provides a static compute method. This class is specialized for diffe...
static SGVector< T > compute(Matrix m, bool no_diag)
static SGVector< T > compute(SGMatrix< T > m, bool no_diag)
static void compute(SGMatrix< T > mat, SGVector< T > result, bool no_diag)
Eigen::Matrix< T, Eigen::Dynamic, Eigen::Dynamic > MatrixXt
Generic class rowwise_sum which provides a static compute method. This class is specialized for diffe...
static T compute(Matrix m, bool no_diag)
static void compute(Block< SGMatrix< T > > b, SGVector< T > result, bool no_diag)
static T compute(SGMatrix< T > mat, bool no_diag)
static T compute(SGMatrix< T > mat, bool no_diag)
Generic class Block which wraps a matrix class and contains block specific information, providing a uniform way to deal with matrix blocks for all supported backend matrices.
static void compute(SGMatrix< T > mat, SGVector< T > result, bool no_diag)
Eigen::Matrix< T, Eigen::Dynamic, 1 > VectorXt
Generic class sum which provides a static compute method. This class is specialized for different typ...
static T compute(Matrix m, bool no_diag)
static SGVector< T > compute(Block< SGMatrix< T > > b, bool no_diag)
static SGVector< T > compute(Matrix m, bool no_diag)
all of classes and functions are contained in the shogun namespace
Eigen::Matrix< T, Eigen::Dynamic, Eigen::Dynamic > MatrixXt
static T compute(Block< SGMatrix< T > > b, bool no_diag)
static SGVector< T > compute(SGMatrix< T > m, bool no_diag)
Eigen::Matrix< T, Eigen::Dynamic, 1 > VectorXt
static T compute(Block< SGMatrix< T > > b, bool no_diag)
static void compute(Block< SGMatrix< T > > b, SGVector< T > result, bool no_diag)
static SGVector< T > compute(Block< SGMatrix< T > > b, bool no_diag)
Eigen::Matrix< T, Eigen::Dynamic, Eigen::Dynamic > MatrixXt
Eigen::Matrix< T, Eigen::Dynamic, Eigen::Dynamic > MatrixXt
Block< Matrix > block(Matrix matrix, index_t row_begin, index_t col_begin, index_t row_size, index_t col_size)
Generic class sum symmetric which provides a static compute method. This class is specialized for dif...