Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

performance fix of convertC3C4

Browse Source 
add OCL 1.2 feature for setTo
bug fix of integral
replace the error code -217 with suitable MACRO
simplify tests, no need apply a new context for each test case
add more control for tests in utility.hpp
pull/26/head
niko 13 years ago
parent 44330afc7f
commit b929012583
37 changed files with 717 additions and 1820 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 2
      cmake/OpenCVDetectOpenCL.cmake
  2. 21
      modules/ocl/include/opencv2/ocl/ocl.hpp
  3. 25
      modules/ocl/perf/main.cpp
  4. 1191
      modules/ocl/perf/perf_test_ocl.cpp
  5. 229
      modules/ocl/perf/test_arithm.cpp
  6. 114
      modules/ocl/perf/test_filters.cpp
  7. 12
      modules/ocl/perf/test_haar.cpp
  8. 134
      modules/ocl/perf/test_imgproc.cpp
  9. 56
      modules/ocl/perf/test_matrix_operation.cpp
  10. 32
      modules/ocl/perf/test_split_merge.cpp
  11. 6
      modules/ocl/perf/utility.hpp
  12. 34
      modules/ocl/src/arithm.cpp
  13. 6
      modules/ocl/src/filtering.cpp
  14. 143
      modules/ocl/src/haar.cpp
  15. 8
      modules/ocl/src/imgproc.cpp
  16. 26
      modules/ocl/src/initialization.cpp
  17. 90
      modules/ocl/src/kernels/convertC3C4.cl
  18. 20
      modules/ocl/src/kernels/imgproc_integral.cl
  19. 15
      modules/ocl/src/kernels/imgproc_integral_sum.cl
  20. 53
      modules/ocl/src/matrix_operations.cpp
  21. 4
      modules/ocl/src/split_merge.cpp
  22. 15
      modules/ocl/test/main.cpp
  23. 11
      modules/ocl/test/test_arithm.cpp
  24. 6
      modules/ocl/test/test_blend.cpp
  25. 6
      modules/ocl/test/test_canny.cpp
  26. 6
      modules/ocl/test/test_columnsum.cpp
  27. 6
      modules/ocl/test/test_fft.cpp
  28. 72
      modules/ocl/test/test_filters.cpp
  29. 4
      modules/ocl/test/test_gemm.cpp
  30. 12
      modules/ocl/test/test_haar.cpp
  31. 4
      modules/ocl/test/test_hog.cpp
  32. 86
      modules/ocl/test/test_imgproc.cpp
  33. 12
      modules/ocl/test/test_match_template.cpp
  34. 40
      modules/ocl/test/test_matrix_operation.cpp
  35. 10
      modules/ocl/test/test_pyrdown.cpp
  36. 6
      modules/ocl/test/test_pyrup.cpp
  37. 20
      modules/ocl/test/test_split_merge.cpp

2
cmake/OpenCVDetectOpenCL.cmake
Unescape View File

@ -2,7 +2,7 @@ if(APPLE)
set(OPENCL_FOUND YES)
set(OPENCL_LIBRARIES "-framework OpenCL")
else()
#find_package(OpenCL QUIET)
find_package(OpenCL QUIET)
if(WITH_OPENCLAMDFFT)
find_path(CLAMDFFT_INCLUDE_DIR
NAMES clAmdFft.h)

21
modules/ocl/include/opencv2/ocl/ocl.hpp
Unescape View File

@ -144,32 +144,15 @@ namespace cv
//! assignment operator. Perfom blocking upload to device.
oclMat &operator = (const Mat &m);
/* Fixme! To be supported in OpenCL later. */
#if 0
//! returns lightweight DevMem2D_ structure for passing to nvcc-compiled code.
// Contains just image size, data ptr and step.
template <class T> operator DevMem2D_<T>() const;
template <class T> operator PtrStep_<T>() const;
#endif
//! pefroms blocking upload data to oclMat.
void upload(const cv::Mat &m);
/* Fixme! To be supported in OpenCL later. */
#if 0
//! upload async
void upload(const CudaMem &m, Stream &stream);
#endif
//! downloads data from device to host memory. Blocking calls.
operator Mat() const;
void download(cv::Mat &m) const;
/* Fixme! To be supported in OpenCL later. */
#if 0
//! download async
void download(CudaMem &m, Stream &stream) const;
#endif
//! returns a new oclMatrix header for the specified row
oclMat row(int y) const;
@ -855,10 +838,6 @@ namespace cv
int minNeighbors, int flags, CvSize minSize = cvSize(0, 0), CvSize maxSize = cvSize(0, 0));
};
///////////////////////////////////////////////////////jhp_benchmark////////////////////////////////////////////////////
void benchmark_copy_vectorize(const oclMat &src, oclMat &dst);
void benchmark_copy_offset_stride(const oclMat &src, oclMat &dst);
void benchmark_ILP();
//! computes vertical sum, supports only CV_32FC1 images
CV_EXPORTS void columnSum(const oclMat& src, oclMat& sum);

25
modules/ocl/perf/main.cpp
Unescape View File

@ -74,29 +74,26 @@ void print_info()
}
#if PERF_TEST_OCL
int main(int argc, char** argv)
{
static std::vector<Info> ocl_info;
ocl::getDevice(ocl_info);
run_perf_test();
return 0;
}
#else
int main(int argc, char** argv)
{
std::vector<cv::ocl::Info> oclinfo;
TS::ptr()->init("ocl");
InitGoogleTest(&argc, argv);
print_info();
int devnums = getDevice(oclinfo);
if(devnums<1)
{
std::cout << "no device found\n";
return -1;
}
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
setBinpath(CLBINPATH);
return RUN_ALL_TESTS();
}
#endif // PERF_TEST_OCL
#else // HAVE_OPENC
#else // DON'T HAVE_OPENCL
int main()
{

1191
modules/ocl/perf/perf_test_ocl.cpp
View File

File diff suppressed because it is too large Load Diff

229
modules/ocl/perf/test_arithm.cpp
Unescape View File

@ -89,7 +89,7 @@ PARAM_TEST_CASE(ArithmTestBase, MatType, bool)
cv::Mat mask_roi;
cv::Mat dst_roi;
cv::Mat dst1_roi; //bak
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
cv::ocl::oclMat gdst1_whole; //bak
@ -119,11 +119,11 @@ PARAM_TEST_CASE(ArithmTestBase, MatType, bool)
cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
val = cv::Scalar(rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0));
int devnums = getDevice(oclinfo);
CV_Assert(devnums>0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums>0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -193,7 +193,7 @@ TEST_P(Lut, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -235,7 +235,7 @@ TEST_P(Lut, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
// src2x = rng.uniform( 0,mat2.cols - 256);
@ -275,7 +275,7 @@ TEST_P(Exp, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -313,7 +313,7 @@ TEST_P(Exp, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -341,7 +341,7 @@ TEST_P(Log, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -377,7 +377,7 @@ TEST_P(Log, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -408,7 +408,7 @@ TEST_P(Add, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -446,7 +446,7 @@ TEST_P(Add, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -468,7 +468,7 @@ TEST_P(Add, Mat_Mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -506,7 +506,7 @@ TEST_P(Add, Mat_Mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -528,7 +528,7 @@ TEST_P(Add, Scalar)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -563,7 +563,7 @@ TEST_P(Add, Scalar)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -584,7 +584,7 @@ TEST_P(Add, Scalar_Mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -620,7 +620,7 @@ TEST_P(Add, Scalar_Mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -646,7 +646,7 @@ TEST_P(Sub, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -683,7 +683,7 @@ TEST_P(Sub, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -705,7 +705,7 @@ TEST_P(Sub, Mat_Mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -743,7 +743,7 @@ TEST_P(Sub, Mat_Mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -765,7 +765,7 @@ TEST_P(Sub, Scalar)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -801,7 +801,7 @@ TEST_P(Sub, Scalar)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -822,7 +822,7 @@ TEST_P(Sub, Scalar_Mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -859,7 +859,7 @@ TEST_P(Sub, Scalar_Mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -885,7 +885,7 @@ TEST_P(Mul, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -922,7 +922,7 @@ TEST_P(Mul, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -944,7 +944,7 @@ TEST_P(Mul, Mat_Scalar)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -982,7 +982,7 @@ TEST_P(Mul, Mat_Scalar)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
cv::RNG& rng = TS::ptr()->get_rng();
@ -1009,7 +1009,7 @@ TEST_P(Div, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1046,7 +1046,7 @@ TEST_P(Div, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1068,7 +1068,7 @@ TEST_P(Div, Mat_Scalar)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1106,7 +1106,7 @@ TEST_P(Div, Mat_Scalar)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
cv::RNG& rng = TS::ptr()->get_rng();
@ -1134,7 +1134,7 @@ TEST_P(Absdiff, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1171,7 +1171,7 @@ TEST_P(Absdiff, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1193,7 +1193,7 @@ TEST_P(Absdiff, Mat_Scalar)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1229,7 +1229,7 @@ TEST_P(Absdiff, Mat_Scalar)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1254,7 +1254,7 @@ TEST_P(CartToPolar, angleInDegree)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1295,7 +1295,7 @@ TEST_P(CartToPolar, angleInDegree)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1319,7 +1319,7 @@ TEST_P(CartToPolar, angleInRadians)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1359,7 +1359,7 @@ TEST_P(CartToPolar, angleInRadians)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1386,7 +1386,7 @@ TEST_P(PolarToCart, angleInDegree)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1427,7 +1427,7 @@ TEST_P(PolarToCart, angleInDegree)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1451,7 +1451,7 @@ TEST_P(PolarToCart, angleInRadians)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1492,7 +1492,7 @@ TEST_P(PolarToCart, angleInRadians)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1520,7 +1520,7 @@ TEST_P(Magnitude, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1557,7 +1557,7 @@ TEST_P(Magnitude, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1581,7 +1581,7 @@ TEST_P(Transpose, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1617,7 +1617,7 @@ TEST_P(Transpose, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1641,7 +1641,7 @@ TEST_P(Flip, X)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1677,7 +1677,7 @@ TEST_P(Flip, X)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1698,7 +1698,7 @@ TEST_P(Flip, Y)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1734,7 +1734,7 @@ TEST_P(Flip, Y)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1755,7 +1755,7 @@ TEST_P(Flip, BOTH)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1791,7 +1791,7 @@ TEST_P(Flip, BOTH)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1816,7 +1816,7 @@ TEST_P(MinMax, MAT)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1865,7 +1865,7 @@ TEST_P(MinMax, MAT)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gmat1 = mat1_roi;
@ -1885,7 +1885,7 @@ TEST_P(MinMax, MASK)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1936,7 +1936,7 @@ TEST_P(MinMax, MASK)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gmat1 = mat1_roi;
@ -1960,7 +1960,7 @@ TEST_P(MinMaxLoc, MAT)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2020,7 +2020,7 @@ TEST_P(MinMaxLoc, MAT)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gmat1 = mat1_roi;
@ -2044,7 +2044,7 @@ TEST_P(MinMaxLoc, MASK)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2106,7 +2106,7 @@ TEST_P(MinMaxLoc, MASK)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gmat1 = mat1_roi;
@ -2132,7 +2132,7 @@ TEST_P(Sum, MAT)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2163,7 +2163,7 @@ TEST_P(Sum, MAT)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gmat1 = mat1_roi;
@ -2192,7 +2192,7 @@ TEST_P(CountNonZero, MAT)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2223,7 +2223,7 @@ TEST_P(CountNonZero, MAT)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gmat1 = mat1_roi;
@ -2253,7 +2253,7 @@ TEST_P(Phase, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2290,7 +2290,7 @@ TEST_P(Phase, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2318,7 +2318,7 @@ TEST_P(Bitwise_and, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2355,7 +2355,7 @@ TEST_P(Bitwise_and, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2379,7 +2379,7 @@ TEST_P(Bitwise_and, Mat_Mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2417,7 +2417,7 @@ TEST_P(Bitwise_and, Mat_Mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2441,7 +2441,7 @@ TEST_P(Bitwise_and, Scalar)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2477,7 +2477,7 @@ TEST_P(Bitwise_and, Scalar)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2500,7 +2500,7 @@ TEST_P(Bitwise_and, Scalar_Mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2537,7 +2537,7 @@ TEST_P(Bitwise_and, Scalar_Mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2567,7 +2567,7 @@ TEST_P(Bitwise_or, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2604,7 +2604,7 @@ TEST_P(Bitwise_or, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2628,7 +2628,7 @@ TEST_P(Bitwise_or, Mat_Mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2666,7 +2666,7 @@ TEST_P(Bitwise_or, Mat_Mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2689,7 +2689,7 @@ TEST_P(Bitwise_or, Scalar)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2725,7 +2725,7 @@ TEST_P(Bitwise_or, Scalar)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2747,7 +2747,7 @@ TEST_P(Bitwise_or, Scalar_Mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2784,7 +2784,7 @@ TEST_P(Bitwise_or, Scalar_Mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2812,7 +2812,7 @@ TEST_P(Bitwise_xor, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2849,7 +2849,7 @@ TEST_P(Bitwise_xor, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2872,7 +2872,7 @@ TEST_P(Bitwise_xor, Mat_Mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2910,7 +2910,7 @@ TEST_P(Bitwise_xor, Mat_Mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2934,7 +2934,7 @@ TEST_P(Bitwise_xor, Scalar)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -2970,7 +2970,7 @@ TEST_P(Bitwise_xor, Scalar)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -2992,7 +2992,7 @@ TEST_P(Bitwise_xor, Scalar_Mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -3029,7 +3029,7 @@ TEST_P(Bitwise_xor, Scalar_Mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -3057,7 +3057,7 @@ TEST_P(Bitwise_not, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -3093,7 +3093,7 @@ TEST_P(Bitwise_not, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -3138,7 +3138,7 @@ PARAM_TEST_CASE ( CompareTestBase, MatType, bool)
cv::Mat mask_roi;
cv::Mat dst_roi;
cv::Mat dst1_roi; //bak
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
cv::ocl::oclMat gdst1_whole; //bak
@ -3169,11 +3169,11 @@ PARAM_TEST_CASE ( CompareTestBase, MatType, bool)
cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
val = cv::Scalar(rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0));
int devnums = getDevice(oclinfo);
CV_Assert(devnums>0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums>0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -3235,7 +3235,7 @@ TEST_P(Compare, Mat)
}
int cmp_codes[] = {CMP_EQ, CMP_GT, CMP_GE, CMP_LT, CMP_LE, CMP_NE};
//const char* cmp_str[] = {"CMP_EQ", "CMP_GT", "CMP_GE", "CMP_LT", "CMP_LE", "CMP_NE"};
const char* cmp_str[] = {"CMP_EQ", "CMP_GT", "CMP_GE", "CMP_LT", "CMP_LE", "CMP_NE"};
int cmp_num = sizeof(cmp_codes) / sizeof(int);
for (int i = 0; i < cmp_num; ++i)
{
@ -3247,7 +3247,7 @@ TEST_P(Compare, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=1;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -3278,13 +3278,14 @@ TEST_P(Compare, Mat)
totalgputick_kernel=t2+totalgputick_kernel;
}
cout<<cmp_str[i] <<endl;
if(k==0){cout<<"no roi\n";}else{cout<<"with roi\n";};
cout << "average cpu runtime is " << totalcputick/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
cout << "average gpu runtime is " << totalgputick/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -3315,7 +3316,7 @@ TEST_P(Pow, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -3351,7 +3352,7 @@ TEST_P(Pow, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
double p=4.5;
@ -3377,7 +3378,7 @@ TEST_P(MagnitudeSqr, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -3418,7 +3419,7 @@ TEST_P(MagnitudeSqr, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
cv::ocl::oclMat clmat1(mat1),clmat2(mat2),cldst;
@ -3441,7 +3442,7 @@ TEST_P(AddWeighted, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -3482,7 +3483,7 @@ TEST_P(AddWeighted, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
double alpha=2.0,beta=1.0,gama=3.0;
@ -3541,7 +3542,7 @@ TEST_P(AddWeighted, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
#else
//for(int j = 0; j < 2; j ++)
//for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
// {
double alpha=2.0,beta=1.0,gama=3.0;
cv::ocl::oclMat clmat1(mat1),clmat2(mat2),cldst;

114
modules/ocl/perf/test_filters.cpp
Unescape View File

@ -85,7 +85,7 @@ PARAM_TEST_CASE(FilterTestBase, MatType, bool)
cv::Mat mask_roi;
cv::Mat dst_roi;
cv::Mat dst1_roi; //bak
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
cv::ocl::oclMat gdst1_whole; //bak
@ -174,7 +174,7 @@ PARAM_TEST_CASE(Blur, MatType, cv::Size, int)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -193,11 +193,11 @@ PARAM_TEST_CASE(Blur, MatType, cv::Size, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath(CLBINPATH);
}
@ -237,7 +237,7 @@ TEST_P(Blur, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -275,7 +275,7 @@ TEST_P(Blur, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -329,11 +329,11 @@ PARAM_TEST_CASE(LaplacianTestBase, MatType, int)
mat = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -375,7 +375,7 @@ TEST_P(Laplacian, Accuracy)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -413,7 +413,7 @@ TEST_P(Laplacian, Accuracy)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -474,11 +474,11 @@ PARAM_TEST_CASE(ErodeDilateBase, MatType, bool)
dst = randomMat(rng, size, type, 5, 16, false);
// rng.fill(kernel, cv::RNG::UNIFORM, cv::Scalar::all(0), cv::Scalar::all(3));
kernel = randomMat(rng, Size(3,3), CV_8UC1, 0, 3, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -521,7 +521,7 @@ TEST_P(Erode, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -560,7 +560,7 @@ TEST_P(Erode, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -588,7 +588,7 @@ TEST_P(Dilate, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -625,7 +625,7 @@ TEST_P(Dilate, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -661,7 +661,7 @@ PARAM_TEST_CASE(Sobel, MatType, int, int, int, int)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -683,11 +683,11 @@ PARAM_TEST_CASE(Sobel, MatType, int, int, int, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -726,7 +726,7 @@ TEST_P(Sobel, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -764,7 +764,7 @@ TEST_P(Sobel, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -800,7 +800,7 @@ PARAM_TEST_CASE(Scharr, MatType, int, int, int)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -821,11 +821,11 @@ PARAM_TEST_CASE(Scharr, MatType, int, int, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -863,7 +863,7 @@ TEST_P(Scharr, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -901,7 +901,7 @@ TEST_P(Scharr, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -941,7 +941,7 @@ PARAM_TEST_CASE(GaussianBlur, MatType, cv::Size, int)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -963,11 +963,11 @@ PARAM_TEST_CASE(GaussianBlur, MatType, cv::Size, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -1006,7 +1006,7 @@ TEST_P(GaussianBlur, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1045,7 +1045,7 @@ TEST_P(GaussianBlur, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1078,19 +1078,19 @@ INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine(Values(CV_8UC1, CV_8UC4, CV_32FC1
INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine(Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(false)));
INSTANTIATE_TEST_CASE_P(Filter, Sobel, Combine(Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4),
Values(1, 2), Values(0, 1), Values(3, 5, 7), Values((MatType)cv::BORDER_CONSTANT,
(MatType)cv::BORDER_REPLICATE, (MatType)cv::BORDER_REFLECT, (MatType)cv::BORDER_REFLECT_101)));
INSTANTIATE_TEST_CASE_P(Filter, Sobel, Combine(Values(CV_8UC1, CV_32FC1),
Values(1, 2), Values(0, 1), Values(3, 5), Values((MatType)cv::BORDER_CONSTANT,
(MatType)cv::BORDER_REPLICATE)));
INSTANTIATE_TEST_CASE_P(Filter, Scharr, Combine(
Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(0, 1), Values(0, 1),
Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE, (MatType)cv::BORDER_REFLECT, (MatType)cv::BORDER_REFLECT_101)));
Values(CV_8UC1, CV_32FC1), Values(0, 1), Values(0, 1),
Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE)));
INSTANTIATE_TEST_CASE_P(Filter, GaussianBlur, Combine(
Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4),
Values(cv::Size(3, 3), cv::Size(5, 5), cv::Size(7, 7)),
Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE, (MatType)cv::BORDER_REFLECT, (MatType)cv::BORDER_REFLECT_101)));
Values(CV_8UC1, CV_32FC1),
Values(cv::Size(3, 3), cv::Size(5, 5)),
Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE)));
#endif // HAVE_OPENCL

12
modules/ocl/perf/test_haar.cpp
Unescape View File

@ -57,7 +57,7 @@ struct getRect { Rect operator ()(const CvAvgComp& e) const { return e.rect; } }
PARAM_TEST_CASE(HaarTestBase, int, int)
{
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
cv::ocl::OclCascadeClassifier cascade, nestedCascade;
cv::CascadeClassifier cpucascade, cpunestedCascade;
// Mat img;
@ -85,11 +85,11 @@ PARAM_TEST_CASE(HaarTestBase, int, int)
return;
}
int devnums = getDevice(oclinfo);
CV_Assert(devnums>0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath("E:\\");
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums>0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath("E:\\");
}
};

134
modules/ocl/perf/test_imgproc.cpp
Unescape View File

@ -103,7 +103,7 @@ PARAM_TEST_CASE(ImgprocTestBase, MatType,MatType,MatType,MatType,MatType, bool)
cv::Mat mask_roi;
cv::Mat dst_roi;
cv::Mat dst1_roi; //bak
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl mat
cv::ocl::oclMat clmat1;
cv::ocl::oclMat clmat2;
@ -128,11 +128,11 @@ PARAM_TEST_CASE(ImgprocTestBase, MatType,MatType,MatType,MatType,MatType, bool)
cv::RNG& rng = TS::ptr()->get_rng();
cv::Size size(MWIDTH, MHEIGHT);
double min = 1,max = 20;
int devnums = getDevice(oclinfo);
CV_Assert(devnums>0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums>0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath(CLBINPATH);
if(type1!=nulltype)
{
mat1 = randomMat(rng, size, type1, min, max, false);
@ -289,7 +289,7 @@ TEST_P(equalizeHist, MatType)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -328,7 +328,7 @@ TEST_P(equalizeHist, MatType)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
if(type1!=nulltype)
@ -370,7 +370,7 @@ TEST_P(bilateralFilter, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -409,7 +409,7 @@ TEST_P(bilateralFilter, Mat)
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
if(type1!=nulltype)
@ -450,7 +450,7 @@ TEST_P(CopyMakeBorder, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -488,7 +488,7 @@ TEST_P(CopyMakeBorder, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
if(type1!=nulltype)
@ -516,7 +516,7 @@ TEST_P(cornerMinEigenVal, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -555,7 +555,7 @@ TEST_P(cornerMinEigenVal, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
int blockSize = 7, apertureSize= 1 + 2 * (rand() % 4);
@ -584,7 +584,7 @@ TEST_P(cornerHarris, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -624,7 +624,7 @@ TEST_P(cornerHarris, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
double kk = 2;
@ -655,7 +655,7 @@ TEST_P(integral, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -694,7 +694,7 @@ TEST_P(integral, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
if(type1!=nulltype)
@ -735,7 +735,7 @@ PARAM_TEST_CASE(WarpTestBase, MatType, int)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -755,11 +755,11 @@ PARAM_TEST_CASE(WarpTestBase, MatType, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath(CLBINPATH);
}
void Has_roi(int b)
{
@ -816,7 +816,7 @@ TEST_P(WarpAffine, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -854,7 +854,7 @@ TEST_P(WarpAffine, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -889,7 +889,7 @@ TEST_P(WarpPerspective, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -927,7 +927,7 @@ TEST_P(WarpPerspective, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -969,7 +969,7 @@ PARAM_TEST_CASE(Resize, MatType, cv::Size, double, double, int)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -1003,11 +1003,11 @@ PARAM_TEST_CASE(Resize, MatType, cv::Size, double, double, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, dsize, type, 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath(CLBINPATH);
}
void Has_roi(int b)
{
@ -1053,7 +1053,7 @@ TEST_P(Resize, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1091,7 +1091,7 @@ TEST_P(Resize, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -1127,7 +1127,7 @@ PARAM_TEST_CASE(Threshold, MatType, ThreshOp)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -1146,11 +1146,11 @@ PARAM_TEST_CASE(Threshold, MatType, ThreshOp)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath(CLBINPATH);
}
void Has_roi(int b)
{
@ -1191,7 +1191,7 @@ TEST_P(Threshold, Mat)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -1232,7 +1232,7 @@ TEST_P(Threshold, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
double maxVal = randomDouble(20.0, 127.0);
@ -1277,7 +1277,7 @@ PARAM_TEST_CASE(meanShiftTestBase, MatType, MatType, int, int, cv::TermCriteria)
cv::ocl::oclMat gdst;
cv::ocl::oclMat gdstCoor;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl mat with roi
cv::ocl::oclMat gsrc_roi;
cv::ocl::oclMat gdst_roi;
@ -1300,11 +1300,11 @@ PARAM_TEST_CASE(meanShiftTestBase, MatType, MatType, int, int, cv::TermCriteria)
dst = randomMat(rng, size, type, 5, 16, false);
dstCoor = randomMat(rng, size, typeCoor, 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -1380,7 +1380,7 @@ TEST_P(meanShiftFiltering, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
@ -1438,7 +1438,7 @@ TEST_P(meanShiftProc, Mat)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
@ -1482,21 +1482,21 @@ INSTANTIATE_TEST_CASE_P(ImgprocTestBase, equalizeHist, Combine(
// NULL_TYPE,
// Values(false))); // Values(false) is the reserved parameter
//INSTANTIATE_TEST_CASE_P(ImgprocTestBase, cornerMinEigenVal, Combine(
// Values(CV_8UC1,CV_32FC1),
// NULL_TYPE,
// ONE_TYPE(CV_32FC1),
// NULL_TYPE,
// NULL_TYPE,
// Values(false))); // Values(false) is the reserved parameter
//
//INSTANTIATE_TEST_CASE_P(ImgprocTestBase, cornerHarris, Combine(
// Values(CV_8UC1,CV_32FC1),
// NULL_TYPE,
// ONE_TYPE(CV_32FC1),
// NULL_TYPE,
// NULL_TYPE,
// Values(false))); // Values(false) is the reserved parameter
INSTANTIATE_TEST_CASE_P(ImgprocTestBase, cornerMinEigenVal, Combine(
Values(CV_8UC1,CV_32FC1),
NULL_TYPE,
ONE_TYPE(CV_32FC1),
NULL_TYPE,
NULL_TYPE,
Values(false))); // Values(false) is the reserved parameter
INSTANTIATE_TEST_CASE_P(ImgprocTestBase, cornerHarris, Combine(
Values(CV_8UC1,CV_32FC1),
NULL_TYPE,
ONE_TYPE(CV_32FC1),
NULL_TYPE,
NULL_TYPE,
Values(false))); // Values(false) is the reserved parameter
INSTANTIATE_TEST_CASE_P(ImgprocTestBase, integral, Combine(

56
modules/ocl/perf/test_matrix_operation.cpp
Unescape View File

@ -72,7 +72,7 @@ PARAM_TEST_CASE(ConvertToTestBase, MatType, MatType)
//src mat with roi
cv::Mat mat_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -90,11 +90,11 @@ PARAM_TEST_CASE(ConvertToTestBase, MatType, MatType)
mat = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -141,7 +141,7 @@ TEST_P(ConvertTo, Accuracy)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -177,7 +177,7 @@ TEST_P(ConvertTo, Accuracy)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -216,7 +216,7 @@ PARAM_TEST_CASE(CopyToTestBase, MatType, bool)
cv::Mat mat_roi;
cv::Mat mask_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -237,11 +237,11 @@ PARAM_TEST_CASE(CopyToTestBase, MatType, bool)
mask = randomMat(rng, size, CV_8UC1, 0, 2, false);
cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -293,7 +293,7 @@ TEST_P(CopyTo, Without_mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -329,7 +329,7 @@ TEST_P(CopyTo, Without_mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -352,7 +352,7 @@ TEST_P(CopyTo, With_mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -389,7 +389,7 @@ TEST_P(CopyTo, With_mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gdst_whole = dst;
@ -425,7 +425,7 @@ PARAM_TEST_CASE(SetToTestBase, MatType, bool)
//src mat with roi
cv::Mat mat_roi;
cv::Mat mask_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gmat_whole;
@ -445,11 +445,11 @@ PARAM_TEST_CASE(SetToTestBase, MatType, bool)
cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
val = cv::Scalar(rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0));
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -495,7 +495,7 @@ TEST_P(SetTo, Without_mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -529,7 +529,7 @@ TEST_P(SetTo, Without_mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gmat_whole = mat;
@ -550,7 +550,7 @@ TEST_P(SetTo, With_mask)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -586,7 +586,7 @@ TEST_P(SetTo, With_mask)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gmat_whole = mat;

32
modules/ocl/perf/test_split_merge.cpp
Unescape View File

@ -87,7 +87,7 @@ PARAM_TEST_CASE(MergeTestBase, MatType, int)
//dst mat with roi
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -111,11 +111,11 @@ PARAM_TEST_CASE(MergeTestBase, MatType, int)
mat3 = randomMat(rng, size, CV_MAKETYPE(type, 1), 5, 16, false);
mat4 = randomMat(rng, size, CV_MAKETYPE(type, 1), 5, 16, false);
dst = randomMat(rng, size, CV_MAKETYPE(type, channels), 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//setBinpath(CLBINPATH);
}
void Has_roi(int b)
{
@ -174,7 +174,7 @@ TEST_P(Merge, Accuracy)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -223,7 +223,7 @@ TEST_P(Merge, Accuracy)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
gmat1 = mat1_roi;
@ -281,7 +281,7 @@ PARAM_TEST_CASE(SplitTestBase, MatType, int)
cv::Mat dst2_roi;
cv::Mat dst3_roi;
cv::Mat dst4_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst1_whole;
cv::ocl::oclMat gdst2_whole;
@ -308,11 +308,11 @@ PARAM_TEST_CASE(SplitTestBase, MatType, int)
dst2 = randomMat(rng, size, CV_MAKETYPE(type, 1), 5, 16, false);
dst3 = randomMat(rng, size, CV_MAKETYPE(type, 1), 5, 16, false);
dst4 = randomMat(rng, size, CV_MAKETYPE(type, 1), 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
setBinpath(CLBINPATH);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//setBinpath(CLBINPATH);
}
void Has_roi(int b)
@ -370,7 +370,7 @@ TEST_P(Split, Accuracy)
double t0=0;
double t1=0;
double t2=0;
for(int k=0;k<2;k++){
for(int k=LOOPROISTART;k<LOOPROIEND;k++){
totalcputick=0;
totalgputick=0;
totalgputick_kernel=0;
@ -422,7 +422,7 @@ TEST_P(Split, Accuracy)
cout << "average gpu runtime without data transfer is " << totalgputick_kernel/((double)cvGetTickFrequency()* LOOP_TIMES *1000.) << "ms" << endl;
}
#else
for(int j = 0; j < 2; j ++)
for(int j = LOOPROISTART; j < LOOPROIEND; j ++)
{
Has_roi(j);
cv::Mat dev_dst[4] = {dst1_roi, dst2_roi, dst3_roi, dst4_roi};

6
modules/ocl/perf/utility.hpp
Unescape View File

@ -47,9 +47,11 @@
#else
#define LOOP_TIMES 1
#endif
#define MWIDTH 2557
#define MHEIGHT 2579
#define MWIDTH 256
#define MHEIGHT 256
#define CLBINPATH ".\\"
#define LOOPROISTART 0
#define LOOPROIEND 1
int randomInt(int minVal, int maxVal);
double randomDouble(double minVal, double maxVal);

34
modules/ocl/src/arithm.cpp
Unescape View File

@ -307,7 +307,7 @@ void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string
{
if(src1.clCxt -> impl -> double_support ==0 && src1.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -370,7 +370,7 @@ void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst, const o
{
if(src1.clCxt -> impl -> double_support ==0 && src1.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -467,7 +467,7 @@ void arithmetic_scalar_run(const oclMat &src1, const Scalar &src2, oclMat &dst,
{
if(src1.clCxt -> impl -> double_support ==0 && src1.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -537,7 +537,7 @@ void arithmetic_scalar_run(const oclMat &src, oclMat &dst, string kernelName, co
{
if(src.clCxt -> impl -> double_support ==0 && src.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -640,7 +640,7 @@ void cv::ocl::divide(double scalar, const oclMat &src, oclMat &dst)
{
if(src.clCxt -> impl -> double_support ==0)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -803,7 +803,7 @@ Scalar cv::ocl::sum(const oclMat &src)
{
if(src.clCxt->impl->double_support==0 && src.depth()==CV_64F)
{
CV_Error(-217,"select device don't support double");
CV_Error(CV_GpuNotSupported,"select device don't support double");
}
static sumFunc functab[2] =
{
@ -948,7 +948,7 @@ void cv::ocl::minMax(const oclMat &src, double *minVal, double *maxVal, const oc
CV_Assert(src.channels() == 1);
if(src.clCxt->impl->double_support==0 && src.depth()==CV_64F)
{
CV_Error(-217,"select device don't support double");
CV_Error(CV_GpuNotSupported,"select device don't support double");
}
static minMaxFunc functab[8] =
{
@ -1032,7 +1032,7 @@ void arithmetic_flip_rows_run(const oclMat &src, oclMat &dst, string kernelName)
{
if(src.clCxt -> impl -> double_support ==0 && src.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -1081,7 +1081,7 @@ void arithmetic_flip_cols_run(const oclMat &src, oclMat &dst, string kernelName,
{
if(src.clCxt -> impl -> double_support ==0 && src.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -1261,7 +1261,7 @@ void arithmetic_exp_log_run(const oclMat &src, oclMat &dst, string kernelName, c
Context *clCxt = src.clCxt;
if(clCxt -> impl -> double_support ==0 && src.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
//int channels = dst.channels();
@ -1302,7 +1302,7 @@ void arithmetic_magnitude_phase_run(const oclMat &src1, const oclMat &src2, oclM
{
if(src1.clCxt -> impl -> double_support ==0 && src1.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -1350,7 +1350,7 @@ void arithmetic_phase_run(const oclMat &src1, const oclMat &src2, oclMat &dst, s
{
if(src1.clCxt -> impl -> double_support ==0 && src1.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -1414,7 +1414,7 @@ void arithmetic_cartToPolar_run(const oclMat &src1, const oclMat &src2, oclMat &
{
if(src1.clCxt -> impl -> double_support ==0 && src1.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -1469,7 +1469,7 @@ void arithmetic_ptc_run(const oclMat &src1, const oclMat &src2, oclMat &dst1, oc
{
if(src1.clCxt -> impl -> double_support ==0 && src1.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -1651,7 +1651,7 @@ void cv::ocl::minMaxLoc(const oclMat &src, double *minVal, double *maxVal,
{
if(src.clCxt->impl->double_support==0 && src.depth()==CV_64F)
{
CV_Error(-217,"select device don't support double");
CV_Error(CV_GpuNotSupported,"select device don't support double");
}
static minMaxLocFunc functab[2] =
{
@ -1698,7 +1698,7 @@ int cv::ocl::countNonZero(const oclMat &src)
size_t groupnum = src.clCxt->impl->maxComputeUnits;
if(src.clCxt->impl->double_support == 0 && src.depth()==CV_64F)
{
CV_Error(-217,"select device don't support double");
CV_Error(CV_GpuNotSupported,"select device don't support double");
}
CV_Assert(groupnum != 0);
groupnum = groupnum * 2;
@ -2122,7 +2122,7 @@ void transpose_run(const oclMat &src, oclMat &dst, string kernelName)
{
if(src.clCxt -> impl -> double_support ==0 && src.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}

6
modules/ocl/src/filtering.cpp
Unescape View File

@ -371,7 +371,7 @@ void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel, Size &ksize, c
sprintf(s, "-D VAL=FLT_MAX -D GENTYPE=float4");
break;
default:
CV_Error(-217,"unsupported type");
CV_Error(CV_StsUnsupportedFormat,"unsupported type");
}
char compile_option[128];
sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D ERODE %s", anchor.x, anchor.y, localThreads[0], localThreads[1],s);
@ -443,7 +443,7 @@ void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel, Size &ksize,
sprintf(s, "-D VAL=-FLT_MAX -D GENTYPE=float4");
break;
default:
CV_Error(-217,"unsupported type");
CV_Error(CV_StsUnsupportedFormat,"unsupported type");
}
char compile_option[128];
sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D DILATE %s", anchor.x, anchor.y, localThreads[0], localThreads[1],s);
@ -1501,7 +1501,7 @@ void cv::ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize, d
{
if(src.clCxt -> impl -> double_support ==0 && src.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}

143
modules/ocl/src/haar.cpp
Unescape View File

@ -976,8 +976,8 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS
cl_mem nodebuffer;
cl_mem candidatebuffer;
cl_mem scaleinfobuffer;
cl_kernel kernel;
kernel = openCLGetKernelFromSource(gimg.clCxt, &haarobjectdetect, "gpuRunHaarClassifierCascade");
//cl_kernel kernel;
//kernel = openCLGetKernelFromSource(gimg.clCxt, &haarobjectdetect, "gpuRunHaarClassifierCascade");
cv::Rect roi, roi2;
cv::Mat imgroi, imgroisq;
cv::ocl::oclMat resizeroi, gimgroi, gimgroisq;
@ -1060,23 +1060,22 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS
//openCLVerifyCall(status);
//openCLSafeCall(clEnqueueWriteBuffer(gsum.clCxt->clCmdQueue,cascadebuffer,1,0,sizeof(GpuHidHaarClassifierCascade),gcascade,0,NULL,NULL));
stagebuffer = clCreateBuffer(gsum.clCxt->impl->clContext, CL_MEM_READ_ONLY, sizeof(GpuHidHaarStageClassifier) * gcascade->count, NULL, &status);
openCLVerifyCall(status);
stagebuffer = openCLCreateBuffer(gsum.clCxt, CL_MEM_READ_ONLY, sizeof(GpuHidHaarStageClassifier) * gcascade->count);
//openCLVerifyCall(status);
openCLSafeCall(clEnqueueWriteBuffer(gsum.clCxt->impl->clCmdQueue, stagebuffer, 1, 0, sizeof(GpuHidHaarStageClassifier)*gcascade->count, stage, 0, NULL, NULL));
//classifierbuffer = clCreateBuffer(gsum.clCxt->clContext,CL_MEM_READ_ONLY,sizeof(GpuHidHaarClassifier)*totalclassifier,NULL,&status);
//status = clEnqueueWriteBuffer(gsum.clCxt->clCmdQueue,classifierbuffer,1,0,sizeof(GpuHidHaarClassifier)*totalclassifier,classifier,0,NULL,NULL);
nodebuffer = clCreateBuffer(gsum.clCxt->impl->clContext, CL_MEM_READ_ONLY,
nodenum * sizeof(GpuHidHaarTreeNode), NULL, &status);
openCLVerifyCall(status);
nodebuffer = openCLCreateBuffer(gsum.clCxt, CL_MEM_READ_ONLY,nodenum * sizeof(GpuHidHaarTreeNode));
//openCLVerifyCall(status);
openCLSafeCall(clEnqueueWriteBuffer(gsum.clCxt->impl->clCmdQueue, nodebuffer, 1, 0,
nodenum * sizeof(GpuHidHaarTreeNode),
node, 0, NULL, NULL));
candidatebuffer = clCreateBuffer(gsum.clCxt->impl->clContext, CL_MEM_WRITE_ONLY, 4 * sizeof(int) * outputsz, NULL, &status);
openCLVerifyCall(status);
scaleinfobuffer = clCreateBuffer(gsum.clCxt->impl->clContext, CL_MEM_READ_ONLY, sizeof(detect_piramid_info) * loopcount, NULL, &status);
openCLVerifyCall(status);
candidatebuffer = openCLCreateBuffer(gsum.clCxt, CL_MEM_WRITE_ONLY, 4 * sizeof(int) * outputsz);
//openCLVerifyCall(status);
scaleinfobuffer = openCLCreateBuffer(gsum.clCxt, CL_MEM_READ_ONLY, sizeof(detect_piramid_info) * loopcount);
//openCLVerifyCall(status);
openCLSafeCall(clEnqueueWriteBuffer(gsum.clCxt->impl->clCmdQueue, scaleinfobuffer, 1, 0, sizeof(detect_piramid_info)*loopcount, scaleinfo, 0, NULL, NULL));
//flag = 1;
//}
@ -1105,8 +1104,27 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS
int argcount = 0;
//int grpnumperline = ((m + localThreads[0] - 1) / localThreads[0]);
//int totalgrp = ((n + localThreads[1] - 1) / localThreads[1])*grpnumperline;
openCLVerifyKernel(gsum.clCxt, kernel, &blocksize, globalThreads, localThreads);
// openCLVerifyKernel(gsum.clCxt, kernel, &blocksize, globalThreads, localThreads);
//openCLSafeCall(clSetKernelArg(kernel,argcount++,sizeof(cl_mem),(void*)&cascadebuffer));
vector<pair<size_t,const void *> > args;
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&stagebuffer ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&scaleinfobuffer ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&nodebuffer ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&gsum.data ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&gsqsum.data ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&candidatebuffer ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&pixelstep ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&loopcount ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&startstage ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&splitstage ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&endstage ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&startnode ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&splitnode ));
args.push_back ( make_pair(sizeof(cl_int4) , (void *)&p ));
args.push_back ( make_pair(sizeof(cl_int4) , (void *)&pq ));
args.push_back ( make_pair(sizeof(cl_float) , (void *)&correction ));
/*
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_mem), (void *)&stagebuffer));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_mem), (void *)&scaleinfobuffer));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_mem), (void *)&nodebuffer));
@ -1122,19 +1140,20 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_int), (void *)&splitnode));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_int4), (void *)&p));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_int4), (void *)&pq));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_float), (void *)&correction));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_float), (void *)&correction));*/
//openCLSafeCall(clSetKernelArg(kernel,argcount++,sizeof(cl_int),(void*)&n));
//openCLSafeCall(clSetKernelArg(kernel,argcount++,sizeof(cl_int),(void*)&grpnumperline));
//openCLSafeCall(clSetKernelArg(kernel,argcount++,sizeof(cl_int),(void*)&totalgrp));
openCLSafeCall(clEnqueueNDRangeKernel(gsum.clCxt->impl->clCmdQueue, kernel, 2, NULL, globalThreads, localThreads, 0, NULL, NULL));
// openCLSafeCall(clEnqueueNDRangeKernel(gsum.clCxt->impl->clCmdQueue, kernel, 2, NULL, globalThreads, localThreads, 0, NULL, NULL));
openCLSafeCall(clFinish(gsum.clCxt->impl->clCmdQueue));
//t = (double)cvGetTickCount() - t;
// openCLSafeCall(clFinish(gsum.clCxt->impl->clCmdQueue));
openCLExecuteKernel(gsum.clCxt, &haarobjectdetect, "gpuRunHaarClassifierCascade", globalThreads, localThreads, args, -1, -1);
//t = (double)cvGetTickCount() - t;
//printf( "detection time = %g ms\n", t/((double)cvGetTickFrequency()*1000.) );
//t = (double)cvGetTickCount();
openCLSafeCall(clEnqueueReadBuffer(gsum.clCxt->impl->clCmdQueue, candidatebuffer, 1, 0, 4 * sizeof(int)*outputsz, candidate, 0, NULL, NULL));
//openCLSafeCall(clEnqueueReadBuffer(gsum.clCxt->impl->clCmdQueue, candidatebuffer, 1, 0, 4 * sizeof(int)*outputsz, candidate, 0, NULL, NULL));
openCLReadBuffer( gsum.clCxt, candidatebuffer, candidate, 4 * sizeof(int)*outputsz );
for(int i = 0; i < outputsz; i++)
if(candidate[4*i+2] != 0)
@ -1149,7 +1168,7 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS
openCLSafeCall(clReleaseMemObject(scaleinfobuffer));
openCLSafeCall(clReleaseMemObject(nodebuffer));
openCLSafeCall(clReleaseMemObject(candidatebuffer));
openCLSafeCall(clReleaseKernel(kernel));
// openCLSafeCall(clReleaseKernel(kernel));
//t = (double)cvGetTickCount() - t;
//printf( "release time = %g ms\n", t/((double)cvGetTickFrequency()*1000.) );
}
@ -1212,19 +1231,19 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS
int outputsz = 256 * globalThreads[0] / localThreads[0];
int nodenum = (datasize - sizeof(GpuHidHaarClassifierCascade) -
sizeof(GpuHidHaarStageClassifier) * gcascade->count - sizeof(GpuHidHaarClassifier) * totalclassifier) / sizeof(GpuHidHaarTreeNode);
nodebuffer = clCreateBuffer(gsum.clCxt->impl->clContext, CL_MEM_READ_ONLY,
nodenum * sizeof(GpuHidHaarTreeNode), NULL, &status);
openCLVerifyCall(status);
nodebuffer = openCLCreateBuffer(gsum.clCxt, CL_MEM_READ_ONLY,
nodenum * sizeof(GpuHidHaarTreeNode));
//openCLVerifyCall(status);
openCLSafeCall(clEnqueueWriteBuffer(gsum.clCxt->impl->clCmdQueue, nodebuffer, 1, 0,
nodenum * sizeof(GpuHidHaarTreeNode),
node, 0, NULL, NULL));
cl_mem newnodebuffer = clCreateBuffer(gsum.clCxt->impl->clContext, CL_MEM_READ_WRITE,
loopcount * nodenum * sizeof(GpuHidHaarTreeNode), NULL, &status);
cl_mem newnodebuffer = openCLCreateBuffer(gsum.clCxt, CL_MEM_READ_WRITE,
loopcount * nodenum * sizeof(GpuHidHaarTreeNode));
int startstage = 0;
int endstage = gcascade->count;
cl_kernel kernel;
kernel = openCLGetKernelFromSource(gsum.clCxt, &haarobjectdetect_scaled2, "gpuRunHaarClassifierCascade_scaled2");
cl_kernel kernel2 = openCLGetKernelFromSource(gimg.clCxt, &haarobjectdetect_scaled2, "gpuscaleclassifier");
//cl_kernel kernel;
//kernel = openCLGetKernelFromSource(gsum.clCxt, &haarobjectdetect_scaled2, "gpuRunHaarClassifierCascade_scaled2");
//cl_kernel kernel2 = openCLGetKernelFromSource(gimg.clCxt, &haarobjectdetect_scaled2, "gpuscaleclassifier");
for(int i = 0; i < loopcount; i++)
{
sz = sizev[i];
@ -1251,34 +1270,48 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS
int startnodenum = nodenum * i;
int argcounts = 0;
float factor2 = (float)factor;
/*
openCLSafeCall(clSetKernelArg(kernel2, argcounts++, sizeof(cl_mem), (void *)&nodebuffer));
openCLSafeCall(clSetKernelArg(kernel2, argcounts++, sizeof(cl_mem), (void *)&newnodebuffer));
openCLSafeCall(clSetKernelArg(kernel2, argcounts++, sizeof(cl_float), (void *)&factor2));
openCLSafeCall(clSetKernelArg(kernel2, argcounts++, sizeof(cl_float), (void *)&correction[i]));
openCLSafeCall(clSetKernelArg(kernel2, argcounts++, sizeof(cl_int), (void *)&startnodenum));
size_t globalThreads2[1] = {nodenum};
clEnqueueNDRangeKernel(gsum.clCxt->impl->clCmdQueue, kernel2, 1, NULL, globalThreads2, 0, 0, NULL, NULL);
clFinish(gsum.clCxt->impl->clCmdQueue);
*/
vector<pair<size_t,const void *> > args1;
args1.push_back ( make_pair(sizeof(cl_mem) , (void *)&nodebuffer ));
args1.push_back ( make_pair(sizeof(cl_mem) , (void *)&newnodebuffer ));
args1.push_back ( make_pair(sizeof(cl_float) , (void *)&factor2 ));
args1.push_back ( make_pair(sizeof(cl_float) , (void *)&correction[i] ));
args1.push_back ( make_pair(sizeof(cl_int) , (void *)&startnodenum ));
size_t globalThreads2[3] = {nodenum,1,1};
size_t localThreads2[3] = {256,1,1};
openCLExecuteKernel(gsum.clCxt, &haarobjectdetect_scaled2, "gpuscaleclassifier", globalThreads2, NULL/*localThreads2*/, args1, -1, -1);
//clEnqueueNDRangeKernel(gsum.clCxt->impl->clCmdQueue, kernel2, 1, NULL, globalThreads2, 0, 0, NULL, NULL);
//clFinish(gsum.clCxt->impl->clCmdQueue);
}
clReleaseKernel(kernel2);
//clReleaseKernel(kernel2);
int step = gsum.step / 4;
int startnode = 0;
int splitstage = 3;
int splitnode = stage[0].count + stage[1].count + stage[2].count;
stagebuffer = clCreateBuffer(gsum.clCxt->impl->clContext, CL_MEM_READ_ONLY, sizeof(GpuHidHaarStageClassifier) * gcascade->count, NULL, &status);
openCLVerifyCall(status);
stagebuffer = openCLCreateBuffer(gsum.clCxt, CL_MEM_READ_ONLY, sizeof(GpuHidHaarStageClassifier) * gcascade->count);
//openCLVerifyCall(status);
openCLSafeCall(clEnqueueWriteBuffer(gsum.clCxt->impl->clCmdQueue, stagebuffer, 1, 0, sizeof(GpuHidHaarStageClassifier)*gcascade->count, stage, 0, NULL, NULL));
candidatebuffer = clCreateBuffer(gsum.clCxt->impl->clContext, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, 4 * sizeof(int) * outputsz, NULL, &status);
openCLVerifyCall(status);
scaleinfobuffer = clCreateBuffer(gsum.clCxt->impl->clContext, CL_MEM_READ_ONLY, sizeof(detect_piramid_info) * loopcount, NULL, &status);
openCLVerifyCall(status);
candidatebuffer = openCLCreateBuffer(gsum.clCxt, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, 4 * sizeof(int) * outputsz);
//openCLVerifyCall(status);
scaleinfobuffer = openCLCreateBuffer(gsum.clCxt, CL_MEM_READ_ONLY, sizeof(detect_piramid_info) * loopcount);
//openCLVerifyCall(status);
openCLSafeCall(clEnqueueWriteBuffer(gsum.clCxt->impl->clCmdQueue, scaleinfobuffer, 1, 0, sizeof(detect_piramid_info)*loopcount, scaleinfo, 0, NULL, NULL));
pbuffer = clCreateBuffer(gsum.clCxt->impl->clContext, CL_MEM_READ_ONLY, sizeof(cl_int4) * loopcount, NULL, &status);
pbuffer = openCLCreateBuffer(gsum.clCxt, CL_MEM_READ_ONLY, sizeof(cl_int4) * loopcount);
openCLSafeCall(clEnqueueWriteBuffer(gsum.clCxt->impl->clCmdQueue, pbuffer, 1, 0, sizeof(cl_int4)*loopcount, p, 0, NULL, NULL));
correctionbuffer = clCreateBuffer(gsum.clCxt->impl->clContext, CL_MEM_READ_ONLY, sizeof(cl_float) * loopcount, NULL, &status);
correctionbuffer = openCLCreateBuffer(gsum.clCxt, CL_MEM_READ_ONLY, sizeof(cl_float) * loopcount);
openCLSafeCall(clEnqueueWriteBuffer(gsum.clCxt->impl->clCmdQueue, correctionbuffer, 1, 0, sizeof(cl_float)*loopcount, correction, 0, NULL, NULL));
int argcount = 0;
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_mem), (void *)&stagebuffer));
//int argcount = 0;
/*openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_mem), (void *)&stagebuffer));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_mem), (void *)&scaleinfobuffer));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_mem), (void *)&newnodebuffer));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_mem), (void *)&gsum.data));
@ -1293,10 +1326,30 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_int), (void *)&splitnode));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_mem), (void *)&pbuffer));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_mem), (void *)&correctionbuffer));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_int), (void *)&nodenum));
openCLSafeCall(clEnqueueNDRangeKernel(gsum.clCxt->impl->clCmdQueue, kernel, 2, NULL, globalThreads, localThreads, 0, NULL, NULL));
openCLSafeCall(clFinish(gsum.clCxt->impl->clCmdQueue));
openCLSafeCall(clSetKernelArg(kernel, argcount++, sizeof(cl_int), (void *)&nodenum));*/
vector<pair<size_t,const void *> > args;
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&stagebuffer ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&scaleinfobuffer ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&newnodebuffer ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&gsum.data ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&gsqsum.data ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&candidatebuffer ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&step ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&loopcount ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&startstage ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&splitstage ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&endstage ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&startnode ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&splitnode ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&pbuffer ));
args.push_back ( make_pair(sizeof(cl_mem) , (void *)&correctionbuffer ));
args.push_back ( make_pair(sizeof(cl_int) , (void *)&nodenum ));
openCLExecuteKernel(gsum.clCxt, &haarobjectdetect_scaled2, "gpuRunHaarClassifierCascade_scaled2", globalThreads, localThreads, args, -1, -1);
//openCLSafeCall(clEnqueueNDRangeKernel(gsum.clCxt->impl->clCmdQueue, kernel, 2, NULL, globalThreads, localThreads, 0, NULL, NULL));
//openCLSafeCall(clFinish(gsum.clCxt->impl->clCmdQueue));
//openCLSafeCall(clEnqueueReadBuffer(gsum.clCxt->clCmdQueue,candidatebuffer,1,0,4*sizeof(int)*outputsz,candidate,0,NULL,NULL));
candidate = (int *)clEnqueueMapBuffer(gsum.clCxt->impl->clCmdQueue, candidatebuffer, 1, CL_MAP_READ, 0, 4 * sizeof(int), 0, 0, 0, &status);

8
modules/ocl/src/imgproc.cpp
Unescape View File

@ -672,7 +672,7 @@ namespace cv
break;
}
default:
CV_Error(-217, "Unsupported source type");
CV_Error(CV_StsUnsupportedFormat, "Unsupported source type");
}
}
@ -898,7 +898,7 @@ namespace cv
CV_Assert(src.type() == CV_8UC1);
if(src.clCxt->impl->double_support == 0 && src.depth() ==CV_64F)
{
CV_Error(-217,"select device don't support double");
CV_Error(CV_GpuNotSupported,"select device don't support double");
}
int vlen = 4;
int offset = src.offset / vlen;
@ -1080,7 +1080,7 @@ namespace cv
{
if(src.clCxt->impl->double_support == 0 && src.depth() ==CV_64F)
{
CV_Error(-217,"select device don't support double");
CV_Error(CV_GpuNotSupported,"select device don't support double");
}
oclMat Dx, Dy;
CV_Assert(borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT);
@ -1093,7 +1093,7 @@ namespace cv
{
if(src.clCxt->impl->double_support == 0 && src.depth() ==CV_64F)
{
CV_Error(-217,"select device don't support double");
CV_Error(CV_GpuNotSupported,"select device don't support double");
}
oclMat Dx, Dy;
CV_Assert(borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT);

26
modules/ocl/src/initialization.cpp
Unescape View File

@ -257,7 +257,7 @@ namespace cv
_devicetype = CL_DEVICE_TYPE_ALL;
break;
default:
CV_Error(-217,"Unkown device type");
CV_Error(CV_GpuApiCallError,"Unkown device type");
}
int devcienums = 0;
// Platform info
@ -456,7 +456,7 @@ namespace cv
char **binaries = (char **)malloc( sizeof(char *) * numDevices );
if(binaries == NULL)
{
CV_Error(-217,"Failed to allocate host memory.(binaries)\r\n");
CV_Error(CV_StsNoMem,"Failed to allocate host memory.(binaries)\r\n");
}
for(i = 0; i < numDevices; i++)
@ -466,7 +466,7 @@ namespace cv
binaries[i] = (char *)malloc( sizeof(char) * binarySizes[i]);
if(binaries[i] == NULL)
{
CV_Error(-217,"Failed to allocate host memory.(binaries[i])\r\n");
CV_Error(CV_StsNoMem,"Failed to allocate host memory.(binaries[i])\r\n");
}
}
else
@ -498,7 +498,7 @@ namespace cv
{
char *temp;
sprintf(temp, "Failed to load kernel file : %s\r\n", fileName);
CV_Error(-217, temp);
CV_Error(CV_GpuApiCallError, temp);
}
else
{
@ -661,14 +661,16 @@ namespace cv
cl_kernel kernel;
kernel = openCLGetKernelFromSource(clCxt, source, kernelName, build_options);
globalThreads[0] = divUp(globalThreads[0], localThreads[0]) * localThreads[0];
globalThreads[1] = divUp(globalThreads[1], localThreads[1]) * localThreads[1];
globalThreads[2] = divUp(globalThreads[2], localThreads[2]) * localThreads[2];
size_t blockSize = localThreads[0] * localThreads[1] * localThreads[2];
cv::ocl::openCLVerifyKernel(clCxt, kernel, &blockSize, globalThreads, localThreads);
if ( localThreads != NULL)
{
globalThreads[0] = divUp(globalThreads[0], localThreads[0]) * localThreads[0];
globalThreads[1] = divUp(globalThreads[1], localThreads[1]) * localThreads[1];
globalThreads[2] = divUp(globalThreads[2], localThreads[2]) * localThreads[2];
size_t blockSize = localThreads[0] * localThreads[1] * localThreads[2];
cv::ocl::openCLVerifyKernel(clCxt, kernel, &blockSize, globalThreads, localThreads);
}
for(int i = 0; i < args.size(); i ++)
openCLSafeCall(clSetKernelArg(kernel, i, args[i].first, args[i].second));

90
modules/ocl/src/kernels/convertC3C4.cl
Unescape View File

@ -33,10 +33,13 @@
//
//
//#pragma OPENCL EXTENSION cl_amd_printf : enable
#define WORKGROUPSIZE 256
__kernel void convertC3C4(__global const GENTYPE4 * restrict src, __global GENTYPE4 *dst, int cols, int rows,
int dstStep_in_piexl,int pixel_end)
{
int id = get_global_id(0);
//read data from source
//int pixel_end = mul24(cols -1 , rows -1);
int3 pixelid = (int3)(mul24(id,3),mad24(id,3,1),mad24(id,3,2));
pixelid = clamp(pixelid,0,pixel_end);
@ -51,18 +54,35 @@ __kernel void convertC3C4(__global const GENTYPE4 * restrict src, __global GENTY
outpix2 = (GENTYPE4)(pixel1.z,pixel1.w,pixel2.x,0);
outpix3 = (GENTYPE4)(pixel2.y,pixel2.z,pixel2.w,0);
int4 outy = (id<<2)/cols;
int4 outx = (id<<2)%cols;
outx.y++;
outx.z+=2;
outx.w+=3;
outy = select(outy,outy+1,outx>=cols);
outx = select(outx,outx-cols,outx>=cols);
//outpix3 = select(outpix3, outpix0, (uchar4)(outy.w>=rows));
//outpix2 = select(outpix2, outpix0, (uchar4)(outy.z>=rows));
//outpix1 = select(outpix1, outpix0, (uchar4)(outy.y>=rows));
//outx = select(outx,(int4)outx.x,outy>=rows);
//outy = select(outy,(int4)outy.x,outy>=rows);
//permutate the data in LDS to avoid global memory conflict
__local GENTYPE4 rearrange[WORKGROUPSIZE*4];
int lid = get_local_id(0)<<2;
rearrange[lid++] = outpix0;
rearrange[lid++] = outpix1;
rearrange[lid++] = outpix2;
rearrange[lid] = outpix3;
lid = get_local_id(0);
barrier(CLK_LOCAL_MEM_FENCE);
outpix0 = rearrange[lid];
lid+=WORKGROUPSIZE;
outpix1 = rearrange[lid];
lid+=WORKGROUPSIZE;
outpix2 = rearrange[lid];
lid+=WORKGROUPSIZE;
outpix3 = rearrange[lid];
//calculate output index
int4 outx, outy;
int4 startid = mad24((int)get_group_id(0),WORKGROUPSIZE*4,(int)get_local_id(0));
startid.y+=WORKGROUPSIZE;
startid.z+=WORKGROUPSIZE*2;
startid.w+=WORKGROUPSIZE*3;
outx = startid%(int4)cols;
outy = startid/(int4)cols;
int4 addr = mad24(outy,dstStep_in_piexl,outx);
if(outx.w<cols && outy.w<rows)
{
@ -102,6 +122,7 @@ __kernel void convertC4C3(__global const GENTYPE4 * restrict src, __global GENTY
x4 = select(x4,x4-(int4)cols,x4>=(int4)cols);
int4 addr = mad24(y4,(int4)srcStep_in_pixel,x4);
GENTYPE4 pixel0,pixel1,pixel2,pixel3, outpixel1, outpixel2;
//read data from src
pixel0 = src[addr.x];
pixel1 = src[addr.y];
pixel2 = src[addr.z];
@ -116,23 +137,40 @@ __kernel void convertC4C3(__global const GENTYPE4 * restrict src, __global GENTY
outpixel2.y = pixel3.x;
outpixel2.z = pixel3.y;
outpixel2.w = pixel3.z;
int4 outaddr = mul24(id>>2 , 3);
outaddr.y++;
outaddr.z+=2;
//printf("%d ",outaddr.z);
if(outaddr.z <= pixel_end)
//permutate the data in LDS to avoid global memory conflict
__local GENTYPE4 rearrange[WORKGROUPSIZE*3];
int lid = mul24((int)get_local_id(0),3);
rearrange[lid++] = pixel0;
rearrange[lid++] = outpixel1;
rearrange[lid] = outpixel2;
barrier(CLK_LOCAL_MEM_FENCE);
lid = get_local_id(0);
pixel0 = rearrange[lid];
lid+=WORKGROUPSIZE;
outpixel1 = rearrange[lid];
lid+=WORKGROUPSIZE;
outpixel2 = rearrange[lid];
//calcultate output index
int3 startid = mad24((int)get_group_id(0),WORKGROUPSIZE*3,(int)get_local_id(0));
startid.y+=WORKGROUPSIZE;
startid.z+=WORKGROUPSIZE*2;
//id = mul24(id>>2 , 3);
if(startid.z <= pixel_end)
{
dst[outaddr.x] = pixel0;
dst[outaddr.y] = outpixel1;
dst[outaddr.z] = outpixel2;
dst[startid.x] = pixel0;
dst[startid.y] = outpixel1;
dst[startid.z] = outpixel2;
}
else if(outaddr.y <= pixel_end)
else if(startid.y <= pixel_end)
{
dst[outaddr.x] = pixel0;
dst[outaddr.y] = outpixel1;
dst[startid.x] = pixel0;
dst[startid.y] = outpixel1;
}
else if(outaddr.x <= pixel_end)
else if(startid.x <= pixel_end)
{
dst[outaddr.x] = pixel0;
}
dst[startid.x] = pixel0;
}
}

20
modules/ocl/src/kernels/imgproc_integral.cl
Unescape View File

@ -74,8 +74,8 @@ kernel void integral_cols(__global uchar4 *src,__global int *sum ,__global float
src_t[0] = (i + lid < rows ? convert_int4(src[src_offset + (lid+i) * src_step + gid]) : 0);
src_t[1] = (i + lid < rows ? convert_int4(src[src_offset + (lid+i) * src_step + gid + 1]) : 0);
sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]);
sqsum_t[0] = (i == 0 ? 0 : lm_sqsum[0][LSIZE_2 + LOG_LSIZE]);
sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]);
sqsum_t[0] = (i == 0 ? 0 : lm_sqsum[0][LSIZE_2 + LOG_LSIZE]);
sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]);
sqsum_t[1] = (i == 0 ? 0 : lm_sqsum[1][LSIZE_2 + LOG_LSIZE]);
barrier(CLK_LOCAL_MEM_FENCE);
@ -102,6 +102,7 @@ kernel void integral_cols(__global uchar4 *src,__global int *sum ,__global float
}
offset <<= 1;
}
barrier(CLK_LOCAL_MEM_FENCE);
if(lid < 2)
{
lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
@ -124,8 +125,9 @@ kernel void integral_cols(__global uchar4 *src,__global int *sum ,__global float
lm_sqsum[lid >> 7][ai] = lm_sqsum[lid >> 7][bi] - lm_sqsum[lid >> 7][ai];
}
}
if(lid > 0 & (i+lid) <= rows){
int loc_s0 = gid * dst_step + i + lid - 1 - pre_invalid * dst_step / 4, loc_s1 = loc_s0 + dst_step ;
barrier(CLK_LOCAL_MEM_FENCE);
int loc_s0 = gid * dst_step + i + lid - 1 - pre_invalid * dst_step / 4, loc_s1 = loc_s0 + dst_step ;
if(lid > 0 && (i+lid) <= rows){
lm_sum[0][bf_loc] += sum_t[0];
lm_sum[1][bf_loc] += sum_t[1];
lm_sqsum[0][bf_loc] += sqsum_t[0];
@ -200,6 +202,7 @@ kernel void integral_rows(__global int4 *srcsum,__global float4 * srcsqsum,__glo
}
offset <<= 1;
}
barrier(CLK_LOCAL_MEM_FENCE);
if(lid < 2)
{
lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
@ -222,7 +225,7 @@ kernel void integral_rows(__global int4 *srcsum,__global float4 * srcsqsum,__glo
lm_sqsum[lid >> 7][ai] = lm_sqsum[lid >> 7][bi] - lm_sqsum[lid >> 7][ai];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
if(gid == 0 && (i + lid) <= rows)
{
sum[sum_offset + i + lid] = 0;
@ -239,10 +242,9 @@ kernel void integral_rows(__global int4 *srcsum,__global float4 * srcsqsum,__glo
sqsum[sqsum_offset + loc1 + k * sqsum_step / 4] = 0;
}
}
if(lid > 0 & (i+lid) <= rows){
int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ;
int loc_sq0 = sqsum_offset + gid * 2 * sqsum_step + sqsum_step / 4 + i + lid, loc_sq1 = loc_sq0 + sqsum_step ;
int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ;
int loc_sq0 = sqsum_offset + gid * 2 * sqsum_step + sqsum_step / 4 + i + lid, loc_sq1 = loc_sq0 + sqsum_step ;
if(lid > 0 && (i+lid) <= rows){
lm_sum[0][bf_loc] += sum_t[0];
lm_sum[1][bf_loc] += sum_t[1];
lm_sqsum[0][bf_loc] += sqsum_t[0];

15
modules/ocl/src/kernels/imgproc_integral_sum.cl
Unescape View File

@ -94,6 +94,7 @@ kernel void integral_cols(__global uchar4 *src,__global int *sum ,
}
offset <<= 1;
}
barrier(CLK_LOCAL_MEM_FENCE);
if(lid < 2)
{
lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
@ -112,7 +113,8 @@ kernel void integral_cols(__global uchar4 *src,__global int *sum ,
lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai];
}
}
if(lid > 0 & (i+lid) <= rows){
barrier(CLK_LOCAL_MEM_FENCE);
if(lid > 0 && (i+lid) <= rows){
int loc_s0 = gid * dst_step + i + lid - 1 - pre_invalid * dst_step / 4, loc_s1 = loc_s0 + dst_step ;
lm_sum[0][bf_loc] += sum_t[0];
lm_sum[1][bf_loc] += sum_t[1];
@ -172,6 +174,7 @@ kernel void integral_rows(__global int4 *srcsum,__global int *sum ,
}
offset <<= 1;
}
barrier(CLK_LOCAL_MEM_FENCE);
if(lid < 2)
{
lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0;
@ -190,7 +193,7 @@ kernel void integral_rows(__global int4 *srcsum,__global int *sum ,
lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
if(gid == 0 && (i + lid) <= rows)
{
sum[sum_offset + i + lid] = 0;
@ -205,16 +208,16 @@ kernel void integral_rows(__global int4 *srcsum,__global int *sum ,
}
}
if(lid > 0 & (i+lid) <= rows){
if(lid > 0 && (i+lid) <= rows){
int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ;
lm_sum[0][bf_loc] += sum_t[0];
lm_sum[1][bf_loc] += sum_t[1];
lm_sum[0][bf_loc] += sum_t[0];
lm_sum[1][bf_loc] += sum_t[1];
sum_p = (__local int*)(&(lm_sum[0][bf_loc]));
for(int k = 0; k < 4; k++)
{
if(gid * 8 + k >= cols) break;
sum[loc_s0 + k * sum_step / 4] = sum_p[k];
}
}
sum_p = (__local int*)(&(lm_sum[1][bf_loc]));
for(int k = 0; k < 4; k++)
{

53
modules/ocl/src/matrix_operations.cpp
Unescape View File

@ -157,7 +157,7 @@ void convert_C3C4(const cl_mem &src, oclMat &dst, int srcStep)
sprintf(compile_option, "-D GENTYPE4=double4");
break;
default:
CV_Error(-217,"unknown depth");
CV_Error(CV_StsUnsupportedFormat,"unknown depth");
}
vector< pair<size_t, const void *> > args;
args.push_back( make_pair( sizeof(cl_mem), (void *)&src));
@ -205,7 +205,7 @@ void convert_C4C3(const oclMat &src, cl_mem &dst, int dstStep)
sprintf(compile_option, "-D GENTYPE4=double4");
break;
default:
CV_Error(-217,"unknown depth");
CV_Error(CV_StsUnsupportedFormat,"unknown depth");
}
vector< pair<size_t, const void *> > args;
@ -517,7 +517,7 @@ void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kern
args.push_back( make_pair( sizeof(cl_uchar4) , (void *)&val.uval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 1:
@ -536,7 +536,7 @@ void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kern
args.push_back( make_pair( sizeof(cl_char4) , (void *)&val.cval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 2:
@ -555,7 +555,7 @@ void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kern
args.push_back( make_pair( sizeof(cl_ushort4) , (void *)&val.usval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 3:
@ -574,7 +574,7 @@ void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kern
args.push_back( make_pair( sizeof(cl_short4) , (void *)&val.shval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 4:
@ -600,7 +600,7 @@ void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kern
args.push_back( make_pair( sizeof(cl_int4) , (void *)&val.ival ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 5:
@ -619,7 +619,7 @@ void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kern
args.push_back( make_pair( sizeof(cl_float4) , (void *)&val.fval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 6:
@ -638,12 +638,28 @@ void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kern
args.push_back( make_pair( sizeof(cl_double4) , (void *)&val.dval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
default:
CV_Error(-217,"unknown depth");
CV_Error(CV_StsUnsupportedFormat,"unknown depth");
}
#if CL_VERSION_1_2
if(dst.offset==0 && dst.cols==dst.wholecols)
{
clEnqueueFillBuffer(dst.clCxt->impl->clCmdQueue,(cl_mem)dst.data,args[0].second,args[0].first,0,dst.step*dst.rows,0,NULL,NULL);
}
else
{
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&step_in_pixel ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&offset_in_pixel));
openCLExecuteKernel(dst.clCxt , &operator_setTo, kernelName, globalThreads,
localThreads, args, -1, -1,compile_option);
}
#else
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows ));
@ -651,6 +667,7 @@ void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kern
args.push_back( make_pair( sizeof(cl_int) , (void *)&offset_in_pixel));
openCLExecuteKernel(dst.clCxt , &operator_setTo, kernelName, globalThreads,
localThreads, args, -1, -1,compile_option);
#endif
}
void set_to_withmask_run(const oclMat &dst, const Scalar &scalar, const oclMat &mask, string kernelName)
@ -696,7 +713,7 @@ void set_to_withmask_run(const oclMat &dst, const Scalar &scalar, const oclMat &
args.push_back( make_pair( sizeof(cl_uchar4) , (void *)&val.uval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 1:
@ -715,7 +732,7 @@ void set_to_withmask_run(const oclMat &dst, const Scalar &scalar, const oclMat &
args.push_back( make_pair( sizeof(cl_char4) , (void *)&val.cval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 2:
@ -734,7 +751,7 @@ void set_to_withmask_run(const oclMat &dst, const Scalar &scalar, const oclMat &
args.push_back( make_pair( sizeof(cl_ushort4) , (void *)&val.usval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 3:
@ -753,7 +770,7 @@ void set_to_withmask_run(const oclMat &dst, const Scalar &scalar, const oclMat &
args.push_back( make_pair( sizeof(cl_short4) , (void *)&val.shval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 4:
@ -772,7 +789,7 @@ void set_to_withmask_run(const oclMat &dst, const Scalar &scalar, const oclMat &
args.push_back( make_pair( sizeof(cl_int4) , (void *)&val.ival ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 5:
@ -791,7 +808,7 @@ void set_to_withmask_run(const oclMat &dst, const Scalar &scalar, const oclMat &
args.push_back( make_pair( sizeof(cl_float4) , (void *)&val.fval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
case 6:
@ -810,11 +827,11 @@ void set_to_withmask_run(const oclMat &dst, const Scalar &scalar, const oclMat &
args.push_back( make_pair( sizeof(cl_double4) , (void *)&val.dval ));
break;
default:
CV_Error(-217,"unsupported channels");
CV_Error(CV_StsUnsupportedFormat,"unsupported channels");
}
break;
default:
CV_Error(-217,"unknown depth");
CV_Error(CV_StsUnsupportedFormat,"unknown depth");
}
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols ));

4
modules/ocl/src/split_merge.cpp
Unescape View File

@ -160,7 +160,7 @@ namespace cv
{
if(mat_dst.clCxt -> impl -> double_support ==0 && mat_dst.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}
@ -299,7 +299,7 @@ namespace cv
if(mat_src.clCxt -> impl -> double_support ==0 && mat_src.type() == CV_64F)
{
CV_Error(-217,"Selected device don't support double\r\n");
CV_Error(CV_GpuNotSupported,"Selected device don't support double\r\n");
return;
}

15
modules/ocl/test/main.cpp
Unescape View File

@ -74,13 +74,6 @@ void print_info()
}
#if PERF_TEST_OCL
int main(int argc, char **argv)
{
run_perf_test();
return 0;
}
#else
int main(int argc, char **argv)
{
TS::ptr()->init("ocl");
@ -89,16 +82,16 @@ int main(int argc, char **argv)
print_info();
std::vector<cv::ocl::Info> oclinfo;
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
if(devnums<1){
int devnums = getDevice(oclinfo);
if(devnums<1)
{
std::cout << "no device found\n";
return -1;
}
return RUN_ALL_TESTS();
}
#endif // PERF_TEST_OCL
#else // HAVE_OPENC
#else // DON'T HAVE_OPENCL
int main()
{

11
modules/ocl/test/test_arithm.cpp
Unescape View File

@ -94,7 +94,7 @@ PARAM_TEST_CASE(ArithmTestBase, MatType, bool)
cv::Mat mask_roi;
cv::Mat dst_roi;
cv::Mat dst1_roi; //bak
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
cv::ocl::oclMat gdst1_whole; //bak
@ -115,6 +115,7 @@ PARAM_TEST_CASE(ArithmTestBase, MatType, bool)
cv::Size size(MWIDTH, MHEIGHT);
mat1 = randomMat(rng, size, type, 5, 16, false);
//mat2 = randomMat(rng, size, type, 5, 16, false);
mat2 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
dst1 = randomMat(rng, size, type, 5, 16, false);
@ -124,10 +125,10 @@ PARAM_TEST_CASE(ArithmTestBase, MatType, bool)
val = cv::Scalar(rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0));
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()

6
modules/ocl/test/test_blend.cpp
Unescape View File

@ -33,7 +33,7 @@ void blendLinearGold(const cv::Mat& img1, const cv::Mat& img2, const cv::Mat& we
PARAM_TEST_CASE(Blend, cv::Size, MatType/*, UseRoi*/)
{
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
cv::Size size;
int type;
bool useRoi;
@ -45,8 +45,8 @@ PARAM_TEST_CASE(Blend, cv::Size, MatType/*, UseRoi*/)
type = GET_PARAM(1);
/*useRoi = GET_PARAM(3);*/
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
}
};

6
modules/ocl/test/test_canny.cpp
Unescape View File

@ -64,13 +64,13 @@ PARAM_TEST_CASE(Canny, AppertureSize, L2gradient)
bool useL2gradient;
cv::Mat edges_gold;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
virtual void SetUp()
{
apperture_size = GET_PARAM(0);
useL2gradient = GET_PARAM(1);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
}
};

6
modules/ocl/test/test_columnsum.cpp
Unescape View File

@ -60,14 +60,14 @@ PARAM_TEST_CASE(ColumnSum, cv::Size, bool )
cv::Size size;
cv::Mat src;
bool useRoi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
virtual void SetUp()
{
size = GET_PARAM(0);
useRoi = GET_PARAM(1);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
}
};

6
modules/ocl/test/test_fft.cpp
Unescape View File

@ -52,11 +52,11 @@ PARAM_TEST_CASE(Dft, cv::Size, bool)
{
cv::Size dft_size;
bool dft_rows;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
virtual void SetUp()
{
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//int devnums = getDevice(oclinfo);
// CV_Assert(devnums > 0);
dft_size = GET_PARAM(0);
dft_rows = GET_PARAM(1);
}

72
modules/ocl/test/test_filters.cpp
Unescape View File

@ -85,7 +85,7 @@ PARAM_TEST_CASE(FilterTestBase, MatType, bool)
cv::Mat mask_roi;
cv::Mat dst_roi;
cv::Mat dst1_roi; //bak
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
cv::ocl::oclMat gdst1_whole; //bak
@ -185,7 +185,7 @@ PARAM_TEST_CASE(Blur, MatType, cv::Size, int)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -204,10 +204,10 @@ PARAM_TEST_CASE(Blur, MatType, cv::Size, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -285,7 +285,7 @@ PARAM_TEST_CASE(LaplacianTestBase, MatType, int)
//src mat with roi
cv::Mat mat_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -304,10 +304,10 @@ PARAM_TEST_CASE(LaplacianTestBase, MatType, int)
mat = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -388,7 +388,7 @@ PARAM_TEST_CASE(ErodeDilateBase, MatType, bool)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -409,10 +409,10 @@ PARAM_TEST_CASE(ErodeDilateBase, MatType, bool)
// rng.fill(kernel, cv::RNG::UNIFORM, cv::Scalar::all(0), cv::Scalar::all(3));
kernel = randomMat(rng, Size(3, 3), CV_8UC1, 0, 3, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -529,7 +529,7 @@ PARAM_TEST_CASE(Sobel, MatType, int, int, int, int)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -552,10 +552,10 @@ PARAM_TEST_CASE(Sobel, MatType, int, int, int, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -633,7 +633,7 @@ PARAM_TEST_CASE(Scharr, MatType, int, int, int)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -656,10 +656,10 @@ PARAM_TEST_CASE(Scharr, MatType, int, int, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -740,7 +740,7 @@ PARAM_TEST_CASE(GaussianBlur, MatType, cv::Size, int)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -763,10 +763,10 @@ PARAM_TEST_CASE(GaussianBlur, MatType, cv::Size, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -841,18 +841,18 @@ INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine(Values(CV_8UC1, CV_8UC4, CV_32FC
INSTANTIATE_TEST_CASE_P(Filter, Sobel, Combine(Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4),
Values(1, 2), Values(0, 1), Values(3, 5, 7), Values((MatType)cv::BORDER_CONSTANT,
(MatType)cv::BORDER_REPLICATE, (MatType)cv::BORDER_REFLECT, (MatType)cv::BORDER_REFLECT_101)));
Values(1, 2), Values(0, 1), Values(3, 5), Values((MatType)cv::BORDER_CONSTANT,
(MatType)cv::BORDER_REPLICATE)));
INSTANTIATE_TEST_CASE_P(Filter, Scharr, Combine(
Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(0, 1), Values(0, 1),
Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE, (MatType)cv::BORDER_REFLECT, (MatType)cv::BORDER_REFLECT_101)));
Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE)));
INSTANTIATE_TEST_CASE_P(Filter, GaussianBlur, Combine(
Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4),
Values(cv::Size(3, 3), cv::Size(5, 5), cv::Size(7, 7)),
Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE, (MatType)cv::BORDER_REFLECT, (MatType)cv::BORDER_REFLECT_101)));
Values(cv::Size(3, 3), cv::Size(5, 5)),
Values((MatType)cv::BORDER_CONSTANT, (MatType)cv::BORDER_REPLICATE)));

4
modules/ocl/test/test_gemm.cpp
Unescape View File

@ -53,13 +53,13 @@ PARAM_TEST_CASE(Gemm, int, cv::Size, int)
int type;
cv::Size mat_size;
int flags;
vector<cv::ocl::Info> info;
//vector<cv::ocl::Info> info;
virtual void SetUp()
{
type = GET_PARAM(0);
mat_size = GET_PARAM(1);
flags = GET_PARAM(2);
cv::ocl::getDevice(info);
//cv::ocl::getDevice(info);
}
};

12
modules/ocl/test/test_haar.cpp
Unescape View File

@ -63,7 +63,7 @@ struct getRect
PARAM_TEST_CASE(HaarTestBase, int, int)
{
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
cv::ocl::OclCascadeClassifier cascade, nestedCascade;
cv::CascadeClassifier cpucascade, cpunestedCascade;
// Mat img;
@ -91,11 +91,11 @@ PARAM_TEST_CASE(HaarTestBase, int, int)
return;
}
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
cv::ocl::setBinpath("E:\\");
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
//cv::ocl::setBinpath("E:\\");
}
};

4
modules/ocl/test/test_hog.cpp
Unescape View File

@ -53,12 +53,12 @@ PARAM_TEST_CASE(HOG,cv::Size,int)
{
cv::Size winSize;
int type;
vector<cv::ocl::Info> info;
//vector<cv::ocl::Info> info;
virtual void SetUp()
{
winSize = GET_PARAM(0);
type = GET_PARAM(1);
cv::ocl::getDevice(info);
//cv::ocl::getDevice(info);
}
};

86
modules/ocl/test/test_imgproc.cpp
Unescape View File

@ -328,7 +328,7 @@ PARAM_TEST_CASE(ImgprocTestBase, MatType, MatType, MatType, MatType, MatType, bo
cv::Mat mask_roi;
cv::Mat dst_roi;
cv::Mat dst1_roi; //bak
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl mat
cv::ocl::oclMat clmat1;
cv::ocl::oclMat clmat2;
@ -353,10 +353,10 @@ PARAM_TEST_CASE(ImgprocTestBase, MatType, MatType, MatType, MatType, MatType, bo
cv::RNG &rng = TS::ptr()->get_rng();
cv::Size size(MWIDTH, MHEIGHT);
double min = 1, max = 20;
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
if(type1 != nulltype)
{
@ -661,7 +661,7 @@ PARAM_TEST_CASE(WarpTestBase, MatType, int)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -681,10 +681,10 @@ PARAM_TEST_CASE(WarpTestBase, MatType, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -807,7 +807,7 @@ PARAM_TEST_CASE(Remap, MatType, MatType, MatType, int, int)
cv::Mat map1;
cv::Mat map2;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
int src_roicols;
int src_roirows;
@ -853,8 +853,8 @@ PARAM_TEST_CASE(Remap, MatType, MatType, MatType, int, int)
bordertype = GET_PARAM(4);
// borderValue = GET_PARAM(6);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
cv::RNG& rng = TS::ptr()->get_rng();
//cv::Size size = cv::Size(20, 20);
@ -1006,7 +1006,7 @@ PARAM_TEST_CASE(Resize, MatType, cv::Size, double, double, int)
int dstx;
int dsty;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
@ -1045,10 +1045,10 @@ PARAM_TEST_CASE(Resize, MatType, cv::Size, double, double, int)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, dsize, type, 5, 16, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -1133,7 +1133,7 @@ PARAM_TEST_CASE(Threshold, MatType, ThreshOp)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -1152,10 +1152,10 @@ PARAM_TEST_CASE(Threshold, MatType, ThreshOp)
mat1 = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -1240,7 +1240,7 @@ PARAM_TEST_CASE(meanShiftTestBase, MatType, MatType, int, int, cv::TermCriteria)
cv::ocl::oclMat gdst;
cv::ocl::oclMat gdstCoor;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl mat with roi
cv::ocl::oclMat gsrc_roi;
cv::ocl::oclMat gdst_roi;
@ -1263,10 +1263,10 @@ PARAM_TEST_CASE(meanShiftTestBase, MatType, MatType, int, int, cv::TermCriteria)
dst = randomMat(rng, size, type, 5, 16, false);
dstCoor = randomMat(rng, size, typeCoor, 5, 16, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -1378,21 +1378,21 @@ INSTANTIATE_TEST_CASE_P(ImgprocTestBase, equalizeHist, Combine(
// NULL_TYPE,
// Values(false))); // Values(false) is the reserved parameter
//
//INSTANTIATE_TEST_CASE_P(ImgprocTestBase, cornerMinEigenVal, Combine(
// Values(CV_8UC1,CV_32FC1),
// NULL_TYPE,
// ONE_TYPE(CV_32FC1),
// NULL_TYPE,
// NULL_TYPE,
// Values(false))); // Values(false) is the reserved parameter
//
//INSTANTIATE_TEST_CASE_P(ImgprocTestBase, cornerHarris, Combine(
// Values(CV_8UC1,CV_32FC1),
// NULL_TYPE,
// ONE_TYPE(CV_32FC1),
// NULL_TYPE,
// NULL_TYPE,
// Values(false))); // Values(false) is the reserved parameter
INSTANTIATE_TEST_CASE_P(ImgprocTestBase, cornerMinEigenVal, Combine(
Values(CV_8UC1,CV_32FC1),
NULL_TYPE,
ONE_TYPE(CV_32FC1),
NULL_TYPE,
NULL_TYPE,
Values(false))); // Values(false) is the reserved parameter
INSTANTIATE_TEST_CASE_P(ImgprocTestBase, cornerHarris, Combine(
Values(CV_8UC1,CV_32FC1),
NULL_TYPE,
ONE_TYPE(CV_32FC1),
NULL_TYPE,
NULL_TYPE,
Values(false))); // Values(false) is the reserved parameter
INSTANTIATE_TEST_CASE_P(ImgprocTestBase, integral, Combine(

12
modules/ocl/test/test_match_template.cpp
Unescape View File

@ -60,7 +60,7 @@ PARAM_TEST_CASE(MatchTemplate8U, cv::Size, TemplateSize, Channels, TemplateMetho
cv::Size templ_size;
int cn;
int method;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
virtual void SetUp()
{
@ -68,8 +68,8 @@ PARAM_TEST_CASE(MatchTemplate8U, cv::Size, TemplateSize, Channels, TemplateMetho
templ_size = GET_PARAM(1);
cn = GET_PARAM(2);
method = GET_PARAM(3);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
}
};
@ -112,7 +112,7 @@ PARAM_TEST_CASE(MatchTemplate32F, cv::Size, TemplateSize, Channels, TemplateMeth
cv::Size templ_size;
int cn;
int method;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
virtual void SetUp()
{
@ -120,8 +120,8 @@ PARAM_TEST_CASE(MatchTemplate32F, cv::Size, TemplateSize, Channels, TemplateMeth
templ_size = GET_PARAM(1);
cn = GET_PARAM(2);
method = GET_PARAM(3);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
}
};

40
modules/ocl/test/test_matrix_operation.cpp
Unescape View File

@ -72,7 +72,7 @@ PARAM_TEST_CASE(ConvertToTestBase, MatType, MatType)
//src mat with roi
cv::Mat mat_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -91,10 +91,10 @@ PARAM_TEST_CASE(ConvertToTestBase, MatType, MatType)
mat = randomMat(rng, size, type, 5, 16, false);
dst = randomMat(rng, size, type, 5, 16, false);
//std::vector<cv::ocl::Info> oclinfo;
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -175,7 +175,7 @@ PARAM_TEST_CASE(CopyToTestBase, MatType, bool)
cv::Mat mat_roi;
cv::Mat mask_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -197,10 +197,10 @@ PARAM_TEST_CASE(CopyToTestBase, MatType, bool)
cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -301,7 +301,7 @@ PARAM_TEST_CASE(SetToTestBase, MatType, bool)
//src mat with roi
cv::Mat mat_roi;
cv::Mat mask_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gmat_whole;
@ -322,10 +322,10 @@ PARAM_TEST_CASE(SetToTestBase, MatType, bool)
cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
val = cv::Scalar(rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0));
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -417,7 +417,7 @@ PARAM_TEST_CASE(convertC3C4, MatType, cv::Size)
//src mat with roi
cv::Mat mat1_roi;
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -433,10 +433,10 @@ PARAM_TEST_CASE(convertC3C4, MatType, cv::Size)
//dst = randomMat(rng, size, type, 5, 16, false);
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[1]);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[1]);
}
void random_roi()

10
modules/ocl/test/test_pyrdown.cpp
Unescape View File

@ -89,7 +89,7 @@ PARAM_TEST_CASE(PyrDown, MatType, bool)
cv::Mat mask_roi;
cv::Mat dst_roi;
cv::Mat dst1_roi; //bak
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
cv::ocl::oclMat gdst1_whole; //bak
@ -119,10 +119,10 @@ PARAM_TEST_CASE(PyrDown, MatType, bool)
val = cv::Scalar(rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0));
int devnums = getDevice(oclinfo);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void Cleanup()

6
modules/ocl/test/test_pyrup.cpp
Unescape View File

@ -53,12 +53,12 @@ PARAM_TEST_CASE(PyrUp,cv::Size,int)
{
cv::Size size;
int type;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
virtual void SetUp()
{
int devnums = cv::ocl::getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//int devnums = cv::ocl::getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
size = GET_PARAM(0);
type = GET_PARAM(1);
}

20
modules/ocl/test/test_split_merge.cpp
Unescape View File

@ -87,7 +87,7 @@ PARAM_TEST_CASE(MergeTestBase, MatType, int)
//dst mat with roi
cv::Mat dst_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst_whole;
@ -112,10 +112,10 @@ PARAM_TEST_CASE(MergeTestBase, MatType, int)
mat4 = randomMat(rng, size, CV_MAKETYPE(type, 1), 5, 16, false);
dst = randomMat(rng, size, CV_MAKETYPE(type, channels), 5, 16, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()
@ -240,7 +240,7 @@ PARAM_TEST_CASE(SplitTestBase, MatType, int)
cv::Mat dst2_roi;
cv::Mat dst3_roi;
cv::Mat dst4_roi;
std::vector<cv::ocl::Info> oclinfo;
//std::vector<cv::ocl::Info> oclinfo;
//ocl dst mat for testing
cv::ocl::oclMat gdst1_whole;
cv::ocl::oclMat gdst2_whole;
@ -268,10 +268,10 @@ PARAM_TEST_CASE(SplitTestBase, MatType, int)
dst3 = randomMat(rng, size, CV_MAKETYPE(type, 1), 5, 16, false);
dst4 = randomMat(rng, size, CV_MAKETYPE(type, 1), 5, 16, false);
int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
CV_Assert(devnums > 0);
//if you want to use undefault device, set it here
//setDevice(oclinfo[0]);
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
//CV_Assert(devnums > 0);
////if you want to use undefault device, set it here
////setDevice(oclinfo[0]);
}
void random_roi()

Write Preview
Loading…
Cancel
Save