From 1be77dd2f3f522ce0473fb313fc06f9ceb4cb7fa Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Fri, 11 Oct 2013 00:00:01 +0400 Subject: [PATCH] enabled accuracy tests for the functions that use inside AMD Blas/Fft --- modules/core/include/opencv2/core/types_c.h | 3 ++- modules/ocl/CMakeLists.txt | 2 +- modules/ocl/src/fft.cpp | 2 +- modules/ocl/src/gemm.cpp | 4 ++-- modules/ocl/src/opencl/filtering_laplacian.cl | 6 +++--- modules/ocl/src/opencl/stereobp.cl | 4 ++-- modules/ocl/test/test_fft.cpp | 5 ----- modules/ocl/test/test_gemm.cpp | 6 +++--- modules/ocl/test/test_kalman.cpp | 4 ---- modules/ocl/test/test_ml.cpp | 3 --- modules/ocl/test/utility.hpp | 8 +++++--- 11 files changed, 19 insertions(+), 28 deletions(-) diff --git a/modules/core/include/opencv2/core/types_c.h b/modules/core/include/opencv2/core/types_c.h index fca2d4697f..3e5d5b0334 100644 --- a/modules/core/include/opencv2/core/types_c.h +++ b/modules/core/include/opencv2/core/types_c.h @@ -266,7 +266,8 @@ enum { CV_OpenGlNotSupported= -218, CV_OpenGlApiCallError= -219, CV_OpenCLDoubleNotSupported= -220, - CV_OpenCLInitError= -221 + CV_OpenCLInitError= -221, + CV_OpenCLNoAMDBlasFft= -222 }; /****************************************************************************************\ diff --git a/modules/ocl/CMakeLists.txt b/modules/ocl/CMakeLists.txt index 0303d68e78..21e0b30858 100644 --- a/modules/ocl/CMakeLists.txt +++ b/modules/ocl/CMakeLists.txt @@ -4,5 +4,5 @@ if(NOT HAVE_OPENCL) endif() set(the_description "OpenCL-accelerated Computer Vision") -ocv_define_module(ocl opencv_core opencv_imgproc opencv_features2d opencv_objdetect opencv_video opencv_calib3d opencv_ml) +ocv_define_module(ocl opencv_core opencv_imgproc opencv_features2d opencv_objdetect opencv_video opencv_calib3d opencv_ml "${OPENCL_LIBRARIES}") ocv_warnings_disable(CMAKE_CXX_FLAGS -Wshadow) diff --git a/modules/ocl/src/fft.cpp b/modules/ocl/src/fft.cpp index e39a4443c4..93029bc85f 100644 --- a/modules/ocl/src/fft.cpp +++ b/modules/ocl/src/fft.cpp @@ -50,7 +50,7 @@ using namespace cv::ocl; #if !defined HAVE_CLAMDFFT void cv::ocl::dft(const oclMat&, oclMat&, Size, int) { - CV_Error(CV_StsNotImplemented, "OpenCL DFT is not implemented"); + CV_Error(CV_OpenCLNoAMDBlasFft, "OpenCL DFT is not implemented"); } namespace cv { namespace ocl { void fft_teardown(); diff --git a/modules/ocl/src/gemm.cpp b/modules/ocl/src/gemm.cpp index 837fd1fa30..bd3e7e5b8f 100644 --- a/modules/ocl/src/gemm.cpp +++ b/modules/ocl/src/gemm.cpp @@ -58,12 +58,12 @@ void clBlasTeardown(); void cv::ocl::gemm(const oclMat&, const oclMat&, double, const oclMat&, double, oclMat&, int) { - CV_Error(CV_StsNotImplemented, "OpenCL BLAS is not implemented"); + CV_Error(CV_OpenCLNoAMDBlasFft, "OpenCL BLAS is not implemented"); } void cv::ocl::clBlasSetup() { - CV_Error(CV_StsNotImplemented, "OpenCL BLAS is not implemented"); + CV_Error(CV_OpenCLNoAMDBlasFft, "OpenCL BLAS is not implemented"); } void cv::ocl::clBlasTeardown() diff --git a/modules/ocl/src/opencl/filtering_laplacian.cl b/modules/ocl/src/opencl/filtering_laplacian.cl index f7430d5332..3c0cc0de38 100644 --- a/modules/ocl/src/opencl/filtering_laplacian.cl +++ b/modules/ocl/src/opencl/filtering_laplacian.cl @@ -211,7 +211,7 @@ __kernel void filter2D( barrier(CLK_LOCAL_MEM_FENCE); if(globalRow < rows && globalCol < cols) { - T_SUM sum = (T_SUM)SUM_ZERO; + T_SUM sum = (T_SUM)(SUM_ZERO); int filterIdx = 0; for(int i = 0; i < FILTER_SIZE; i++) { @@ -291,7 +291,7 @@ __kernel void filter2D_3x3( T_IMG data = src[mad24(selected_row, src_step, selected_cols)]; int con = selected_row >= 0 && selected_row < wholerows && selected_cols >= 0 && selected_cols < wholecols; - data = con ? data : 0; + data = con ? data : (T_IMG)(0); local_data[mad24(i, LOCAL_MEM_STEP, lX)] = data; if(lX < (ANX << 1)) @@ -300,7 +300,7 @@ __kernel void filter2D_3x3( data = src[mad24(selected_row, src_step, selected_cols)]; con = selected_row >= 0 && selected_row < wholerows && selected_cols >= 0 && selected_cols < wholecols; - data = con ? data : 0; + data = con ? data : (T_IMG)(0); local_data[mad24(i, LOCAL_MEM_STEP, lX) + groupX_size] = data; } #else diff --git a/modules/ocl/src/opencl/stereobp.cl b/modules/ocl/src/opencl/stereobp.cl index 1d523e7885..24bf55cb21 100644 --- a/modules/ocl/src/opencl/stereobp.cl +++ b/modules/ocl/src/opencl/stereobp.cl @@ -290,7 +290,7 @@ void message(__global T *us_, __global T *ds_, __global T *ls_, __global T *rs_, minimum += cmax_disc_term; - float4 sum = 0; + float4 sum = (float4)(0); prev = convert_float4(t_dst[CNDISP - 1]); for (int disp = CNDISP - 2; disp >= 0; disp--) { @@ -308,7 +308,7 @@ void message(__global T *us_, __global T *ds_, __global T *ls_, __global T *rs_, t_dst[CNDISP - 1] = saturate_cast4(dst_reg); sum += dst_reg; - sum /= CNDISP; + sum /= (float4)(CNDISP); #pragma unroll for(int i = 0, idx = 0; i < CNDISP; ++i, idx+=msg_disp_step) { diff --git a/modules/ocl/test/test_fft.cpp b/modules/ocl/test/test_fft.cpp index 6a9878feb5..fc8a0f72af 100644 --- a/modules/ocl/test/test_fft.cpp +++ b/modules/ocl/test/test_fft.cpp @@ -47,8 +47,6 @@ using namespace std; -#ifdef HAVE_CLAMDFFT - //////////////////////////////////////////////////////////////////////////// // Dft @@ -102,9 +100,6 @@ OCL_TEST_P(Dft, R2CthenC2R) EXPECT_MAT_NEAR(a, d_c, a.size().area() * 1e-4); } - INSTANTIATE_TEST_CASE_P(OCL_ImgProc, Dft, testing::Combine( testing::Values(cv::Size(2, 3), cv::Size(5, 4), cv::Size(25, 20), cv::Size(512, 1), cv::Size(1024, 768)), testing::Values(0, (int)cv::DFT_ROWS, (int)cv::DFT_SCALE) )); - -#endif // HAVE_CLAMDFFT diff --git a/modules/ocl/test/test_gemm.cpp b/modules/ocl/test/test_gemm.cpp index 376090d9d1..68dab0ac01 100644 --- a/modules/ocl/test/test_gemm.cpp +++ b/modules/ocl/test/test_gemm.cpp @@ -42,12 +42,13 @@ // //M*/ - #include "test_precomp.hpp" + using namespace std; -#ifdef HAVE_CLAMDBLAS + //////////////////////////////////////////////////////////////////////////// // GEMM + PARAM_TEST_CASE(Gemm, int, cv::Size, int) { int type; @@ -81,4 +82,3 @@ INSTANTIATE_TEST_CASE_P(ocl_gemm, Gemm, testing::Combine( testing::Values(CV_32FC1, CV_32FC2/*, CV_64FC1, CV_64FC2*/), testing::Values(cv::Size(20, 20), cv::Size(300, 300)), testing::Values(0, (int)cv::GEMM_1_T, (int)cv::GEMM_2_T, (int)(cv::GEMM_1_T + cv::GEMM_2_T)))); -#endif diff --git a/modules/ocl/test/test_kalman.cpp b/modules/ocl/test/test_kalman.cpp index 6cbeab03f6..f02df6af71 100644 --- a/modules/ocl/test/test_kalman.cpp +++ b/modules/ocl/test/test_kalman.cpp @@ -46,8 +46,6 @@ #ifdef HAVE_OPENCL -#ifdef HAVE_CLAMDBLAS - using namespace cv; using namespace cv::ocl; using namespace cvtest; @@ -147,6 +145,4 @@ OCL_TEST_P(Kalman, Accuracy) INSTANTIATE_TEST_CASE_P(OCL_Video, Kalman, Combine(Values(3, 7), Values(30))); -#endif // HAVE_CLAMDBLAS - #endif // HAVE_OPENCL diff --git a/modules/ocl/test/test_ml.cpp b/modules/ocl/test/test_ml.cpp index 12518f4fb5..52d40bb11c 100644 --- a/modules/ocl/test/test_ml.cpp +++ b/modules/ocl/test/test_ml.cpp @@ -128,8 +128,6 @@ INSTANTIATE_TEST_CASE_P(OCL_ML, KNN, Combine(Values(6, 5), Values(Size(200, 400) ////////////////////////////////SVM///////////////////////////////////////////////// -#ifdef HAVE_CLAMDBLAS - PARAM_TEST_CASE(SVM_OCL, int, int, int) { cv::Size size; @@ -307,6 +305,5 @@ INSTANTIATE_TEST_CASE_P(OCL_ML, SVM_OCL, testing::Combine( Values((int)CvSVM::C_SVC, (int)CvSVM::NU_SVC, (int)CvSVM::ONE_CLASS, (int)CvSVM::NU_SVR), Values(2, 3, 4) )); -#endif // HAVE_CLAMDBLAS #endif // HAVE_OPENCL diff --git a/modules/ocl/test/utility.hpp b/modules/ocl/test/utility.hpp index 0a91d57d83..7501b19fc0 100644 --- a/modules/ocl/test/utility.hpp +++ b/modules/ocl/test/utility.hpp @@ -291,10 +291,12 @@ CV_FLAGS(DftFlags, DFT_INVERSE, DFT_SCALE, DFT_ROWS, DFT_COMPLEX_OUTPUT, DFT_REA } \ catch (const cv::Exception & ex) \ { \ - if (ex.code != CV_OpenCLDoubleNotSupported) \ - throw; \ - else \ + if (ex.code == CV_OpenCLDoubleNotSupported)\ std::cout << "Test skipped (selected device does not support double)" << std::endl; \ + else if (ex.code == CV_OpenCLNoAMDBlasFft) \ + std::cout << "Test skipped (AMD Blas / Fft libraries are not available)" << std::endl; \ + else \ + throw; \ } \ } \ \