From b295add508e071b4ebe26cd1ac86d7cc3e2470fb Mon Sep 17 00:00:00 2001 From: rok-cesnovar Date: Mon, 4 Jun 2018 22:10:16 +0200 Subject: [PATCH 01/14] added GPU checks --- stan/math/gpu/err/check_matrix_gpu.hpp | 240 ++++++++++++++++++ .../kernels/check_diagonal_zeros_kernel.cl | 16 ++ stan/math/gpu/kernels/check_nan_kernel.cl | 17 ++ .../gpu/kernels/check_symmetric_kernel.cl | 19 ++ stan/math/gpu/opencl_context.hpp | 12 + test/unit/math/gpu/opencl_check_test.cpp | 85 +++++++ 6 files changed, 389 insertions(+) create mode 100644 stan/math/gpu/err/check_matrix_gpu.hpp create mode 100644 stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl create mode 100644 stan/math/gpu/kernels/check_nan_kernel.cl create mode 100644 stan/math/gpu/kernels/check_symmetric_kernel.cl create mode 100644 test/unit/math/gpu/opencl_check_test.cpp diff --git a/stan/math/gpu/err/check_matrix_gpu.hpp b/stan/math/gpu/err/check_matrix_gpu.hpp new file mode 100644 index 00000000000..e813e320962 --- /dev/null +++ b/stan/math/gpu/err/check_matrix_gpu.hpp @@ -0,0 +1,240 @@ +#ifndef STAN_MATH_PRIM_MAT_ERR_CHECK_GPU_HPP +#define STAN_MATH_PRIM_MAT_ERR_CHECK_GPU_HPP +#ifdef STAN_OPENCL +#include +#include +#include +#include +#include +#include +#include + +namespace stan { + namespace math { + /** + * Check if the specified matrix on the GPU is square. + * + * This check allows 0x0 matrices. + * + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y Matrix on the GPU to test + * + * @throw std::invalid_argument if the matrix + * is not square + */ + inline void + check_square(const char* function, + const char* name, + const matrix_gpu& y) { + check_size_match(function, + "Expecting a square matrix; rows of ", name, y.rows(), + "columns of ", name, y.cols()); + } + /** + * Check if the specified matrix on the GPU has NaN values + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y Matrix to test + * @param pos_def_check Test if matrix y is positive definite + * + * @throw std::domain_error if + * any element of the matrix is NaN. + */ + inline void + check_nan(const char* function, + const char* name, + const matrix_gpu& y, + const int pos_def_check = 0) { + if (y.size() == 0) return; + + cl::Kernel kernel_check_nan = opencl_context.get_kernel("check_nan"); + cl::CommandQueue cmd_queue = opencl_context.queue(); + cl::Context& ctx = opencl_context.context(); + int nan_flag = 0; + try { + cl::Buffer buffer_nan_flag(ctx, CL_MEM_READ_WRITE, + sizeof(int)); + + cmd_queue.enqueueWriteBuffer(buffer_nan_flag, CL_TRUE, 0, + sizeof(int), &nan_flag); + + kernel_check_nan.setArg(0, y.buffer()); + kernel_check_nan.setArg(1, y.rows()); + kernel_check_nan.setArg(2, y.cols()); + kernel_check_nan.setArg(3, buffer_nan_flag); + + cmd_queue.enqueueNDRangeKernel(kernel_check_nan, + cl::NullRange, cl::NDRange(y.rows(), y.cols()), cl::NullRange); + + cmd_queue.enqueueReadBuffer(buffer_nan_flag, CL_TRUE, 0, + sizeof(int), &nan_flag); + } catch (const cl::Error& e) { + check_opencl_error("nan_check", e); + } + // if NaN values were found in the matrix + if (nan_flag) { + // in case this is a part of the positive definite check + // output is different + if ( pos_def_check ) { + domain_error(function, name, + "is not positive definite", ""); + } else { + domain_error(function, name, + "has NaN values", ""); + } + } + } + /** + * Check if the specified matrix on the GPU is symmetric + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y Matrix to test + * + * @throw std::domain_error if + * the matrix is not symmetric. + */ + inline void + check_symmetric(const char* function, + const char* name, + const matrix_gpu& y) { + if (y.size() == 0) return; + check_square(function, name, y); + cl::Kernel kernel_check_symmetric = + opencl_context.get_kernel("check_symmetric"); + cl::CommandQueue cmd_queue = opencl_context.queue(); + cl::Context& ctx = opencl_context.context(); + int symmetric_flag = 0; + try { + cl::Buffer buffer_symmetric_flag(ctx, CL_MEM_READ_WRITE, + sizeof(int)); + + cmd_queue.enqueueWriteBuffer(buffer_symmetric_flag, CL_TRUE, 0, + sizeof(int), &symmetric_flag); + + kernel_check_symmetric.setArg(0, y.buffer()); + kernel_check_symmetric.setArg(1, y.rows()); + kernel_check_symmetric.setArg(2, y.cols()); + kernel_check_symmetric.setArg(3, buffer_symmetric_flag); + kernel_check_symmetric.setArg(4, math::CONSTRAINT_TOLERANCE); + + cmd_queue.enqueueNDRangeKernel(kernel_check_symmetric, + cl::NullRange, cl::NDRange(y.rows(), y.cols()), cl::NullRange); + + cmd_queue.enqueueReadBuffer(buffer_symmetric_flag, CL_TRUE, 0, + sizeof(int), &symmetric_flag); + } catch (const cl::Error& e) { + check_opencl_error("symmetric_check", e); + } + // if the matrix is not symmetric + if (symmetric_flag) { + domain_error(function, name, "is not symmetric", ""); + } + } + /** + * Check if the specified matrix on the GPU has zeros on the diagonal + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y Matrix to test + * @param pos_def_check Test if matrix y is positive definite + * + * @throw std::domain_error if + * any diagonal element of the matrix is zero. + */ + inline void + check_diagonal_zeros(const char* function, + const char* name, + const matrix_gpu& y, + const int pos_def_check = 0) { + if (y.size() == 0) return; + + cl::Kernel kernel_check_diagonal_zeros = + opencl_context.get_kernel("check_diagonal_zeros"); + cl::CommandQueue cmd_queue = opencl_context.queue(); + cl::Context ctx = opencl_context.context(); + int flag = 0; + try { + cl::Buffer buffer_flag(ctx, CL_MEM_READ_WRITE, + sizeof(int)); + + cmd_queue.enqueueWriteBuffer(buffer_flag, CL_TRUE, 0, + sizeof(int), &flag); + + kernel_check_diagonal_zeros.setArg(0, y.buffer()); + kernel_check_diagonal_zeros.setArg(1, y.rows()); + kernel_check_diagonal_zeros.setArg(2, y.cols()); + kernel_check_diagonal_zeros.setArg(3, buffer_flag); + + cmd_queue.enqueueNDRangeKernel(kernel_check_diagonal_zeros, + cl::NullRange, cl::NDRange(y.rows(), y.cols()), cl::NullRange); + + cmd_queue.enqueueReadBuffer(buffer_flag, CL_TRUE, 0, + sizeof(int), &flag); + } catch (const cl::Error& e) { + check_opencl_error("diag_zeros_check", e); + } + // if zeros were found on the diagonal + if (flag) { + if ( pos_def_check ) { + domain_error(function, name, + "is not positive definite", ""); + } else { + domain_error(function, name, + "has zeros on the diagonal.", ""); + } + } + } + /** + * Check if the specified matrix on the GPU has NaN values + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y Matrix to test + * + * @throw std::domain_error if + * any element of the matrix is NaN. + */ + inline void + check_positive_definite(const char* function, + const char* name, + const matrix_gpu& y) { + if (y.size() == 0) return; + check_nan(function, name, y, 1); + check_diagonal_zeros(function, name, y, 1); + } + /** + * Check if the two matrices on the GPU are of the same size. + * + * This function checks the runtime sizes only. + * + * + * @param function Function name (for error messages) + * @param name1 Variable name for the first matrix (for error messages) + * @param y1 First matrix + * @param name2 Variable name for the second matrix (for error messages) + * @param y2 Second matrix + * + * @throw std::invalid_argument + * if the dimensions of the matrices do not match + */ + inline void check_matching_dims(const char* function, + const char* name1, + const matrix_gpu& y1, + const char* name2, + const matrix_gpu& y2) { + check_size_match(function, + "Rows of ", name1, y1.rows(), + "rows of ", name2, y2.rows()); + check_size_match(function, + "Columns of ", name1, y1.cols(), + "columns of ", name2, y2.cols()); + } + + } +} +#endif +#endif diff --git a/stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl b/stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl new file mode 100644 index 00000000000..941cce7fe27 --- /dev/null +++ b/stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl @@ -0,0 +1,16 @@ +R"( +#ifndef A +#define A(i,j) A[j*rows+i] +#endif +__kernel void check_diagonal_zeros( + __global double *A, + int rows, + int cols, + __global int *flag) { + const int i = get_global_id(0); + if( i < rows && i < cols ) { + if (A(i,i) == 0) { + flag[0] = 1; + } + } +};)" diff --git a/stan/math/gpu/kernels/check_nan_kernel.cl b/stan/math/gpu/kernels/check_nan_kernel.cl new file mode 100644 index 00000000000..859926a662c --- /dev/null +++ b/stan/math/gpu/kernels/check_nan_kernel.cl @@ -0,0 +1,17 @@ +R"( +#ifndef A +#define A(i,j) A[j*rows+i] +#endif +__kernel void check_nan( + __global double *A, + int rows, + int cols, + __global int *flag) { + const int i = get_global_id(0); + const int j = get_global_id(1); + if( i < rows && j < cols ) { + if (isnan(A(i,j))) { + flag[0] = 1; + } + } +};)" diff --git a/stan/math/gpu/kernels/check_symmetric_kernel.cl b/stan/math/gpu/kernels/check_symmetric_kernel.cl new file mode 100644 index 00000000000..08f7a2eab48 --- /dev/null +++ b/stan/math/gpu/kernels/check_symmetric_kernel.cl @@ -0,0 +1,19 @@ +R"( +#ifndef A +#define A(i,j) A[j*rows+i] +#endif +__kernel void check_symmetric( + __global double *A, + int rows, + int cols, + __global int *flag, + double tolerance) { + const int i = get_global_id(0); + const int j = get_global_id(1); + if( i < rows && j < cols ) { + double diff = fabs(A(i,j)-A(j,i)); + if ( diff > tolerance ) { + flag[0] = 1; + } + } +};)" diff --git a/stan/math/gpu/opencl_context.hpp b/stan/math/gpu/opencl_context.hpp index bd2d8419603..38f6411aa22 100644 --- a/stan/math/gpu/opencl_context.hpp +++ b/stan/math/gpu/opencl_context.hpp @@ -113,11 +113,23 @@ class opencl_context_base { const char* copy_matrix_kernel = #include ; // NOLINT + const char* check_nan_kernel = +#include + ; // NOLINT + const char* check_diagonal_zeros_kernel = +#include + ; // NOLINT + const char* check_symmetric_kernel = +#include + ; // NOLINT kernel_info["dummy"] = { false, "timing", "__kernel void dummy(__global const int* foo) { };"}; kernel_info["dummy2"] = { false, "timing", "__kernel void dummy2(__global const int* foo) { };"}; kernel_info["copy"] = {false, "basic_matrix", copy_matrix_kernel}; + kernel_info["check_nan"] = {false, "check", check_nan_kernel}; + kernel_info["check_diagonal_zeros"] = {false, "check", check_diagonal_zeros_kernel}; + kernel_info["check_symmetric"] = {false, "check", check_symmetric_kernel}; } protected: diff --git a/test/unit/math/gpu/opencl_check_test.cpp b/test/unit/math/gpu/opencl_check_test.cpp new file mode 100644 index 00000000000..f8f43c366bd --- /dev/null +++ b/test/unit/math/gpu/opencl_check_test.cpp @@ -0,0 +1,85 @@ +#include +#include +#include +#include + +using stan::math::check_nan; + +// ---------- check_nan: matrix tests ---------- +TEST(ErrorHandlingScalarGPU, Check_nan_Matrix) { + const char* function = "check_nan"; + Eigen::Matrix x; + using stan::math::matrix_gpu; + x.resize(3); + x << -1, 0, 1; + matrix_gpu xx(x); + + ASSERT_NO_THROW(check_nan(function, "xx", xx)) + << "check_nan should be true with finite xx"; +} + +TEST(ErrorHandlingScalarGPU, Check_nan_Matrix_quit_nan) { + const char* function = "check_nan"; + Eigen::Matrix x; + using stan::math::matrix_gpu; + + x.resize(3); + x << -1, 0, std::numeric_limits::quiet_NaN(); + matrix_gpu xx(x); + EXPECT_THROW(check_nan(function, "xx", xx), std::domain_error) + << "check_nan should throw exception on NaN"; +} + +TEST(ErrorHandlingScalarGPU, Check_nan_positions) { + const char* function = "check_nan"; + double nan = std::numeric_limits::quiet_NaN(); + using stan::math::matrix_gpu; + Eigen::Matrix x_mat(3); + x_mat << nan, 0, 1; + matrix_gpu xx_mat1(x_mat); + EXPECT_THROW(check_nan(function, "xx_mat1", xx_mat1), + std::domain_error); + + x_mat << 1, nan, 1; + matrix_gpu xx_mat2(x_mat); + EXPECT_THROW(check_nan(function, "xx_mat2", xx_mat2), + std::domain_error); + + x_mat << 1, 0, nan; + matrix_gpu xx_mat3(x_mat); + EXPECT_THROW(check_nan(function, "xx_mat3", xx_mat3), + std::domain_error); +} + +TEST(ErrorHandlingScalarGPU, check_rv_v_symmetric_gpu) { + const char* function = "check_symmetric"; + + stan::math::row_vector_d rv; + stan::math::vector_d v; + rv.resize(3); + v.resize(3); + stan::math::matrix_gpu rvv(rv); + stan::math::matrix_gpu vv(v); + EXPECT_THROW(check_symmetric(function, "rvv_non_symm", rvv), + std::invalid_argument); + EXPECT_THROW(check_symmetric(function, "v_non_symm", vv), + std::invalid_argument); +} + +TEST(ErrorHandlingScalarGPU, check_m_symmetric) { + const char* function = "check_symmetric"; + + stan::math::matrix_d m_ok(3, 3); + stan::math::matrix_d m_fail(3, 3); + m_fail << 1, 2, 3, + 2, 4, -5, + 3, 5, 6; + stan::math::matrix_gpu mm_fail(m_fail); + m_ok << 1, 2, 3, + 2, 4, 5, + 3, 5, 6; + stan::math::matrix_gpu mm_ok(m_ok); + EXPECT_THROW(check_symmetric(function, "m_non_symm", mm_fail), + std::domain_error); + EXPECT_NO_THROW(check_symmetric(function, "m_symm_mat1", mm_ok)); +} From b76fe33d8a95837587acdb30d3cab5d926e93ab2 Mon Sep 17 00:00:00 2001 From: rok-cesnovar Date: Mon, 4 Jun 2018 22:25:42 +0200 Subject: [PATCH 02/14] added ifndef for OPENCL define --- test/unit/math/gpu/opencl_check_test.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/test/unit/math/gpu/opencl_check_test.cpp b/test/unit/math/gpu/opencl_check_test.cpp index f8f43c366bd..09b97198cb0 100644 --- a/test/unit/math/gpu/opencl_check_test.cpp +++ b/test/unit/math/gpu/opencl_check_test.cpp @@ -1,3 +1,4 @@ +#ifdef STAN_OPENCL #include #include #include @@ -83,3 +84,4 @@ TEST(ErrorHandlingScalarGPU, check_m_symmetric) { std::domain_error); EXPECT_NO_THROW(check_symmetric(function, "m_symm_mat1", mm_ok)); } +#endif From 60c791bb9d1d8c44d91073c2b33db17a81b0448a Mon Sep 17 00:00:00 2001 From: Stan Jenkins Date: Mon, 4 Jun 2018 21:19:56 -0400 Subject: [PATCH 03/14] [Jenkins] auto-formatting by clang-format version 6.0.0 (tags/google/stable/2017-11-14) --- stan/math/gpu/err/check_matrix_gpu.hpp | 426 +++++++++++------------ stan/math/gpu/opencl_context.hpp | 3 +- test/unit/math/gpu/opencl_check_test.cpp | 21 +- 3 files changed, 213 insertions(+), 237 deletions(-) diff --git a/stan/math/gpu/err/check_matrix_gpu.hpp b/stan/math/gpu/err/check_matrix_gpu.hpp index e813e320962..5c9a5dd678c 100644 --- a/stan/math/gpu/err/check_matrix_gpu.hpp +++ b/stan/math/gpu/err/check_matrix_gpu.hpp @@ -10,231 +10,213 @@ #include namespace stan { - namespace math { - /** - * Check if the specified matrix on the GPU is square. - * - * This check allows 0x0 matrices. - * - * - * @param function Function name (for error messages) - * @param name Variable name (for error messages) - * @param y Matrix on the GPU to test - * - * @throw std::invalid_argument if the matrix - * is not square - */ - inline void - check_square(const char* function, - const char* name, - const matrix_gpu& y) { - check_size_match(function, - "Expecting a square matrix; rows of ", name, y.rows(), - "columns of ", name, y.cols()); - } - /** - * Check if the specified matrix on the GPU has NaN values - * - * @param function Function name (for error messages) - * @param name Variable name (for error messages) - * @param y Matrix to test - * @param pos_def_check Test if matrix y is positive definite - * - * @throw std::domain_error if - * any element of the matrix is NaN. - */ - inline void - check_nan(const char* function, - const char* name, - const matrix_gpu& y, - const int pos_def_check = 0) { - if (y.size() == 0) return; - - cl::Kernel kernel_check_nan = opencl_context.get_kernel("check_nan"); - cl::CommandQueue cmd_queue = opencl_context.queue(); - cl::Context& ctx = opencl_context.context(); - int nan_flag = 0; - try { - cl::Buffer buffer_nan_flag(ctx, CL_MEM_READ_WRITE, - sizeof(int)); - - cmd_queue.enqueueWriteBuffer(buffer_nan_flag, CL_TRUE, 0, - sizeof(int), &nan_flag); - - kernel_check_nan.setArg(0, y.buffer()); - kernel_check_nan.setArg(1, y.rows()); - kernel_check_nan.setArg(2, y.cols()); - kernel_check_nan.setArg(3, buffer_nan_flag); - - cmd_queue.enqueueNDRangeKernel(kernel_check_nan, - cl::NullRange, cl::NDRange(y.rows(), y.cols()), cl::NullRange); - - cmd_queue.enqueueReadBuffer(buffer_nan_flag, CL_TRUE, 0, - sizeof(int), &nan_flag); - } catch (const cl::Error& e) { - check_opencl_error("nan_check", e); - } - // if NaN values were found in the matrix - if (nan_flag) { - // in case this is a part of the positive definite check - // output is different - if ( pos_def_check ) { - domain_error(function, name, - "is not positive definite", ""); - } else { - domain_error(function, name, - "has NaN values", ""); - } - } - } - /** - * Check if the specified matrix on the GPU is symmetric - * - * @param function Function name (for error messages) - * @param name Variable name (for error messages) - * @param y Matrix to test - * - * @throw std::domain_error if - * the matrix is not symmetric. - */ - inline void - check_symmetric(const char* function, - const char* name, - const matrix_gpu& y) { - if (y.size() == 0) return; - check_square(function, name, y); - cl::Kernel kernel_check_symmetric = - opencl_context.get_kernel("check_symmetric"); - cl::CommandQueue cmd_queue = opencl_context.queue(); - cl::Context& ctx = opencl_context.context(); - int symmetric_flag = 0; - try { - cl::Buffer buffer_symmetric_flag(ctx, CL_MEM_READ_WRITE, - sizeof(int)); - - cmd_queue.enqueueWriteBuffer(buffer_symmetric_flag, CL_TRUE, 0, - sizeof(int), &symmetric_flag); - - kernel_check_symmetric.setArg(0, y.buffer()); - kernel_check_symmetric.setArg(1, y.rows()); - kernel_check_symmetric.setArg(2, y.cols()); - kernel_check_symmetric.setArg(3, buffer_symmetric_flag); - kernel_check_symmetric.setArg(4, math::CONSTRAINT_TOLERANCE); - - cmd_queue.enqueueNDRangeKernel(kernel_check_symmetric, - cl::NullRange, cl::NDRange(y.rows(), y.cols()), cl::NullRange); - - cmd_queue.enqueueReadBuffer(buffer_symmetric_flag, CL_TRUE, 0, - sizeof(int), &symmetric_flag); - } catch (const cl::Error& e) { - check_opencl_error("symmetric_check", e); - } - // if the matrix is not symmetric - if (symmetric_flag) { - domain_error(function, name, "is not symmetric", ""); - } - } - /** - * Check if the specified matrix on the GPU has zeros on the diagonal - * - * @param function Function name (for error messages) - * @param name Variable name (for error messages) - * @param y Matrix to test - * @param pos_def_check Test if matrix y is positive definite - * - * @throw std::domain_error if - * any diagonal element of the matrix is zero. - */ - inline void - check_diagonal_zeros(const char* function, - const char* name, - const matrix_gpu& y, - const int pos_def_check = 0) { - if (y.size() == 0) return; - - cl::Kernel kernel_check_diagonal_zeros = - opencl_context.get_kernel("check_diagonal_zeros"); - cl::CommandQueue cmd_queue = opencl_context.queue(); - cl::Context ctx = opencl_context.context(); - int flag = 0; - try { - cl::Buffer buffer_flag(ctx, CL_MEM_READ_WRITE, - sizeof(int)); - - cmd_queue.enqueueWriteBuffer(buffer_flag, CL_TRUE, 0, - sizeof(int), &flag); - - kernel_check_diagonal_zeros.setArg(0, y.buffer()); - kernel_check_diagonal_zeros.setArg(1, y.rows()); - kernel_check_diagonal_zeros.setArg(2, y.cols()); - kernel_check_diagonal_zeros.setArg(3, buffer_flag); - - cmd_queue.enqueueNDRangeKernel(kernel_check_diagonal_zeros, - cl::NullRange, cl::NDRange(y.rows(), y.cols()), cl::NullRange); - - cmd_queue.enqueueReadBuffer(buffer_flag, CL_TRUE, 0, - sizeof(int), &flag); - } catch (const cl::Error& e) { - check_opencl_error("diag_zeros_check", e); - } - // if zeros were found on the diagonal - if (flag) { - if ( pos_def_check ) { - domain_error(function, name, - "is not positive definite", ""); - } else { - domain_error(function, name, - "has zeros on the diagonal.", ""); - } - } - } - /** - * Check if the specified matrix on the GPU has NaN values - * - * @param function Function name (for error messages) - * @param name Variable name (for error messages) - * @param y Matrix to test - * - * @throw std::domain_error if - * any element of the matrix is NaN. - */ - inline void - check_positive_definite(const char* function, - const char* name, - const matrix_gpu& y) { - if (y.size() == 0) return; - check_nan(function, name, y, 1); - check_diagonal_zeros(function, name, y, 1); +namespace math { +/** + * Check if the specified matrix on the GPU is square. + * + * This check allows 0x0 matrices. + * + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y Matrix on the GPU to test + * + * @throw std::invalid_argument if the matrix + * is not square + */ +inline void check_square(const char* function, const char* name, + const matrix_gpu& y) { + check_size_match(function, "Expecting a square matrix; rows of ", name, + y.rows(), "columns of ", name, y.cols()); +} +/** + * Check if the specified matrix on the GPU has NaN values + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y Matrix to test + * @param pos_def_check Test if matrix y is positive definite + * + * @throw std::domain_error if + * any element of the matrix is NaN. + */ +inline void check_nan(const char* function, const char* name, + const matrix_gpu& y, const int pos_def_check = 0) { + if (y.size() == 0) + return; + + cl::Kernel kernel_check_nan = opencl_context.get_kernel("check_nan"); + cl::CommandQueue cmd_queue = opencl_context.queue(); + cl::Context& ctx = opencl_context.context(); + int nan_flag = 0; + try { + cl::Buffer buffer_nan_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); + + cmd_queue.enqueueWriteBuffer(buffer_nan_flag, CL_TRUE, 0, sizeof(int), + &nan_flag); + + kernel_check_nan.setArg(0, y.buffer()); + kernel_check_nan.setArg(1, y.rows()); + kernel_check_nan.setArg(2, y.cols()); + kernel_check_nan.setArg(3, buffer_nan_flag); + + cmd_queue.enqueueNDRangeKernel(kernel_check_nan, cl::NullRange, + cl::NDRange(y.rows(), y.cols()), + cl::NullRange); + + cmd_queue.enqueueReadBuffer(buffer_nan_flag, CL_TRUE, 0, sizeof(int), + &nan_flag); + } catch (const cl::Error& e) { + check_opencl_error("nan_check", e); + } + // if NaN values were found in the matrix + if (nan_flag) { + // in case this is a part of the positive definite check + // output is different + if (pos_def_check) { + domain_error(function, name, "is not positive definite", ""); + } else { + domain_error(function, name, "has NaN values", ""); } - /** - * Check if the two matrices on the GPU are of the same size. - * - * This function checks the runtime sizes only. - * - * - * @param function Function name (for error messages) - * @param name1 Variable name for the first matrix (for error messages) - * @param y1 First matrix - * @param name2 Variable name for the second matrix (for error messages) - * @param y2 Second matrix - * - * @throw std::invalid_argument - * if the dimensions of the matrices do not match - */ - inline void check_matching_dims(const char* function, - const char* name1, - const matrix_gpu& y1, - const char* name2, - const matrix_gpu& y2) { - check_size_match(function, - "Rows of ", name1, y1.rows(), - "rows of ", name2, y2.rows()); - check_size_match(function, - "Columns of ", name1, y1.cols(), - "columns of ", name2, y2.cols()); + } +} +/** + * Check if the specified matrix on the GPU is symmetric + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y Matrix to test + * + * @throw std::domain_error if + * the matrix is not symmetric. + */ +inline void check_symmetric(const char* function, const char* name, + const matrix_gpu& y) { + if (y.size() == 0) + return; + check_square(function, name, y); + cl::Kernel kernel_check_symmetric + = opencl_context.get_kernel("check_symmetric"); + cl::CommandQueue cmd_queue = opencl_context.queue(); + cl::Context& ctx = opencl_context.context(); + int symmetric_flag = 0; + try { + cl::Buffer buffer_symmetric_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); + + cmd_queue.enqueueWriteBuffer(buffer_symmetric_flag, CL_TRUE, 0, sizeof(int), + &symmetric_flag); + + kernel_check_symmetric.setArg(0, y.buffer()); + kernel_check_symmetric.setArg(1, y.rows()); + kernel_check_symmetric.setArg(2, y.cols()); + kernel_check_symmetric.setArg(3, buffer_symmetric_flag); + kernel_check_symmetric.setArg(4, math::CONSTRAINT_TOLERANCE); + + cmd_queue.enqueueNDRangeKernel(kernel_check_symmetric, cl::NullRange, + cl::NDRange(y.rows(), y.cols()), + cl::NullRange); + + cmd_queue.enqueueReadBuffer(buffer_symmetric_flag, CL_TRUE, 0, sizeof(int), + &symmetric_flag); + } catch (const cl::Error& e) { + check_opencl_error("symmetric_check", e); + } + // if the matrix is not symmetric + if (symmetric_flag) { + domain_error(function, name, "is not symmetric", ""); + } +} +/** + * Check if the specified matrix on the GPU has zeros on the diagonal + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y Matrix to test + * @param pos_def_check Test if matrix y is positive definite + * + * @throw std::domain_error if + * any diagonal element of the matrix is zero. + */ +inline void check_diagonal_zeros(const char* function, const char* name, + const matrix_gpu& y, + const int pos_def_check = 0) { + if (y.size() == 0) + return; + + cl::Kernel kernel_check_diagonal_zeros + = opencl_context.get_kernel("check_diagonal_zeros"); + cl::CommandQueue cmd_queue = opencl_context.queue(); + cl::Context ctx = opencl_context.context(); + int flag = 0; + try { + cl::Buffer buffer_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); + + cmd_queue.enqueueWriteBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &flag); + + kernel_check_diagonal_zeros.setArg(0, y.buffer()); + kernel_check_diagonal_zeros.setArg(1, y.rows()); + kernel_check_diagonal_zeros.setArg(2, y.cols()); + kernel_check_diagonal_zeros.setArg(3, buffer_flag); + + cmd_queue.enqueueNDRangeKernel(kernel_check_diagonal_zeros, cl::NullRange, + cl::NDRange(y.rows(), y.cols()), + cl::NullRange); + + cmd_queue.enqueueReadBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &flag); + } catch (const cl::Error& e) { + check_opencl_error("diag_zeros_check", e); + } + // if zeros were found on the diagonal + if (flag) { + if (pos_def_check) { + domain_error(function, name, "is not positive definite", ""); + } else { + domain_error(function, name, "has zeros on the diagonal.", ""); } - } } +/** + * Check if the specified matrix on the GPU has NaN values + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y Matrix to test + * + * @throw std::domain_error if + * any element of the matrix is NaN. + */ +inline void check_positive_definite(const char* function, const char* name, + const matrix_gpu& y) { + if (y.size() == 0) + return; + check_nan(function, name, y, 1); + check_diagonal_zeros(function, name, y, 1); +} +/** + * Check if the two matrices on the GPU are of the same size. + * + * This function checks the runtime sizes only. + * + * + * @param function Function name (for error messages) + * @param name1 Variable name for the first matrix (for error messages) + * @param y1 First matrix + * @param name2 Variable name for the second matrix (for error messages) + * @param y2 Second matrix + * + * @throw std::invalid_argument + * if the dimensions of the matrices do not match + */ +inline void check_matching_dims(const char* function, const char* name1, + const matrix_gpu& y1, const char* name2, + const matrix_gpu& y2) { + check_size_match(function, "Rows of ", name1, y1.rows(), "rows of ", name2, + y2.rows()); + check_size_match(function, "Columns of ", name1, y1.cols(), "columns of ", + name2, y2.cols()); +} + +} // namespace math +} // namespace stan #endif #endif diff --git a/stan/math/gpu/opencl_context.hpp b/stan/math/gpu/opencl_context.hpp index 38f6411aa22..b1d20a919a7 100644 --- a/stan/math/gpu/opencl_context.hpp +++ b/stan/math/gpu/opencl_context.hpp @@ -128,7 +128,8 @@ class opencl_context_base { false, "timing", "__kernel void dummy2(__global const int* foo) { };"}; kernel_info["copy"] = {false, "basic_matrix", copy_matrix_kernel}; kernel_info["check_nan"] = {false, "check", check_nan_kernel}; - kernel_info["check_diagonal_zeros"] = {false, "check", check_diagonal_zeros_kernel}; + kernel_info["check_diagonal_zeros"] + = {false, "check", check_diagonal_zeros_kernel}; kernel_info["check_symmetric"] = {false, "check", check_symmetric_kernel}; } diff --git a/test/unit/math/gpu/opencl_check_test.cpp b/test/unit/math/gpu/opencl_check_test.cpp index 09b97198cb0..2902b4839ec 100644 --- a/test/unit/math/gpu/opencl_check_test.cpp +++ b/test/unit/math/gpu/opencl_check_test.cpp @@ -16,7 +16,7 @@ TEST(ErrorHandlingScalarGPU, Check_nan_Matrix) { matrix_gpu xx(x); ASSERT_NO_THROW(check_nan(function, "xx", xx)) - << "check_nan should be true with finite xx"; + << "check_nan should be true with finite xx"; } TEST(ErrorHandlingScalarGPU, Check_nan_Matrix_quit_nan) { @@ -28,7 +28,7 @@ TEST(ErrorHandlingScalarGPU, Check_nan_Matrix_quit_nan) { x << -1, 0, std::numeric_limits::quiet_NaN(); matrix_gpu xx(x); EXPECT_THROW(check_nan(function, "xx", xx), std::domain_error) - << "check_nan should throw exception on NaN"; + << "check_nan should throw exception on NaN"; } TEST(ErrorHandlingScalarGPU, Check_nan_positions) { @@ -38,18 +38,15 @@ TEST(ErrorHandlingScalarGPU, Check_nan_positions) { Eigen::Matrix x_mat(3); x_mat << nan, 0, 1; matrix_gpu xx_mat1(x_mat); - EXPECT_THROW(check_nan(function, "xx_mat1", xx_mat1), - std::domain_error); + EXPECT_THROW(check_nan(function, "xx_mat1", xx_mat1), std::domain_error); x_mat << 1, nan, 1; matrix_gpu xx_mat2(x_mat); - EXPECT_THROW(check_nan(function, "xx_mat2", xx_mat2), - std::domain_error); + EXPECT_THROW(check_nan(function, "xx_mat2", xx_mat2), std::domain_error); x_mat << 1, 0, nan; matrix_gpu xx_mat3(x_mat); - EXPECT_THROW(check_nan(function, "xx_mat3", xx_mat3), - std::domain_error); + EXPECT_THROW(check_nan(function, "xx_mat3", xx_mat3), std::domain_error); } TEST(ErrorHandlingScalarGPU, check_rv_v_symmetric_gpu) { @@ -72,13 +69,9 @@ TEST(ErrorHandlingScalarGPU, check_m_symmetric) { stan::math::matrix_d m_ok(3, 3); stan::math::matrix_d m_fail(3, 3); - m_fail << 1, 2, 3, - 2, 4, -5, - 3, 5, 6; + m_fail << 1, 2, 3, 2, 4, -5, 3, 5, 6; stan::math::matrix_gpu mm_fail(m_fail); - m_ok << 1, 2, 3, - 2, 4, 5, - 3, 5, 6; + m_ok << 1, 2, 3, 2, 4, 5, 3, 5, 6; stan::math::matrix_gpu mm_ok(m_ok); EXPECT_THROW(check_symmetric(function, "m_non_symm", mm_fail), std::domain_error); From 6b8d80f169dfda217a41152de13aa4419fc38320 Mon Sep 17 00:00:00 2001 From: Steve Bronder Date: Mon, 4 Jun 2018 21:42:42 -0400 Subject: [PATCH 04/14] Fix docs --- stan/math/gpu/err/check_matrix_gpu.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/stan/math/gpu/err/check_matrix_gpu.hpp b/stan/math/gpu/err/check_matrix_gpu.hpp index e813e320962..2e4a34d0aa3 100644 --- a/stan/math/gpu/err/check_matrix_gpu.hpp +++ b/stan/math/gpu/err/check_matrix_gpu.hpp @@ -189,7 +189,7 @@ namespace stan { } } /** - * Check if the specified matrix on the GPU has NaN values + * Check if the specified matrix on the GPU is positive definite * * @param function Function name (for error messages) * @param name Variable name (for error messages) @@ -210,7 +210,7 @@ namespace stan { * Check if the two matrices on the GPU are of the same size. * * This function checks the runtime sizes only. - * + * * * @param function Function name (for error messages) * @param name1 Variable name for the first matrix (for error messages) From ecb9475926a9fba0de59bb98f418822b29c0422a Mon Sep 17 00:00:00 2001 From: Steve Bronder Date: Mon, 4 Jun 2018 21:45:33 -0400 Subject: [PATCH 05/14] update docs --- stan/math/gpu/err/check_matrix_gpu.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/stan/math/gpu/err/check_matrix_gpu.hpp b/stan/math/gpu/err/check_matrix_gpu.hpp index 5c9a5dd678c..42de0fd7169 100644 --- a/stan/math/gpu/err/check_matrix_gpu.hpp +++ b/stan/math/gpu/err/check_matrix_gpu.hpp @@ -176,7 +176,7 @@ inline void check_diagonal_zeros(const char* function, const char* name, } } /** - * Check if the specified matrix on the GPU has NaN values + * Check if the specified matrix on the GPU is positive definite * * @param function Function name (for error messages) * @param name Variable name (for error messages) From 3e40566bd74208809a0673fd33dcaf7b0cc39c1b Mon Sep 17 00:00:00 2001 From: Steve Bronder Date: Tue, 5 Jun 2018 19:55:10 -0400 Subject: [PATCH 06/14] Adds matrix_gpu<\code> to the docs and moved around the flags and if statements to throw to be within the if statement --- stan/math/gpu/err/check_matrix_gpu.hpp | 97 +++++++++----------------- 1 file changed, 32 insertions(+), 65 deletions(-) diff --git a/stan/math/gpu/err/check_matrix_gpu.hpp b/stan/math/gpu/err/check_matrix_gpu.hpp index 42de0fd7169..f5c93890718 100644 --- a/stan/math/gpu/err/check_matrix_gpu.hpp +++ b/stan/math/gpu/err/check_matrix_gpu.hpp @@ -12,16 +12,16 @@ namespace stan { namespace math { /** - * Check if the specified matrix on the GPU is square. + * Check if the matrix_gpu is square. * * This check allows 0x0 matrices. * * * @param function Function name (for error messages) * @param name Variable name (for error messages) - * @param y Matrix on the GPU to test + * @param y matrix_gpu to test * - * @throw std::invalid_argument if the matrix + * @throw std::invalid_argument if the matrix_gpu * is not square */ inline void check_square(const char* function, const char* name, @@ -30,28 +30,26 @@ inline void check_square(const char* function, const char* name, y.rows(), "columns of ", name, y.cols()); } /** - * Check if the specified matrix on the GPU has NaN values + * Check if the matrix_gpu has NaN values * * @param function Function name (for error messages) * @param name Variable name (for error messages) - * @param y Matrix to test - * @param pos_def_check Test if matrix y is positive definite + * @param y matrix_gpu to test * * @throw std::domain_error if * any element of the matrix is NaN. */ inline void check_nan(const char* function, const char* name, - const matrix_gpu& y, const int pos_def_check = 0) { + const matrix_gpu& y) { if (y.size() == 0) return; cl::Kernel kernel_check_nan = opencl_context.get_kernel("check_nan"); cl::CommandQueue cmd_queue = opencl_context.queue(); cl::Context& ctx = opencl_context.context(); - int nan_flag = 0; try { + int nan_flag = 0; cl::Buffer buffer_nan_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); - cmd_queue.enqueueWriteBuffer(buffer_nan_flag, CL_TRUE, 0, sizeof(int), &nan_flag); @@ -66,26 +64,20 @@ inline void check_nan(const char* function, const char* name, cmd_queue.enqueueReadBuffer(buffer_nan_flag, CL_TRUE, 0, sizeof(int), &nan_flag); + // if NaN values were found in the matrix + if (nan_flag) { + domain_error(function, name, "has NaN values", ""); + } } catch (const cl::Error& e) { check_opencl_error("nan_check", e); } - // if NaN values were found in the matrix - if (nan_flag) { - // in case this is a part of the positive definite check - // output is different - if (pos_def_check) { - domain_error(function, name, "is not positive definite", ""); - } else { - domain_error(function, name, "has NaN values", ""); - } - } } /** - * Check if the specified matrix on the GPU is symmetric + * Check if the matrix_gpu is symmetric * * @param function Function name (for error messages) * @param name Variable name (for error messages) - * @param y Matrix to test + * @param y matrix_gpu to test * * @throw std::domain_error if * the matrix is not symmetric. @@ -99,10 +91,9 @@ inline void check_symmetric(const char* function, const char* name, = opencl_context.get_kernel("check_symmetric"); cl::CommandQueue cmd_queue = opencl_context.queue(); cl::Context& ctx = opencl_context.context(); - int symmetric_flag = 0; try { + int symmetric_flag = 0; cl::Buffer buffer_symmetric_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); - cmd_queue.enqueueWriteBuffer(buffer_symmetric_flag, CL_TRUE, 0, sizeof(int), &symmetric_flag); @@ -118,28 +109,27 @@ inline void check_symmetric(const char* function, const char* name, cmd_queue.enqueueReadBuffer(buffer_symmetric_flag, CL_TRUE, 0, sizeof(int), &symmetric_flag); + // if the matrix is not symmetric + if (symmetric_flag) { + domain_error(function, name, "is not symmetric", ""); + } } catch (const cl::Error& e) { check_opencl_error("symmetric_check", e); } - // if the matrix is not symmetric - if (symmetric_flag) { - domain_error(function, name, "is not symmetric", ""); - } + } /** - * Check if the specified matrix on the GPU has zeros on the diagonal + * Check if the matrix_gpu has zeros on the diagonal * * @param function Function name (for error messages) * @param name Variable name (for error messages) - * @param y Matrix to test - * @param pos_def_check Test if matrix y is positive definite + * @param y matrix_gpu to test * * @throw std::domain_error if * any diagonal element of the matrix is zero. */ inline void check_diagonal_zeros(const char* function, const char* name, - const matrix_gpu& y, - const int pos_def_check = 0) { + const matrix_gpu& y) { if (y.size() == 0) return; @@ -147,10 +137,9 @@ inline void check_diagonal_zeros(const char* function, const char* name, = opencl_context.get_kernel("check_diagonal_zeros"); cl::CommandQueue cmd_queue = opencl_context.queue(); cl::Context ctx = opencl_context.context(); - int flag = 0; try { + int flag = 0; cl::Buffer buffer_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); - cmd_queue.enqueueWriteBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &flag); kernel_check_diagonal_zeros.setArg(0, y.buffer()); @@ -163,46 +152,24 @@ inline void check_diagonal_zeros(const char* function, const char* name, cl::NullRange); cmd_queue.enqueueReadBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &flag); + // if zeros were found on the diagonal + if (flag) { + domain_error(function, name, "has zeros on the diagonal.", ""); + } + } catch (const cl::Error& e) { check_opencl_error("diag_zeros_check", e); } - // if zeros were found on the diagonal - if (flag) { - if (pos_def_check) { - domain_error(function, name, "is not positive definite", ""); - } else { - domain_error(function, name, "has zeros on the diagonal.", ""); - } - } -} -/** - * Check if the specified matrix on the GPU is positive definite - * - * @param function Function name (for error messages) - * @param name Variable name (for error messages) - * @param y Matrix to test - * - * @throw std::domain_error if - * any element of the matrix is NaN. - */ -inline void check_positive_definite(const char* function, const char* name, - const matrix_gpu& y) { - if (y.size() == 0) - return; - check_nan(function, name, y, 1); - check_diagonal_zeros(function, name, y, 1); } + /** - * Check if the two matrices on the GPU are of the same size. - * - * This function checks the runtime sizes only. - * + * Check if two matrix_gpus have the same dimensions. * * @param function Function name (for error messages) * @param name1 Variable name for the first matrix (for error messages) - * @param y1 First matrix + * @param y1 First matrix_gpu * @param name2 Variable name for the second matrix (for error messages) - * @param y2 Second matrix + * @param y2 Second matrix_gpu * * @throw std::invalid_argument * if the dimensions of the matrices do not match From f3d9bacd33a0dd0297fde5364f1d17dbf74d3ac9 Mon Sep 17 00:00:00 2001 From: Stan Jenkins Date: Tue, 5 Jun 2018 19:55:59 -0400 Subject: [PATCH 07/14] [Jenkins] auto-formatting by clang-format version 6.0.0 (tags/google/stable/2017-11-14) --- stan/math/gpu/err/check_matrix_gpu.hpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/stan/math/gpu/err/check_matrix_gpu.hpp b/stan/math/gpu/err/check_matrix_gpu.hpp index f5c93890718..99ddbc768ae 100644 --- a/stan/math/gpu/err/check_matrix_gpu.hpp +++ b/stan/math/gpu/err/check_matrix_gpu.hpp @@ -66,7 +66,7 @@ inline void check_nan(const char* function, const char* name, &nan_flag); // if NaN values were found in the matrix if (nan_flag) { - domain_error(function, name, "has NaN values", ""); + domain_error(function, name, "has NaN values", ""); } } catch (const cl::Error& e) { check_opencl_error("nan_check", e); @@ -116,7 +116,6 @@ inline void check_symmetric(const char* function, const char* name, } catch (const cl::Error& e) { check_opencl_error("symmetric_check", e); } - } /** * Check if the matrix_gpu has zeros on the diagonal @@ -154,7 +153,7 @@ inline void check_diagonal_zeros(const char* function, const char* name, cmd_queue.enqueueReadBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &flag); // if zeros were found on the diagonal if (flag) { - domain_error(function, name, "has zeros on the diagonal.", ""); + domain_error(function, name, "has zeros on the diagonal.", ""); } } catch (const cl::Error& e) { From dd0c6937bbe0e39d2ff9a42be4cea61058060f41 Mon Sep 17 00:00:00 2001 From: Steve Bronder Date: Tue, 5 Jun 2018 20:01:07 -0400 Subject: [PATCH 08/14] lint the OpenCL kernels --- stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl | 6 +++--- stan/math/gpu/kernels/check_nan_kernel.cl | 6 +++--- stan/math/gpu/kernels/check_symmetric_kernel.cl | 8 ++++---- stan/math/gpu/kernels/copy_matrix_kernel.cl | 4 ++-- 4 files changed, 12 insertions(+), 12 deletions(-) diff --git a/stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl b/stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl index 941cce7fe27..db6fb014e91 100644 --- a/stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl +++ b/stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl @@ -1,6 +1,6 @@ R"( #ifndef A -#define A(i,j) A[j*rows+i] +#define A(i, j) A[j * rows + i] #endif __kernel void check_diagonal_zeros( __global double *A, @@ -8,8 +8,8 @@ __kernel void check_diagonal_zeros( int cols, __global int *flag) { const int i = get_global_id(0); - if( i < rows && i < cols ) { - if (A(i,i) == 0) { + if (i < rows && i < cols) { + if (A(i, i) == 0) { flag[0] = 1; } } diff --git a/stan/math/gpu/kernels/check_nan_kernel.cl b/stan/math/gpu/kernels/check_nan_kernel.cl index 859926a662c..f950382ba22 100644 --- a/stan/math/gpu/kernels/check_nan_kernel.cl +++ b/stan/math/gpu/kernels/check_nan_kernel.cl @@ -1,6 +1,6 @@ R"( #ifndef A -#define A(i,j) A[j*rows+i] +#define A(i, j) A[j * rows + i] #endif __kernel void check_nan( __global double *A, @@ -9,8 +9,8 @@ __kernel void check_nan( __global int *flag) { const int i = get_global_id(0); const int j = get_global_id(1); - if( i < rows && j < cols ) { - if (isnan(A(i,j))) { + if (i < rows && j < cols) { + if (isnan(A(i, j))) { flag[0] = 1; } } diff --git a/stan/math/gpu/kernels/check_symmetric_kernel.cl b/stan/math/gpu/kernels/check_symmetric_kernel.cl index 08f7a2eab48..9ec734d7ed9 100644 --- a/stan/math/gpu/kernels/check_symmetric_kernel.cl +++ b/stan/math/gpu/kernels/check_symmetric_kernel.cl @@ -1,6 +1,6 @@ R"( #ifndef A -#define A(i,j) A[j*rows+i] +#define A(i, j) A[j * rows + i] #endif __kernel void check_symmetric( __global double *A, @@ -10,9 +10,9 @@ __kernel void check_symmetric( double tolerance) { const int i = get_global_id(0); const int j = get_global_id(1); - if( i < rows && j < cols ) { - double diff = fabs(A(i,j)-A(j,i)); - if ( diff > tolerance ) { + if (i < rows && j < cols) { + double diff = fabs(A(i, j) - A(j, i)); + if (diff > tolerance) { flag[0] = 1; } } diff --git a/stan/math/gpu/kernels/copy_matrix_kernel.cl b/stan/math/gpu/kernels/copy_matrix_kernel.cl index 0034ef5f9b1..3ca7b6e1318 100644 --- a/stan/math/gpu/kernels/copy_matrix_kernel.cl +++ b/stan/math/gpu/kernels/copy_matrix_kernel.cl @@ -1,6 +1,6 @@ R"( -#define A(i, j) A[j*rows+i] -#define B(i, j) B[j*rows+i] +#define A(i, j) A[j * rows + i] +#define B(i, j) B[j * rows + i] __kernel void copy( __global double *A, __global double *B, From 0134e845be20bfeb0bb1f8889f1a5a8999befa7e Mon Sep 17 00:00:00 2001 From: Steve Bronder Date: Tue, 5 Jun 2018 20:21:21 -0400 Subject: [PATCH 09/14] remove 0 by 0 being accepted info from docs --- stan/math/gpu/err/check_matrix_gpu.hpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/stan/math/gpu/err/check_matrix_gpu.hpp b/stan/math/gpu/err/check_matrix_gpu.hpp index 99ddbc768ae..37b8d6eabcb 100644 --- a/stan/math/gpu/err/check_matrix_gpu.hpp +++ b/stan/math/gpu/err/check_matrix_gpu.hpp @@ -14,9 +14,6 @@ namespace math { /** * Check if the matrix_gpu is square. * - * This check allows 0x0 matrices. - * - * * @param function Function name (for error messages) * @param name Variable name (for error messages) * @param y matrix_gpu to test From bb91d35a974e559ad5f0baa6c6bab75ed5a7a28d Mon Sep 17 00:00:00 2001 From: Steve Bronder Date: Tue, 5 Jun 2018 20:46:09 -0400 Subject: [PATCH 10/14] cpplint error --- stan/math/gpu/err/check_matrix_gpu.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/stan/math/gpu/err/check_matrix_gpu.hpp b/stan/math/gpu/err/check_matrix_gpu.hpp index 37b8d6eabcb..178c939fe15 100644 --- a/stan/math/gpu/err/check_matrix_gpu.hpp +++ b/stan/math/gpu/err/check_matrix_gpu.hpp @@ -152,7 +152,6 @@ inline void check_diagonal_zeros(const char* function, const char* name, if (flag) { domain_error(function, name, "has zeros on the diagonal.", ""); } - } catch (const cl::Error& e) { check_opencl_error("diag_zeros_check", e); } From 5271e2574fd6b198ba769a95b4319e6bd1ba084e Mon Sep 17 00:00:00 2001 From: rok-cesnovar Date: Fri, 8 Jun 2018 22:12:43 +0200 Subject: [PATCH 11/14] added quotation to BOOST_LIB_ABS --- make/libraries | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/make/libraries b/make/libraries index 673fabcd68d..8caa2edbd03 100644 --- a/make/libraries +++ b/make/libraries @@ -73,13 +73,13 @@ $(BOOST_LIB)/mpi.so: $(BOOST_LIB)/libboost_serialization.so: $(BOOST_LIB)/mpi.so $(BOOST_LIB)/libboost_serialization.dylib: $(BOOST_LIB)/mpi.so - install_name_tool -add_rpath $(BOOST_LIB_ABS) $(BOOST_LIB)/libboost_serialization.dylib + install_name_tool -add_rpath "$(BOOST_LIB_ABS)" $(BOOST_LIB)/libboost_serialization.dylib install_name_tool -id @rpath/libboost_serialization.dylib $(BOOST_LIB)/libboost_serialization.dylib $(BOOST_LIB)/libboost_mpi.so: $(BOOST_LIB)/mpi.so $(BOOST_LIB)/libboost_mpi.dylib: $(BOOST_LIB)/mpi.so - install_name_tool -add_rpath $(BOOST_LIB_ABS) $(BOOST_LIB)/libboost_mpi.dylib + install_name_tool -add_rpath "$(BOOST_LIB_ABS)" $(BOOST_LIB)/libboost_mpi.dylib install_name_tool -change libboost_serialization.dylib @rpath/libboost_serialization.dylib $(BOOST_LIB)/libboost_mpi.dylib install_name_tool -id @rpath/libboost_mpi.dylib $(BOOST_LIB)/libboost_mpi.dylib From 181c569cbd46219b5735a6dd434eef0c237f268c Mon Sep 17 00:00:00 2001 From: rok-cesnovar Date: Tue, 19 Jun 2018 11:18:48 +0200 Subject: [PATCH 12/14] split gpu checks in to separate files, fixed the header guards, added tests for check diagonal zeros --- stan/math/gpu/err/check_diagonal_zeros.hpp | 55 ++++++ stan/math/gpu/err/check_matching_dims.hpp | 33 ++++ stan/math/gpu/err/check_matrix_gpu.hpp | 184 --------------------- stan/math/gpu/err/check_nan.hpp | 56 +++++++ stan/math/gpu/err/check_square.hpp | 28 ++++ stan/math/gpu/err/check_symmetric.hpp | 58 +++++++ test/unit/math/gpu/opencl_check_test.cpp | 26 ++- 7 files changed, 252 insertions(+), 188 deletions(-) create mode 100644 stan/math/gpu/err/check_diagonal_zeros.hpp create mode 100644 stan/math/gpu/err/check_matching_dims.hpp delete mode 100644 stan/math/gpu/err/check_matrix_gpu.hpp create mode 100644 stan/math/gpu/err/check_nan.hpp create mode 100644 stan/math/gpu/err/check_square.hpp create mode 100644 stan/math/gpu/err/check_symmetric.hpp diff --git a/stan/math/gpu/err/check_diagonal_zeros.hpp b/stan/math/gpu/err/check_diagonal_zeros.hpp new file mode 100644 index 00000000000..53d801a548e --- /dev/null +++ b/stan/math/gpu/err/check_diagonal_zeros.hpp @@ -0,0 +1,55 @@ +#ifndef STAN_MATH_GPU_ERR_CHECK_DIAGONAL_ZEROS_HPP +#define STAN_MATH_GPU_ERR_CHECK_DIAGONAL_ZEROS_HPP +#ifdef STAN_OPENCL +#include +#include + +namespace stan { +namespace math { +/** + * Check if the matrix_gpu has zeros on the diagonal + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y matrix_gpu to test + * + * @throw std::domain_error if + * any diagonal element of the matrix is zero. + */ +inline void check_diagonal_zeros(const char* function, const char* name, + const matrix_gpu& y) { + if (y.size() == 0) + return; + + cl::Kernel kernel_check_diagonal_zeros + = opencl_context.get_kernel("check_diagonal_zeros"); + cl::CommandQueue cmd_queue = opencl_context.queue(); + cl::Context ctx = opencl_context.context(); + try { + int flag = 0; + cl::Buffer buffer_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); + cmd_queue.enqueueWriteBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &flag); + + kernel_check_diagonal_zeros.setArg(0, y.buffer()); + kernel_check_diagonal_zeros.setArg(1, y.rows()); + kernel_check_diagonal_zeros.setArg(2, y.cols()); + kernel_check_diagonal_zeros.setArg(3, buffer_flag); + + cmd_queue.enqueueNDRangeKernel(kernel_check_diagonal_zeros, cl::NullRange, + cl::NDRange(y.rows(), y.cols()), + cl::NullRange); + + cmd_queue.enqueueReadBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &flag); + // if zeros were found on the diagonal + if (flag) { + domain_error(function, name, "has zeros on the diagonal.", ""); + } + } catch (const cl::Error& e) { + check_opencl_error("diag_zeros_check", e); + } +} + +} // namespace math +} // namespace stan +#endif +#endif diff --git a/stan/math/gpu/err/check_matching_dims.hpp b/stan/math/gpu/err/check_matching_dims.hpp new file mode 100644 index 00000000000..56b4e591e5d --- /dev/null +++ b/stan/math/gpu/err/check_matching_dims.hpp @@ -0,0 +1,33 @@ +#ifndef STAN_MATH_GPU_ERR_CHECK_MATCHING_DIMS_HPP +#define STAN_MATH_GPU_ERR_CHECK_MATCHING_DIMS_HPP +#ifdef STAN_OPENCL +#include +#include + +namespace stan { +namespace math { +/** + * Check if two matrix_gpus have the same dimensions. + * + * @param function Function name (for error messages) + * @param name1 Variable name for the first matrix (for error messages) + * @param y1 First matrix_gpu + * @param name2 Variable name for the second matrix (for error messages) + * @param y2 Second matrix_gpu + * + * @throw std::invalid_argument + * if the dimensions of the matrices do not match + */ +inline void check_matching_dims(const char* function, const char* name1, + const matrix_gpu& y1, const char* name2, + const matrix_gpu& y2) { + check_size_match(function, "Rows of ", name1, y1.rows(), "rows of ", name2, + y2.rows()); + check_size_match(function, "Columns of ", name1, y1.cols(), "columns of ", + name2, y2.cols()); +} + +} // namespace math +} // namespace stan +#endif +#endif diff --git a/stan/math/gpu/err/check_matrix_gpu.hpp b/stan/math/gpu/err/check_matrix_gpu.hpp deleted file mode 100644 index 178c939fe15..00000000000 --- a/stan/math/gpu/err/check_matrix_gpu.hpp +++ /dev/null @@ -1,184 +0,0 @@ -#ifndef STAN_MATH_PRIM_MAT_ERR_CHECK_GPU_HPP -#define STAN_MATH_PRIM_MAT_ERR_CHECK_GPU_HPP -#ifdef STAN_OPENCL -#include -#include -#include -#include -#include -#include -#include - -namespace stan { -namespace math { -/** - * Check if the matrix_gpu is square. - * - * @param function Function name (for error messages) - * @param name Variable name (for error messages) - * @param y matrix_gpu to test - * - * @throw std::invalid_argument if the matrix_gpu - * is not square - */ -inline void check_square(const char* function, const char* name, - const matrix_gpu& y) { - check_size_match(function, "Expecting a square matrix; rows of ", name, - y.rows(), "columns of ", name, y.cols()); -} -/** - * Check if the matrix_gpu has NaN values - * - * @param function Function name (for error messages) - * @param name Variable name (for error messages) - * @param y matrix_gpu to test - * - * @throw std::domain_error if - * any element of the matrix is NaN. - */ -inline void check_nan(const char* function, const char* name, - const matrix_gpu& y) { - if (y.size() == 0) - return; - - cl::Kernel kernel_check_nan = opencl_context.get_kernel("check_nan"); - cl::CommandQueue cmd_queue = opencl_context.queue(); - cl::Context& ctx = opencl_context.context(); - try { - int nan_flag = 0; - cl::Buffer buffer_nan_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); - cmd_queue.enqueueWriteBuffer(buffer_nan_flag, CL_TRUE, 0, sizeof(int), - &nan_flag); - - kernel_check_nan.setArg(0, y.buffer()); - kernel_check_nan.setArg(1, y.rows()); - kernel_check_nan.setArg(2, y.cols()); - kernel_check_nan.setArg(3, buffer_nan_flag); - - cmd_queue.enqueueNDRangeKernel(kernel_check_nan, cl::NullRange, - cl::NDRange(y.rows(), y.cols()), - cl::NullRange); - - cmd_queue.enqueueReadBuffer(buffer_nan_flag, CL_TRUE, 0, sizeof(int), - &nan_flag); - // if NaN values were found in the matrix - if (nan_flag) { - domain_error(function, name, "has NaN values", ""); - } - } catch (const cl::Error& e) { - check_opencl_error("nan_check", e); - } -} -/** - * Check if the matrix_gpu is symmetric - * - * @param function Function name (for error messages) - * @param name Variable name (for error messages) - * @param y matrix_gpu to test - * - * @throw std::domain_error if - * the matrix is not symmetric. - */ -inline void check_symmetric(const char* function, const char* name, - const matrix_gpu& y) { - if (y.size() == 0) - return; - check_square(function, name, y); - cl::Kernel kernel_check_symmetric - = opencl_context.get_kernel("check_symmetric"); - cl::CommandQueue cmd_queue = opencl_context.queue(); - cl::Context& ctx = opencl_context.context(); - try { - int symmetric_flag = 0; - cl::Buffer buffer_symmetric_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); - cmd_queue.enqueueWriteBuffer(buffer_symmetric_flag, CL_TRUE, 0, sizeof(int), - &symmetric_flag); - - kernel_check_symmetric.setArg(0, y.buffer()); - kernel_check_symmetric.setArg(1, y.rows()); - kernel_check_symmetric.setArg(2, y.cols()); - kernel_check_symmetric.setArg(3, buffer_symmetric_flag); - kernel_check_symmetric.setArg(4, math::CONSTRAINT_TOLERANCE); - - cmd_queue.enqueueNDRangeKernel(kernel_check_symmetric, cl::NullRange, - cl::NDRange(y.rows(), y.cols()), - cl::NullRange); - - cmd_queue.enqueueReadBuffer(buffer_symmetric_flag, CL_TRUE, 0, sizeof(int), - &symmetric_flag); - // if the matrix is not symmetric - if (symmetric_flag) { - domain_error(function, name, "is not symmetric", ""); - } - } catch (const cl::Error& e) { - check_opencl_error("symmetric_check", e); - } -} -/** - * Check if the matrix_gpu has zeros on the diagonal - * - * @param function Function name (for error messages) - * @param name Variable name (for error messages) - * @param y matrix_gpu to test - * - * @throw std::domain_error if - * any diagonal element of the matrix is zero. - */ -inline void check_diagonal_zeros(const char* function, const char* name, - const matrix_gpu& y) { - if (y.size() == 0) - return; - - cl::Kernel kernel_check_diagonal_zeros - = opencl_context.get_kernel("check_diagonal_zeros"); - cl::CommandQueue cmd_queue = opencl_context.queue(); - cl::Context ctx = opencl_context.context(); - try { - int flag = 0; - cl::Buffer buffer_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); - cmd_queue.enqueueWriteBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &flag); - - kernel_check_diagonal_zeros.setArg(0, y.buffer()); - kernel_check_diagonal_zeros.setArg(1, y.rows()); - kernel_check_diagonal_zeros.setArg(2, y.cols()); - kernel_check_diagonal_zeros.setArg(3, buffer_flag); - - cmd_queue.enqueueNDRangeKernel(kernel_check_diagonal_zeros, cl::NullRange, - cl::NDRange(y.rows(), y.cols()), - cl::NullRange); - - cmd_queue.enqueueReadBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &flag); - // if zeros were found on the diagonal - if (flag) { - domain_error(function, name, "has zeros on the diagonal.", ""); - } - } catch (const cl::Error& e) { - check_opencl_error("diag_zeros_check", e); - } -} - -/** - * Check if two matrix_gpus have the same dimensions. - * - * @param function Function name (for error messages) - * @param name1 Variable name for the first matrix (for error messages) - * @param y1 First matrix_gpu - * @param name2 Variable name for the second matrix (for error messages) - * @param y2 Second matrix_gpu - * - * @throw std::invalid_argument - * if the dimensions of the matrices do not match - */ -inline void check_matching_dims(const char* function, const char* name1, - const matrix_gpu& y1, const char* name2, - const matrix_gpu& y2) { - check_size_match(function, "Rows of ", name1, y1.rows(), "rows of ", name2, - y2.rows()); - check_size_match(function, "Columns of ", name1, y1.cols(), "columns of ", - name2, y2.cols()); -} - -} // namespace math -} // namespace stan -#endif -#endif diff --git a/stan/math/gpu/err/check_nan.hpp b/stan/math/gpu/err/check_nan.hpp new file mode 100644 index 00000000000..27916498e40 --- /dev/null +++ b/stan/math/gpu/err/check_nan.hpp @@ -0,0 +1,56 @@ +#ifndef STAN_MATH_GPU_ERR_CHECK_NAN_HPP +#define STAN_MATH_GPU_ERR_CHECK_NAN_HPP +#ifdef STAN_OPENCL +#include +#include + +namespace stan { +namespace math { +/** + * Check if the matrix_gpu has NaN values + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y matrix_gpu to test + * + * @throw std::domain_error if + * any element of the matrix is NaN. + */ +inline void check_nan(const char* function, const char* name, + const matrix_gpu& y) { + if (y.size() == 0) + return; + + cl::Kernel kernel_check_nan = opencl_context.get_kernel("check_nan"); + cl::CommandQueue cmd_queue = opencl_context.queue(); + cl::Context& ctx = opencl_context.context(); + try { + int nan_flag = 0; + cl::Buffer buffer_nan_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); + cmd_queue.enqueueWriteBuffer(buffer_nan_flag, CL_TRUE, 0, sizeof(int), + &nan_flag); + + kernel_check_nan.setArg(0, y.buffer()); + kernel_check_nan.setArg(1, y.rows()); + kernel_check_nan.setArg(2, y.cols()); + kernel_check_nan.setArg(3, buffer_nan_flag); + + cmd_queue.enqueueNDRangeKernel(kernel_check_nan, cl::NullRange, + cl::NDRange(y.rows(), y.cols()), + cl::NullRange); + + cmd_queue.enqueueReadBuffer(buffer_nan_flag, CL_TRUE, 0, sizeof(int), + &nan_flag); + // if NaN values were found in the matrix + if (nan_flag) { + domain_error(function, name, "has NaN values", ""); + } + } catch (const cl::Error& e) { + check_opencl_error("nan_check", e); + } +} + +} // namespace math +} // namespace stan +#endif +#endif diff --git a/stan/math/gpu/err/check_square.hpp b/stan/math/gpu/err/check_square.hpp new file mode 100644 index 00000000000..eedc78c0552 --- /dev/null +++ b/stan/math/gpu/err/check_square.hpp @@ -0,0 +1,28 @@ +#ifndef STAN_MATH_GPU_ERR_CHECK_SQUARE_HPP +#define STAN_MATH_GPU_ERR_CHECK_SQUARE_HPP +#ifdef STAN_OPENCL +#include +#include + +namespace stan { +namespace math { +/** + * Check if the matrix_gpu is square. + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y matrix_gpu to test + * + * @throw std::invalid_argument if the matrix_gpu + * is not square + */ +inline void check_square(const char* function, const char* name, + const matrix_gpu& y) { + check_size_match(function, "Expecting a square matrix; rows of ", name, + y.rows(), "columns of ", name, y.cols()); +} + +} // namespace math +} // namespace stan +#endif +#endif diff --git a/stan/math/gpu/err/check_symmetric.hpp b/stan/math/gpu/err/check_symmetric.hpp new file mode 100644 index 00000000000..d21bd20625b --- /dev/null +++ b/stan/math/gpu/err/check_symmetric.hpp @@ -0,0 +1,58 @@ +#ifndef STAN_MATH_GPU_ERR_CHECK_SYMMETRIC_HPP +#define STAN_MATH_GPU_ERR_CHECK_SYMMETRIC_HPP +#ifdef STAN_OPENCL +#include +#include + +namespace stan { +namespace math { +/** + * Check if the matrix_gpu is symmetric + * + * @param function Function name (for error messages) + * @param name Variable name (for error messages) + * @param y matrix_gpu to test + * + * @throw std::domain_error if + * the matrix is not symmetric. + */ +inline void check_symmetric(const char* function, const char* name, + const matrix_gpu& y) { + if (y.size() == 0) + return; + check_square(function, name, y); + cl::Kernel kernel_check_symmetric + = opencl_context.get_kernel("check_symmetric"); + cl::CommandQueue cmd_queue = opencl_context.queue(); + cl::Context& ctx = opencl_context.context(); + try { + int symmetric_flag = 0; + cl::Buffer buffer_symmetric_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); + cmd_queue.enqueueWriteBuffer(buffer_symmetric_flag, CL_TRUE, 0, sizeof(int), + &symmetric_flag); + + kernel_check_symmetric.setArg(0, y.buffer()); + kernel_check_symmetric.setArg(1, y.rows()); + kernel_check_symmetric.setArg(2, y.cols()); + kernel_check_symmetric.setArg(3, buffer_symmetric_flag); + kernel_check_symmetric.setArg(4, math::CONSTRAINT_TOLERANCE); + + cmd_queue.enqueueNDRangeKernel(kernel_check_symmetric, cl::NullRange, + cl::NDRange(y.rows(), y.cols()), + cl::NullRange); + + cmd_queue.enqueueReadBuffer(buffer_symmetric_flag, CL_TRUE, 0, sizeof(int), + &symmetric_flag); + // if the matrix is not symmetric + if (symmetric_flag) { + domain_error(function, name, "is not symmetric", ""); + } + } catch (const cl::Error& e) { + check_opencl_error("symmetric_check", e); + } +} + +} // namespace math +} // namespace stan +#endif +#endif diff --git a/test/unit/math/gpu/opencl_check_test.cpp b/test/unit/math/gpu/opencl_check_test.cpp index 2902b4839ec..1c122c9c950 100644 --- a/test/unit/math/gpu/opencl_check_test.cpp +++ b/test/unit/math/gpu/opencl_check_test.cpp @@ -1,13 +1,17 @@ #ifdef STAN_OPENCL #include -#include +#include +#include +#include +#include +#include #include #include using stan::math::check_nan; // ---------- check_nan: matrix tests ---------- -TEST(ErrorHandlingScalarGPU, Check_nan_Matrix) { +TEST(ErrorHandlingScalarGPU, check_nan_Matrix) { const char* function = "check_nan"; Eigen::Matrix x; using stan::math::matrix_gpu; @@ -19,7 +23,7 @@ TEST(ErrorHandlingScalarGPU, Check_nan_Matrix) { << "check_nan should be true with finite xx"; } -TEST(ErrorHandlingScalarGPU, Check_nan_Matrix_quit_nan) { +TEST(ErrorHandlingScalarGPU, check_nan_Matrix_quit_nan) { const char* function = "check_nan"; Eigen::Matrix x; using stan::math::matrix_gpu; @@ -31,7 +35,7 @@ TEST(ErrorHandlingScalarGPU, Check_nan_Matrix_quit_nan) { << "check_nan should throw exception on NaN"; } -TEST(ErrorHandlingScalarGPU, Check_nan_positions) { +TEST(ErrorHandlingScalarGPU, check_nan_positions) { const char* function = "check_nan"; double nan = std::numeric_limits::quiet_NaN(); using stan::math::matrix_gpu; @@ -77,4 +81,18 @@ TEST(ErrorHandlingScalarGPU, check_m_symmetric) { std::domain_error); EXPECT_NO_THROW(check_symmetric(function, "m_symm_mat1", mm_ok)); } + +TEST(ErrorHandlingScalarGPU, check_m_diagonal_zeros) { + const char* function = "check_diagonal_zeros"; + + stan::math::matrix_d m_ok(3, 3); + stan::math::matrix_d m_fail(3, 3); + m_ok << 1, 2, 3, 2, 4, -5, 3, 5, 6; + m_fail << 1, 2, 3, 2, 0, -5, 3, 5, 6; + stan::math::matrix_gpu mm_ok(m_ok); + stan::math::matrix_gpu mm_fail(m_fail); + EXPECT_NO_THROW(check_diagonal_zeros(function, "mm_ok",mm_ok)); + EXPECT_THROW(check_diagonal_zeros(function, "mm_fail",mm_fail), std::domain_error); +} + #endif From 131d95fd94af6728c782ab5db68ae2230c62ca95 Mon Sep 17 00:00:00 2001 From: rok-cesnovar Date: Tue, 19 Jun 2018 11:34:13 +0200 Subject: [PATCH 13/14] change kernel names to is_* --- stan/math/gpu/err/check_diagonal_zeros.hpp | 16 ++++++++++------ stan/math/gpu/err/check_nan.hpp | 2 +- stan/math/gpu/err/check_symmetric.hpp | 6 +++--- .../gpu/kernels/check_diagonal_zeros_kernel.cl | 2 +- stan/math/gpu/kernels/check_nan_kernel.cl | 2 +- stan/math/gpu/kernels/check_symmetric_kernel.cl | 4 ++-- stan/math/gpu/opencl_context.hpp | 6 +++--- test/unit/math/gpu/opencl_check_test.cpp | 5 +++-- 8 files changed, 24 insertions(+), 19 deletions(-) diff --git a/stan/math/gpu/err/check_diagonal_zeros.hpp b/stan/math/gpu/err/check_diagonal_zeros.hpp index 53d801a548e..d09257570a7 100644 --- a/stan/math/gpu/err/check_diagonal_zeros.hpp +++ b/stan/math/gpu/err/check_diagonal_zeros.hpp @@ -22,26 +22,30 @@ inline void check_diagonal_zeros(const char* function, const char* name, return; cl::Kernel kernel_check_diagonal_zeros - = opencl_context.get_kernel("check_diagonal_zeros"); + = opencl_context.get_kernel("is_zero_on_diagonal"); cl::CommandQueue cmd_queue = opencl_context.queue(); cl::Context ctx = opencl_context.context(); try { - int flag = 0; + int zero_on_diagonal_flag = 0; cl::Buffer buffer_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); - cmd_queue.enqueueWriteBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &flag); + cmd_queue.enqueueWriteBuffer(buffer_flag, CL_TRUE, + 0, sizeof(int), + &zero_on_diagonal_flag); kernel_check_diagonal_zeros.setArg(0, y.buffer()); kernel_check_diagonal_zeros.setArg(1, y.rows()); kernel_check_diagonal_zeros.setArg(2, y.cols()); kernel_check_diagonal_zeros.setArg(3, buffer_flag); - cmd_queue.enqueueNDRangeKernel(kernel_check_diagonal_zeros, cl::NullRange, + cmd_queue.enqueueNDRangeKernel(kernel_check_diagonal_zeros, + cl::NullRange, cl::NDRange(y.rows(), y.cols()), cl::NullRange); - cmd_queue.enqueueReadBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &flag); + cmd_queue.enqueueReadBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), + &zero_on_diagonal_flag); // if zeros were found on the diagonal - if (flag) { + if (zero_on_diagonal_flag) { domain_error(function, name, "has zeros on the diagonal.", ""); } } catch (const cl::Error& e) { diff --git a/stan/math/gpu/err/check_nan.hpp b/stan/math/gpu/err/check_nan.hpp index 27916498e40..3b8e36bb317 100644 --- a/stan/math/gpu/err/check_nan.hpp +++ b/stan/math/gpu/err/check_nan.hpp @@ -21,7 +21,7 @@ inline void check_nan(const char* function, const char* name, if (y.size() == 0) return; - cl::Kernel kernel_check_nan = opencl_context.get_kernel("check_nan"); + cl::Kernel kernel_check_nan = opencl_context.get_kernel("is_nan"); cl::CommandQueue cmd_queue = opencl_context.queue(); cl::Context& ctx = opencl_context.context(); try { diff --git a/stan/math/gpu/err/check_symmetric.hpp b/stan/math/gpu/err/check_symmetric.hpp index d21bd20625b..0f2f68043e9 100644 --- a/stan/math/gpu/err/check_symmetric.hpp +++ b/stan/math/gpu/err/check_symmetric.hpp @@ -22,11 +22,11 @@ inline void check_symmetric(const char* function, const char* name, return; check_square(function, name, y); cl::Kernel kernel_check_symmetric - = opencl_context.get_kernel("check_symmetric"); + = opencl_context.get_kernel("is_symmetric"); cl::CommandQueue cmd_queue = opencl_context.queue(); cl::Context& ctx = opencl_context.context(); try { - int symmetric_flag = 0; + int symmetric_flag = 1; cl::Buffer buffer_symmetric_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); cmd_queue.enqueueWriteBuffer(buffer_symmetric_flag, CL_TRUE, 0, sizeof(int), &symmetric_flag); @@ -44,7 +44,7 @@ inline void check_symmetric(const char* function, const char* name, cmd_queue.enqueueReadBuffer(buffer_symmetric_flag, CL_TRUE, 0, sizeof(int), &symmetric_flag); // if the matrix is not symmetric - if (symmetric_flag) { + if (!symmetric_flag) { domain_error(function, name, "is not symmetric", ""); } } catch (const cl::Error& e) { diff --git a/stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl b/stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl index db6fb014e91..64fba761a2b 100644 --- a/stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl +++ b/stan/math/gpu/kernels/check_diagonal_zeros_kernel.cl @@ -2,7 +2,7 @@ R"( #ifndef A #define A(i, j) A[j * rows + i] #endif -__kernel void check_diagonal_zeros( +__kernel void is_zero_on_diagonal( __global double *A, int rows, int cols, diff --git a/stan/math/gpu/kernels/check_nan_kernel.cl b/stan/math/gpu/kernels/check_nan_kernel.cl index f950382ba22..c4bf981d5f9 100644 --- a/stan/math/gpu/kernels/check_nan_kernel.cl +++ b/stan/math/gpu/kernels/check_nan_kernel.cl @@ -2,7 +2,7 @@ R"( #ifndef A #define A(i, j) A[j * rows + i] #endif -__kernel void check_nan( +__kernel void is_nan( __global double *A, int rows, int cols, diff --git a/stan/math/gpu/kernels/check_symmetric_kernel.cl b/stan/math/gpu/kernels/check_symmetric_kernel.cl index 9ec734d7ed9..6940e10a1ae 100644 --- a/stan/math/gpu/kernels/check_symmetric_kernel.cl +++ b/stan/math/gpu/kernels/check_symmetric_kernel.cl @@ -2,7 +2,7 @@ R"( #ifndef A #define A(i, j) A[j * rows + i] #endif -__kernel void check_symmetric( +__kernel void is_symmetric( __global double *A, int rows, int cols, @@ -13,7 +13,7 @@ __kernel void check_symmetric( if (i < rows && j < cols) { double diff = fabs(A(i, j) - A(j, i)); if (diff > tolerance) { - flag[0] = 1; + flag[0] = 0; } } };)" diff --git a/stan/math/gpu/opencl_context.hpp b/stan/math/gpu/opencl_context.hpp index b1d20a919a7..34499d90c62 100644 --- a/stan/math/gpu/opencl_context.hpp +++ b/stan/math/gpu/opencl_context.hpp @@ -127,10 +127,10 @@ class opencl_context_base { kernel_info["dummy2"] = { false, "timing", "__kernel void dummy2(__global const int* foo) { };"}; kernel_info["copy"] = {false, "basic_matrix", copy_matrix_kernel}; - kernel_info["check_nan"] = {false, "check", check_nan_kernel}; - kernel_info["check_diagonal_zeros"] + kernel_info["is_nan"] = {false, "check", check_nan_kernel}; + kernel_info["is_zero_on_diagonal"] = {false, "check", check_diagonal_zeros_kernel}; - kernel_info["check_symmetric"] = {false, "check", check_symmetric_kernel}; + kernel_info["is_symmetric"] = {false, "check", check_symmetric_kernel}; } protected: diff --git a/test/unit/math/gpu/opencl_check_test.cpp b/test/unit/math/gpu/opencl_check_test.cpp index 1c122c9c950..4e20b9182ac 100644 --- a/test/unit/math/gpu/opencl_check_test.cpp +++ b/test/unit/math/gpu/opencl_check_test.cpp @@ -91,8 +91,9 @@ TEST(ErrorHandlingScalarGPU, check_m_diagonal_zeros) { m_fail << 1, 2, 3, 2, 0, -5, 3, 5, 6; stan::math::matrix_gpu mm_ok(m_ok); stan::math::matrix_gpu mm_fail(m_fail); - EXPECT_NO_THROW(check_diagonal_zeros(function, "mm_ok",mm_ok)); - EXPECT_THROW(check_diagonal_zeros(function, "mm_fail",mm_fail), std::domain_error); + EXPECT_NO_THROW(check_diagonal_zeros(function, "mm_ok", mm_ok)); + EXPECT_THROW(check_diagonal_zeros(function, "mm_fail", mm_fail), + std::domain_error); } #endif From b132adcd511a1d7c1f4f8aa003647941604bdf3d Mon Sep 17 00:00:00 2001 From: Stan Jenkins Date: Tue, 19 Jun 2018 14:11:06 -0400 Subject: [PATCH 14/14] [Jenkins] auto-formatting by clang-format version 6.0.0 (tags/google/stable/2017-11-14) --- stan/math/gpu/err/check_diagonal_zeros.hpp | 6 ++---- stan/math/gpu/err/check_symmetric.hpp | 3 +-- test/unit/math/gpu/opencl_check_test.cpp | 2 +- 3 files changed, 4 insertions(+), 7 deletions(-) diff --git a/stan/math/gpu/err/check_diagonal_zeros.hpp b/stan/math/gpu/err/check_diagonal_zeros.hpp index d09257570a7..a5a5283cd73 100644 --- a/stan/math/gpu/err/check_diagonal_zeros.hpp +++ b/stan/math/gpu/err/check_diagonal_zeros.hpp @@ -28,8 +28,7 @@ inline void check_diagonal_zeros(const char* function, const char* name, try { int zero_on_diagonal_flag = 0; cl::Buffer buffer_flag(ctx, CL_MEM_READ_WRITE, sizeof(int)); - cmd_queue.enqueueWriteBuffer(buffer_flag, CL_TRUE, - 0, sizeof(int), + cmd_queue.enqueueWriteBuffer(buffer_flag, CL_TRUE, 0, sizeof(int), &zero_on_diagonal_flag); kernel_check_diagonal_zeros.setArg(0, y.buffer()); @@ -37,8 +36,7 @@ inline void check_diagonal_zeros(const char* function, const char* name, kernel_check_diagonal_zeros.setArg(2, y.cols()); kernel_check_diagonal_zeros.setArg(3, buffer_flag); - cmd_queue.enqueueNDRangeKernel(kernel_check_diagonal_zeros, - cl::NullRange, + cmd_queue.enqueueNDRangeKernel(kernel_check_diagonal_zeros, cl::NullRange, cl::NDRange(y.rows(), y.cols()), cl::NullRange); diff --git a/stan/math/gpu/err/check_symmetric.hpp b/stan/math/gpu/err/check_symmetric.hpp index 0f2f68043e9..17351f39b86 100644 --- a/stan/math/gpu/err/check_symmetric.hpp +++ b/stan/math/gpu/err/check_symmetric.hpp @@ -21,8 +21,7 @@ inline void check_symmetric(const char* function, const char* name, if (y.size() == 0) return; check_square(function, name, y); - cl::Kernel kernel_check_symmetric - = opencl_context.get_kernel("is_symmetric"); + cl::Kernel kernel_check_symmetric = opencl_context.get_kernel("is_symmetric"); cl::CommandQueue cmd_queue = opencl_context.queue(); cl::Context& ctx = opencl_context.context(); try { diff --git a/test/unit/math/gpu/opencl_check_test.cpp b/test/unit/math/gpu/opencl_check_test.cpp index 4e20b9182ac..4b2a5a95a2a 100644 --- a/test/unit/math/gpu/opencl_check_test.cpp +++ b/test/unit/math/gpu/opencl_check_test.cpp @@ -93,7 +93,7 @@ TEST(ErrorHandlingScalarGPU, check_m_diagonal_zeros) { stan::math::matrix_gpu mm_fail(m_fail); EXPECT_NO_THROW(check_diagonal_zeros(function, "mm_ok", mm_ok)); EXPECT_THROW(check_diagonal_zeros(function, "mm_fail", mm_fail), - std::domain_error); + std::domain_error); } #endif