fixed octave computation in SURF_GPU

used random images in gpu filter tests

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 0e6dda9..08c6be5 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -1516,13 +1516,14 @@ class CV_EXPORTS SURF_GPU
 public:\r
     enum KeypointLayout\r
     {\r
-        SF_X = 0,\r
-        SF_Y,\r
-        SF_LAPLACIAN,\r
-        SF_SIZE,\r
-        SF_DIR,\r
-        SF_HESSIAN,\r
-        SF_FEATURE_STRIDE\r
+        X_ROW = 0,\r
+        Y_ROW,\r
+        LAPLACIAN_ROW,\r
+        OCTAVE_ROW,\r
+        SIZE_ROW,\r
+        ANGLE_ROW,\r
+        HESSIAN_ROW,\r
+        ROWS_COUNT\r
     };\r
 \r
     //! the default constructor\r

diff --git a/modules/gpu/src/cuda/surf.cu b/modules/gpu/src/cuda/surf.cu
index 9d265cc..76181bc 100644 (file)
--- a/modules/gpu/src/cuda/surf.cu
+++ b/modules/gpu/src/cuda/surf.cu
@@ -117,7 +117,7 @@ namespace cv { namespace gpu { namespace device
         template <int N> __device__ float icvCalcHaarPatternSum(const float src[][5], int oldSize, int newSize, int y, int x)\r
         {\r
         #if __CUDA_ARCH__ >= 200\r
-            typedef double real_t;        \r
+            typedef double real_t;\r
         #else\r
             typedef float  real_t;\r
         #endif\r
@@ -248,7 +248,7 @@ namespace cv { namespace gpu { namespace device
         template <typename Mask>\r
         __global__ void icvFindMaximaInLayer(const PtrStepf det, const PtrStepf trace, int4* maxPosBuffer, unsigned int* maxCounter)\r
         {\r
-            #if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 110\r
+            #if __CUDA_ARCH__ >= 110\r
 \r
             extern __shared__ float N9[];\r
 \r
@@ -368,10 +368,10 @@ namespace cv { namespace gpu { namespace device
         // INTERPOLATION\r
 \r
         __global__ void icvInterpolateKeypoint(const PtrStepf det, const int4* maxPosBuffer,\r
-            float* featureX, float* featureY, int* featureLaplacian, float* featureSize, float* featureHessian,\r
+            float* featureX, float* featureY, int* featureLaplacian, int* featureOctave, float* featureSize, float* featureHessian,\r
             unsigned int* featureCounter)\r
         {\r
-            #if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 110\r
+            #if __CUDA_ARCH__ >= 110\r
 \r
             const int4 maxPos = maxPosBuffer[blockIdx.x];\r
 \r
@@ -459,6 +459,7 @@ namespace cv { namespace gpu { namespace device
                                 featureX[ind] = px;\r
                                 featureY[ind] = py;\r
                                 featureLaplacian[ind] = maxPos.w;\r
+                                featureOctave[ind] = c_octave;\r
                                 featureSize[ind] = psize;\r
                                 featureHessian[ind] = N9[1][1][1];\r
                             }\r
@@ -471,7 +472,7 @@ namespace cv { namespace gpu { namespace device
         }\r
 \r
         void icvInterpolateKeypoint_gpu(const PtrStepf& det, const int4* maxPosBuffer, unsigned int maxCounter, \r
-            float* featureX, float* featureY, int* featureLaplacian, float* featureSize, float* featureHessian, \r
+            float* featureX, float* featureY, int* featureLaplacian, int* featureOctave, float* featureSize, float* featureHessian, \r
             unsigned int* featureCounter)\r
         {\r
             dim3 threads;\r
@@ -482,7 +483,7 @@ namespace cv { namespace gpu { namespace device
             dim3 grid;\r
             grid.x = maxCounter;\r
 \r
-            icvInterpolateKeypoint<<<grid, threads>>>(det, maxPosBuffer, featureX, featureY, featureLaplacian, featureSize, featureHessian, featureCounter);\r
+            icvInterpolateKeypoint<<<grid, threads>>>(det, maxPosBuffer, featureX, featureY, featureLaplacian, featureOctave, featureSize, featureHessian, featureCounter);\r
             cudaSafeCall( cudaGetLastError() );\r
 \r
             cudaSafeCall( cudaDeviceSynchronize() );\r

diff --git a/modules/gpu/src/filtering.cpp b/modules/gpu/src/filtering.cpp
index 4a5f032..ac61d20 100644 (file)
--- a/modules/gpu/src/filtering.cpp
+++ b/modules/gpu/src/filtering.cpp
@@ -948,7 +948,9 @@ namespace
         {\r
             DeviceInfo devInfo;\r
             int cc = devInfo.majorVersion() * 10 + devInfo.minorVersion();\r
-            CV_Assert(cc >= 20 || ksize <= 16);\r
+            if (ksize > 16 && cc < 20)\r
+                CV_Error(CV_StsNotImplemented, "column linear filter doesn't implemented for kernel size > 16 for device with compute capabilities less than 2.0");\r
+\r
             func(src, dst, kernel.ptr<float>(), ksize, anchor, brd_type, cc, StreamAccessor::getStream(s));\r
         }\r
 \r

diff --git a/modules/gpu/src/surf.cpp b/modules/gpu/src/surf.cpp
index 2bc9f58..dea51a9 100644 (file)
--- a/modules/gpu/src/surf.cpp
+++ b/modules/gpu/src/surf.cpp
@@ -80,7 +80,7 @@ namespace cv { namespace gpu { namespace device
             int img_rows, int img_cols, int octave, bool use_mask, int nLayers);\r
 \r
         void icvInterpolateKeypoint_gpu(const PtrStepf& det, const int4* maxPosBuffer, unsigned int maxCounter,\r
-            float* featureX, float* featureY, int* featureLaplacian, float* featureSize, float* featureHessian,\r
+            float* featureX, float* featureY, int* featureLaplacian, int* featureOctave, float* featureSize, float* featureHessian,\r
             unsigned int* featureCounter);\r
 \r
         void icvCalcOrientation_gpu(const float* featureX, const float* featureY, const float* featureSize, float* featureDir, int nFeatures);\r
@@ -161,7 +161,7 @@ namespace
             ensureSizeIsEnough(img_rows * (surf_.nOctaveLayers + 2), img_cols, CV_32FC1, surf_.trace);\r
 \r
             ensureSizeIsEnough(1, maxCandidates, CV_32SC4, surf_.maxPosBuffer);\r
-            ensureSizeIsEnough(SURF_GPU::SF_FEATURE_STRIDE, maxFeatures, CV_32FC1, keypoints);\r
+            ensureSizeIsEnough(SURF_GPU::ROWS_COUNT, maxFeatures, CV_32FC1, keypoints);\r
             keypoints.setTo(Scalar::all(0));\r
 \r
             for (int octave = 0; octave < surf_.nOctaves; ++octave)\r
@@ -183,9 +183,10 @@ namespace
                 if (maxCounter > 0)\r
                 {\r
                     icvInterpolateKeypoint_gpu(surf_.det, surf_.maxPosBuffer.ptr<int4>(), maxCounter,\r
-                        keypoints.ptr<float>(SURF_GPU::SF_X), keypoints.ptr<float>(SURF_GPU::SF_Y),\r
-                        keypoints.ptr<int>(SURF_GPU::SF_LAPLACIAN), keypoints.ptr<float>(SURF_GPU::SF_SIZE),\r
-                        keypoints.ptr<float>(SURF_GPU::SF_HESSIAN), counters.ptr<unsigned int>());\r
+                        keypoints.ptr<float>(SURF_GPU::X_ROW), keypoints.ptr<float>(SURF_GPU::Y_ROW),\r
+                        keypoints.ptr<int>(SURF_GPU::LAPLACIAN_ROW), keypoints.ptr<int>(SURF_GPU::OCTAVE_ROW), \r
+                        keypoints.ptr<float>(SURF_GPU::SIZE_ROW), keypoints.ptr<float>(SURF_GPU::HESSIAN_ROW), \r
+                        counters.ptr<unsigned int>());\r
                 }\r
             }\r
             unsigned int featureCounter;\r
@@ -195,7 +196,7 @@ namespace
             keypoints.cols = featureCounter;\r
 \r
             if (surf_.upright)\r
-                keypoints.row(SURF_GPU::SF_DIR).setTo(Scalar::all(90.0));\r
+                keypoints.row(SURF_GPU::ANGLE_ROW).setTo(Scalar::all(90.0));\r
             else\r
                 findOrientation(keypoints);\r
         }\r
@@ -205,8 +206,8 @@ namespace
             const int nFeatures = keypoints.cols;\r
             if (nFeatures > 0)\r
             {\r
-                icvCalcOrientation_gpu(keypoints.ptr<float>(SURF_GPU::SF_X), keypoints.ptr<float>(SURF_GPU::SF_Y),\r
-                    keypoints.ptr<float>(SURF_GPU::SF_SIZE), keypoints.ptr<float>(SURF_GPU::SF_DIR), nFeatures);\r
+                icvCalcOrientation_gpu(keypoints.ptr<float>(SURF_GPU::X_ROW), keypoints.ptr<float>(SURF_GPU::Y_ROW),\r
+                    keypoints.ptr<float>(SURF_GPU::SIZE_ROW), keypoints.ptr<float>(SURF_GPU::ANGLE_ROW), nFeatures);\r
             }\r
         }\r
 \r
@@ -216,8 +217,8 @@ namespace
             if (nFeatures > 0)\r
             {\r
                 ensureSizeIsEnough(nFeatures, descriptorSize, CV_32F, descriptors);\r
-                compute_descriptors_gpu(descriptors, keypoints.ptr<float>(SURF_GPU::SF_X), keypoints.ptr<float>(SURF_GPU::SF_Y),\r
-                    keypoints.ptr<float>(SURF_GPU::SF_SIZE), keypoints.ptr<float>(SURF_GPU::SF_DIR), nFeatures);\r
+                compute_descriptors_gpu(descriptors, keypoints.ptr<float>(SURF_GPU::X_ROW), keypoints.ptr<float>(SURF_GPU::Y_ROW),\r
+                    keypoints.ptr<float>(SURF_GPU::SIZE_ROW), keypoints.ptr<float>(SURF_GPU::ANGLE_ROW), nFeatures);\r
             }\r
         }\r
 \r
@@ -266,20 +267,22 @@ void cv::gpu::SURF_GPU::uploadKeypoints(const vector<KeyPoint>& keypoints, GpuMa
         keypointsGPU.release();\r
     else\r
     {\r
-        Mat keypointsCPU(SURF_GPU::SF_FEATURE_STRIDE, static_cast<int>(keypoints.size()), CV_32FC1);\r
+        Mat keypointsCPU(SURF_GPU::ROWS_COUNT, static_cast<int>(keypoints.size()), CV_32FC1);\r
 \r
-        float* kp_x = keypointsCPU.ptr<float>(SURF_GPU::SF_X);\r
-        float* kp_y = keypointsCPU.ptr<float>(SURF_GPU::SF_Y);\r
-        int* kp_laplacian = keypointsCPU.ptr<int>(SURF_GPU::SF_LAPLACIAN);\r
-        float* kp_size = keypointsCPU.ptr<float>(SURF_GPU::SF_SIZE);\r
-        float* kp_dir = keypointsCPU.ptr<float>(SURF_GPU::SF_DIR);\r
-        float* kp_hessian = keypointsCPU.ptr<float>(SURF_GPU::SF_HESSIAN);\r
+        float* kp_x = keypointsCPU.ptr<float>(SURF_GPU::X_ROW);\r
+        float* kp_y = keypointsCPU.ptr<float>(SURF_GPU::Y_ROW);\r
+        int* kp_laplacian = keypointsCPU.ptr<int>(SURF_GPU::LAPLACIAN_ROW);\r
+        int* kp_octave = keypointsCPU.ptr<int>(SURF_GPU::OCTAVE_ROW);\r
+        float* kp_size = keypointsCPU.ptr<float>(SURF_GPU::SIZE_ROW);\r
+        float* kp_dir = keypointsCPU.ptr<float>(SURF_GPU::ANGLE_ROW);\r
+        float* kp_hessian = keypointsCPU.ptr<float>(SURF_GPU::HESSIAN_ROW);\r
 \r
         for (size_t i = 0, size = keypoints.size(); i < size; ++i)\r
         {\r
             const KeyPoint& kp = keypoints[i];\r
             kp_x[i] = kp.pt.x;\r
             kp_y[i] = kp.pt.y;\r
+            kp_octave[i] = kp.octave;\r
             kp_size[i] = kp.size;\r
             kp_dir[i] = kp.angle;\r
             kp_hessian[i] = kp.response;\r
@@ -290,30 +293,6 @@ void cv::gpu::SURF_GPU::uploadKeypoints(const vector<KeyPoint>& keypoints, GpuMa
     }\r
 }\r
 \r
-namespace\r
-{\r
-    int getPointOctave(float size, const SURF_GPU& params)\r
-    {\r
-        int best_octave = 0;\r
-        float min_diff = numeric_limits<float>::max();\r
-        for (int octave = 1; octave < params.nOctaves; ++octave)\r
-        {\r
-            for (int layer = 0; layer < params.nOctaveLayers; ++layer)\r
-            {\r
-                float diff = std::abs(size - (float)calcSize(octave, layer));\r
-                if (min_diff > diff)\r
-                {\r
-                    min_diff = diff;\r
-                    best_octave = octave;\r
-                    if (min_diff == 0)\r
-                        return best_octave;\r
-                }\r
-            }\r
-        }\r
-        return best_octave;\r
-    }\r
-}\r
-\r
 void cv::gpu::SURF_GPU::downloadKeypoints(const GpuMat& keypointsGPU, vector<KeyPoint>& keypoints)\r
 {\r
     const int nFeatures = keypointsGPU.cols;\r
@@ -322,18 +301,19 @@ void cv::gpu::SURF_GPU::downloadKeypoints(const GpuMat& keypointsGPU, vector<Key
         keypoints.clear();\r
     else\r
     {\r
-        CV_Assert(keypointsGPU.type() == CV_32FC1 && keypointsGPU.rows == SF_FEATURE_STRIDE);\r
+        CV_Assert(keypointsGPU.type() == CV_32FC1 && keypointsGPU.rows == ROWS_COUNT);\r
 \r
         Mat keypointsCPU(keypointsGPU);\r
 \r
         keypoints.resize(nFeatures);\r
-\r
-        float* kp_x = keypointsCPU.ptr<float>(SF_X);\r
-        float* kp_y = keypointsCPU.ptr<float>(SF_Y);\r
-        int* kp_laplacian = keypointsCPU.ptr<int>(SF_LAPLACIAN);\r
-        float* kp_size = keypointsCPU.ptr<float>(SF_SIZE);\r
-        float* kp_dir = keypointsCPU.ptr<float>(SF_DIR);\r
-        float* kp_hessian = keypointsCPU.ptr<float>(SF_HESSIAN);\r
+        \r
+        float* kp_x = keypointsCPU.ptr<float>(SURF_GPU::X_ROW);\r
+        float* kp_y = keypointsCPU.ptr<float>(SURF_GPU::Y_ROW);\r
+        int* kp_laplacian = keypointsCPU.ptr<int>(SURF_GPU::LAPLACIAN_ROW);\r
+        int* kp_octave = keypointsCPU.ptr<int>(SURF_GPU::OCTAVE_ROW);\r
+        float* kp_size = keypointsCPU.ptr<float>(SURF_GPU::SIZE_ROW);\r
+        float* kp_dir = keypointsCPU.ptr<float>(SURF_GPU::ANGLE_ROW);\r
+        float* kp_hessian = keypointsCPU.ptr<float>(SURF_GPU::HESSIAN_ROW);\r
 \r
         for (int i = 0; i < nFeatures; ++i)\r
         {\r
@@ -341,10 +321,10 @@ void cv::gpu::SURF_GPU::downloadKeypoints(const GpuMat& keypointsGPU, vector<Key
             kp.pt.x = kp_x[i];\r
             kp.pt.y = kp_y[i];\r
             kp.class_id = kp_laplacian[i];\r
+            kp.octave = kp_octave[i];\r
             kp.size = kp_size[i];\r
             kp.angle = kp_dir[i];\r
             kp.response = kp_hessian[i];\r
-            kp.octave = getPointOctave(kp.size, *this);\r
         }\r
     }\r
 }\r

diff --git a/modules/gpu/test/test_core.cpp b/modules/gpu/test/test_core.cpp
index 348336d..a00c6fa 100644 (file)
--- a/modules/gpu/test/test_core.cpp
+++ b/modules/gpu/test/test_core.cpp
@@ -437,7 +437,7 @@ TEST_P(Multiply_Array, WithScale)
         cv::Mat dst_gold;\r
         cv::multiply(mat1, mat2, dst_gold, scale, depth.second);\r
 \r
-        EXPECT_MAT_NEAR(dst_gold, dst, depth.first >= CV_32F || depth.second >= CV_32F ? 1e-4 : 0.0);\r
+        EXPECT_MAT_NEAR(dst_gold, dst, 1.0);\r
     }\r
 }\r
 \r
@@ -2715,7 +2715,7 @@ TEST_P(Sum, Sqr)
 \r
     cv::Scalar val_gold = sqrSumGold(src);\r
 \r
-    EXPECT_SCALAR_NEAR(val_gold, val, CV_MAT_DEPTH(type) < CV_32F ? 0.0 : 10);\r
+    EXPECT_SCALAR_NEAR(val_gold, val, CV_MAT_DEPTH(type) < CV_32F ? 0.0 : 0.5);\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Core, Sum, testing::Combine(\r

diff --git a/modules/gpu/test/test_features2d.cpp b/modules/gpu/test/test_features2d.cpp
index 14642ab..2121820 100644 (file)
--- a/modules/gpu/test/test_features2d.cpp
+++ b/modules/gpu/test/test_features2d.cpp
@@ -85,7 +85,7 @@ testing::AssertionResult assertKeyPointsEquals(const char* gold_expr, const char
     std::sort(actual.begin(), actual.end(), KeyPointLess());\r
     std::sort(gold.begin(), gold.end(), KeyPointLess());\r
 \r
-    for (size_t i; i < gold.size(); ++i)\r
+    for (size_t i = 0; i < gold.size(); ++i)\r
     {\r
         const cv::KeyPoint& p1 = gold[i];\r
         const cv::KeyPoint& p2 = actual[i];\r

diff --git a/modules/gpu/test/test_filters.cpp b/modules/gpu/test/test_filters.cpp
index 1424165..6b9f23c 100644 (file)
--- a/modules/gpu/test/test_filters.cpp
+++ b/modules/gpu/test/test_filters.cpp
@@ -68,59 +68,45 @@ cv::Mat getInnerROI(cv::InputArray m, int ksize)
 \r
 IMPLEMENT_PARAM_CLASS(Anchor, cv::Point)\r
 \r
-PARAM_TEST_CASE(Blur, cv::gpu::DeviceInfo, KSize, Anchor, UseRoi)\r
+PARAM_TEST_CASE(Blur, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, Anchor, UseRoi)\r
 {\r
     cv::gpu::DeviceInfo devInfo;\r
+    cv::Size size;\r
+    int type;\r
     cv::Size ksize;\r
     cv::Point anchor;\r
     bool useRoi;\r
 \r
-    cv::Mat img;\r
-\r
     virtual void SetUp()\r
     {\r
         devInfo = GET_PARAM(0);\r
-        ksize = GET_PARAM(1);\r
-        anchor = GET_PARAM(2);\r
-        useRoi = GET_PARAM(3);\r
+        size = GET_PARAM(1);\r
+        type = GET_PARAM(2);\r
+        ksize = GET_PARAM(3);\r
+        anchor = GET_PARAM(4);\r
+        useRoi = GET_PARAM(5);\r
 \r
         cv::gpu::setDevice(devInfo.deviceID());\r
-\r
-        img = readImage("stereobp/aloe-L.png");\r
-        ASSERT_FALSE(img.empty());\r
     }\r
 };\r
 \r
-TEST_P(Blur, Gray)\r
-{\r
-    cv::Mat img_gray;\r
-    cv::cvtColor(img, img_gray, CV_BGR2GRAY);\r
-\r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::blur(loadMat(img_gray, useRoi), dst, ksize, anchor);\r
-\r
-    cv::Mat dst_gold;\r
-    cv::blur(img_gray, dst_gold, ksize, anchor);\r
-\r
-    EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0);\r
-}\r
-\r
-TEST_P(Blur, Color)\r
+TEST_P(Blur, Accuracy)\r
 {\r
-    cv::Mat img_rgba;\r
-    cv::cvtColor(img, img_rgba, CV_BGR2BGRA);\r
+    cv::Mat src = randomMat(size, type);\r
 \r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::blur(loadMat(img_rgba, useRoi), dst, ksize, anchor);\r
+    cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
+    cv::gpu::blur(loadMat(src, useRoi), dst, ksize, anchor);\r
 \r
     cv::Mat dst_gold;\r
-    cv::blur(img_rgba, dst_gold, ksize, anchor);\r
+    cv::blur(src, dst_gold, ksize, anchor);\r
 \r
     EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 1.0);\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Filter, Blur, testing::Combine(\r
     ALL_DEVICES,\r
+    DIFFERENT_SIZES,\r
+    testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)),\r
     testing::Values(KSize(3, 3), KSize(5, 5), KSize(7, 7)),\r
     testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))),\r
     WHOLE_SUBMAT));\r
@@ -131,69 +117,52 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Blur, testing::Combine(
 IMPLEMENT_PARAM_CLASS(Deriv_X, int)\r
 IMPLEMENT_PARAM_CLASS(Deriv_Y, int)\r
 \r
-PARAM_TEST_CASE(Sobel, cv::gpu::DeviceInfo, KSize, Deriv_X, Deriv_Y, BorderType, UseRoi)\r
+PARAM_TEST_CASE(Sobel, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, Deriv_X, Deriv_Y, BorderType, UseRoi)\r
 {\r
     cv::gpu::DeviceInfo devInfo;\r
+    cv::Size size;\r
+    int type;\r
     cv::Size ksize;\r
     int dx;\r
     int dy;\r
     int borderType;\r
     bool useRoi;\r
 \r
-    cv::Mat img;\r
-\r
     virtual void SetUp()\r
     {\r
         devInfo = GET_PARAM(0);\r
-        ksize = GET_PARAM(1);\r
-        dx = GET_PARAM(2);\r
-        dy = GET_PARAM(3);\r
-        borderType = GET_PARAM(4);\r
-        useRoi = GET_PARAM(5);\r
+        size = GET_PARAM(1);\r
+        type = GET_PARAM(2);\r
+        ksize = GET_PARAM(3);\r
+        dx = GET_PARAM(4);\r
+        dy = GET_PARAM(5);\r
+        borderType = GET_PARAM(6);\r
+        useRoi = GET_PARAM(7);\r
 \r
         cv::gpu::setDevice(devInfo.deviceID());\r
-\r
-        img = readImage("stereobp/aloe-L.png");\r
-        ASSERT_FALSE(img.empty());\r
     }\r
 };\r
 \r
-TEST_P(Sobel, Gray)\r
-{\r
-    if (dx == 0 && dy == 0)\r
-        return;\r
-\r
-    cv::Mat img_gray;\r
-    cv::cvtColor(img, img_gray, CV_BGR2GRAY);\r
-\r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::Sobel(loadMat(img_gray, useRoi), dst, -1, dx, dy, ksize.width, 1.0, borderType);\r
-\r
-    cv::Mat dst_gold;\r
-    cv::Sobel(img_gray, dst_gold, -1, dx, dy, ksize.width, 1.0, 0.0, borderType);\r
-\r
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
-}\r
-\r
-TEST_P(Sobel, Color)\r
+TEST_P(Sobel, Accuracy)\r
 {\r
     if (dx == 0 && dy == 0)\r
         return;\r
 \r
-    cv::Mat img_rgba;\r
-    cv::cvtColor(img, img_rgba, CV_BGR2BGRA);\r
+    cv::Mat src = randomMat(size, type);\r
 \r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::Sobel(loadMat(img_rgba, useRoi), dst, -1, dx, dy, ksize.width, 1.0, borderType);\r
+    cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
+    cv::gpu::Sobel(loadMat(src, useRoi), dst, -1, dx, dy, ksize.width, 1.0, borderType);\r
 \r
     cv::Mat dst_gold;\r
-    cv::Sobel(img_rgba, dst_gold, -1, dx, dy, ksize.width, 1.0, 0.0, borderType);\r
+    cv::Sobel(src, dst_gold, -1, dx, dy, ksize.width, 1.0, 0.0, borderType);\r
 \r
     EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Filter, Sobel, testing::Combine(\r
     ALL_DEVICES,\r
+    DIFFERENT_SIZES,\r
+    testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)),\r
     testing::Values(KSize(3, 3), KSize(5, 5), KSize(7, 7)),\r
     testing::Values(Deriv_X(0), Deriv_X(1), Deriv_X(2)),\r
     testing::Values(Deriv_Y(0), Deriv_Y(1), Deriv_Y(2)),\r
@@ -206,67 +175,50 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Sobel, testing::Combine(
 /////////////////////////////////////////////////////////////////////////////////////////////////\r
 // Scharr\r
 \r
-PARAM_TEST_CASE(Scharr, cv::gpu::DeviceInfo, Deriv_X, Deriv_Y, BorderType, UseRoi)\r
+PARAM_TEST_CASE(Scharr, cv::gpu::DeviceInfo, cv::Size, MatType, Deriv_X, Deriv_Y, BorderType, UseRoi)\r
 {\r
     cv::gpu::DeviceInfo devInfo;\r
+    cv::Size size;\r
+    int type;\r
     int dx;\r
     int dy;\r
     int borderType;\r
     bool useRoi;\r
 \r
-    cv::Mat img;\r
-\r
     virtual void SetUp()\r
     {\r
         devInfo = GET_PARAM(0);\r
-        dx = GET_PARAM(1);\r
-        dy = GET_PARAM(2);\r
-        borderType = GET_PARAM(3);\r
-        useRoi = GET_PARAM(4);\r
+        size = GET_PARAM(1);\r
+        type = GET_PARAM(2);\r
+        dx = GET_PARAM(3);\r
+        dy = GET_PARAM(4);\r
+        borderType = GET_PARAM(5);\r
+        useRoi = GET_PARAM(6);\r
 \r
         cv::gpu::setDevice(devInfo.deviceID());\r
-\r
-        img = readImage("stereobp/aloe-L.png");\r
-        ASSERT_FALSE(img.empty());\r
     }\r
 };\r
 \r
-TEST_P(Scharr, Gray)\r
-{\r
-    if (dx + dy != 1)\r
-        return;\r
-\r
-    cv::Mat img_gray;\r
-    cv::cvtColor(img, img_gray, CV_BGR2GRAY);\r
-\r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::Scharr(loadMat(img_gray, useRoi), dst, -1, dx, dy, 1.0, borderType);\r
-\r
-    cv::Mat dst_gold;\r
-    cv::Scharr(img_gray, dst_gold, -1, dx, dy, 1.0, 0.0, borderType);\r
-\r
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
-}\r
-\r
-TEST_P(Scharr, Color)\r
+TEST_P(Scharr, Accuracy)\r
 {\r
     if (dx + dy != 1)\r
         return;\r
 \r
-    cv::Mat img_rgba;\r
-    cv::cvtColor(img, img_rgba, CV_BGR2BGRA);\r
+    cv::Mat src = randomMat(size, type);\r
 \r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::Scharr(loadMat(img_rgba, useRoi), dst, -1, dx, dy, 1.0, borderType);\r
+    cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
+    cv::gpu::Scharr(loadMat(src, useRoi), dst, -1, dx, dy, 1.0, borderType);\r
 \r
     cv::Mat dst_gold;\r
-    cv::Scharr(img_rgba, dst_gold, -1, dx, dy, 1.0, 0.0, borderType);\r
+    cv::Scharr(src, dst_gold, -1, dx, dy, 1.0, 0.0, borderType);\r
 \r
     EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Filter, Scharr, testing::Combine(\r
     ALL_DEVICES,\r
+    DIFFERENT_SIZES,\r
+    testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)),\r
     testing::Values(Deriv_X(0), Deriv_X(1)),\r
     testing::Values(Deriv_Y(0), Deriv_Y(1)),\r
     testing::Values(BorderType(cv::BORDER_REFLECT101),\r
@@ -278,63 +230,47 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Scharr, testing::Combine(
 /////////////////////////////////////////////////////////////////////////////////////////////////\r
 // GaussianBlur\r
 \r
-PARAM_TEST_CASE(GaussianBlur, cv::gpu::DeviceInfo, KSize, BorderType, UseRoi)\r
+PARAM_TEST_CASE(GaussianBlur, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, BorderType, UseRoi)\r
 {\r
     cv::gpu::DeviceInfo devInfo;\r
+    cv::Size size;\r
+    int type;\r
     cv::Size ksize;\r
     int borderType;\r
     bool useRoi;\r
 \r
-    cv::Mat img;\r
-    double sigma1, sigma2;\r
-\r
     virtual void SetUp()\r
     {\r
         devInfo = GET_PARAM(0);\r
-        ksize = GET_PARAM(1);\r
-        borderType = GET_PARAM(2);\r
-        useRoi = GET_PARAM(3);\r
+        size = GET_PARAM(1);\r
+        type = GET_PARAM(2);\r
+        ksize = GET_PARAM(3);\r
+        borderType = GET_PARAM(4);\r
+        useRoi = GET_PARAM(5);\r
 \r
         cv::gpu::setDevice(devInfo.deviceID());\r
-\r
-        img = readImage("stereobp/aloe-L.png");\r
-        ASSERT_FALSE(img.empty());\r
-\r
-        sigma1 = randomDouble(0.1, 1.0);\r
-        sigma2 = randomDouble(0.1, 1.0);\r
     }\r
 };\r
 \r
-TEST_P(GaussianBlur, Gray)\r
-{\r
-    cv::Mat img_gray;\r
-    cv::cvtColor(img, img_gray, CV_BGR2GRAY);\r
-\r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::GaussianBlur(loadMat(img_gray, useRoi), dst, ksize, sigma1, sigma2, borderType);\r
-\r
-    cv::Mat dst_gold;\r
-    cv::GaussianBlur(img_gray, dst_gold, ksize, sigma1, sigma2, borderType);\r
-\r
-    EXPECT_MAT_NEAR(dst_gold, dst, 4.0);\r
-}\r
-\r
-TEST_P(GaussianBlur, Color)\r
+TEST_P(GaussianBlur, Accuracy)\r
 {\r
-    cv::Mat img_rgba;\r
-    cv::cvtColor(img, img_rgba, CV_BGR2BGRA);\r
+    cv::Mat src = randomMat(size, type);\r
+    double sigma1 = randomDouble(0.1, 1.0);\r
+    double sigma2 = randomDouble(0.1, 1.0);\r
 \r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::GaussianBlur(loadMat(img_rgba, useRoi), dst, ksize, sigma1, sigma2, borderType);\r
+    cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
+    cv::gpu::GaussianBlur(loadMat(src, useRoi), dst, ksize, sigma1, sigma2, borderType);\r
 \r
     cv::Mat dst_gold;\r
-    cv::GaussianBlur(img_rgba, dst_gold, ksize, sigma1, sigma2, borderType);\r
+    cv::GaussianBlur(src, dst_gold, ksize, sigma1, sigma2, borderType);\r
 \r
     EXPECT_MAT_NEAR(dst_gold, dst, 4.0);\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Filter, GaussianBlur, testing::Combine(\r
     ALL_DEVICES,\r
+    DIFFERENT_SIZES,\r
+    testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)),\r
     testing::Values(KSize(3, 3),\r
                     KSize(5, 5),\r
                     KSize(7, 7),\r
@@ -359,72 +295,46 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, GaussianBlur, testing::Combine(
 /////////////////////////////////////////////////////////////////////////////////////////////////\r
 // Laplacian\r
 \r
-PARAM_TEST_CASE(Laplacian, cv::gpu::DeviceInfo, KSize, UseRoi)\r
+PARAM_TEST_CASE(Laplacian, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, UseRoi)\r
 {\r
     cv::gpu::DeviceInfo devInfo;\r
+    cv::Size size;\r
+    int type;\r
     cv::Size ksize;\r
     bool useRoi;\r
 \r
-    cv::Mat img;\r
-\r
     virtual void SetUp()\r
     {\r
         devInfo = GET_PARAM(0);\r
-        ksize = GET_PARAM(1);\r
-        useRoi = GET_PARAM(2);\r
+        size = GET_PARAM(1);\r
+        type = GET_PARAM(2);\r
+        ksize = GET_PARAM(3);\r
+        useRoi = GET_PARAM(4);\r
 \r
         cv::gpu::setDevice(devInfo.deviceID());\r
-\r
-        img = readImage("stereobp/aloe-L.png");\r
-        ASSERT_FALSE(img.empty());\r
     }\r
 };\r
 \r
-TEST_P(Laplacian, Gray)\r
-{\r
-    cv::Mat img_gray;\r
-    cv::cvtColor(img, img_gray, CV_BGR2GRAY);\r
-\r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::Laplacian(loadMat(img_gray, useRoi), dst, -1, ksize.width);\r
-\r
-    cv::Mat dst_gold;\r
-    cv::Laplacian(img_gray, dst_gold, -1, ksize.width);\r
-\r
-    EXPECT_MAT_NEAR(getInnerROI(dst_gold, cv::Size(3, 3)), getInnerROI(dst, cv::Size(3, 3)), 0.0);\r
-}\r
-\r
-TEST_P(Laplacian, Color)\r
+TEST_P(Laplacian, Accuracy)\r
 {\r
-    cv::Mat img_rgba;\r
-    cv::cvtColor(img, img_rgba, CV_BGR2BGRA);\r
+    cv::Mat src = randomMat(size, type);\r
 \r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::Laplacian(loadMat(img_rgba, useRoi), dst, -1, ksize.width);\r
-\r
-    cv::Mat dst_gold;\r
-    cv::Laplacian(img_rgba, dst_gold, -1, ksize.width);\r
-\r
-    EXPECT_MAT_NEAR(getInnerROI(dst_gold, cv::Size(3, 3)), getInnerROI(dst, cv::Size(3, 3)), 0.0);\r
-}\r
-\r
-TEST_P(Laplacian, Gray_32FC1)\r
-{\r
-    cv::Mat src;\r
-    cv::cvtColor(img, src, CV_BGR2GRAY);\r
-    src.convertTo(src, CV_32F, 1.0 / 255.0);\r
-\r
-    cv::gpu::GpuMat dst;\r
+    cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
     cv::gpu::Laplacian(loadMat(src, useRoi), dst, -1, ksize.width);\r
 \r
     cv::Mat dst_gold;\r
     cv::Laplacian(src, dst_gold, -1, ksize.width);\r
 \r
-    EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+    if (type == CV_32FC1)\r
+        EXPECT_MAT_NEAR(dst_gold, dst, 0.0);\r
+    else\r
+        EXPECT_MAT_NEAR(getInnerROI(dst_gold, cv::Size(3, 3)), getInnerROI(dst, cv::Size(3, 3)), 0.0);\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Filter, Laplacian, testing::Combine(\r
     ALL_DEVICES,\r
+    DIFFERENT_SIZES,\r
+    testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)),\r
     testing::Values(KSize(1, 1), KSize(3, 3)),\r
     WHOLE_SUBMAT));\r
 \r
@@ -433,58 +343,38 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Laplacian, testing::Combine(
 \r
 IMPLEMENT_PARAM_CLASS(Iterations, int)\r
 \r
-PARAM_TEST_CASE(Erode, cv::gpu::DeviceInfo, Anchor, Iterations, UseRoi)\r
+PARAM_TEST_CASE(Erode, cv::gpu::DeviceInfo, cv::Size, MatType, Anchor, Iterations, UseRoi)\r
 {\r
     cv::gpu::DeviceInfo devInfo;\r
+    cv::Size size;\r
+    int type;\r
     cv::Point anchor;\r
     int iterations;\r
     bool useRoi;\r
 \r
-    cv::Mat img;\r
-    cv::Mat kernel;\r
-\r
     virtual void SetUp()\r
     {\r
         devInfo = GET_PARAM(0);\r
-        anchor = GET_PARAM(1);\r
-        iterations = GET_PARAM(2);\r
-        useRoi = GET_PARAM(3);\r
+        size = GET_PARAM(1);\r
+        type = GET_PARAM(2);\r
+        anchor = GET_PARAM(3);\r
+        iterations = GET_PARAM(4);\r
+        useRoi = GET_PARAM(5);\r
 \r
         cv::gpu::setDevice(devInfo.deviceID());\r
-\r
-        img = readImage("stereobp/aloe-L.png");\r
-        ASSERT_FALSE(img.empty());\r
-\r
-        kernel = cv::Mat::ones(3, 3, CV_8U);\r
     }\r
 };\r
 \r
-TEST_P(Erode, Gray)\r
-{\r
-    cv::Mat img_gray;\r
-    cv::cvtColor(img, img_gray, CV_BGR2GRAY);\r
-\r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::erode(loadMat(img_gray, useRoi), dst, kernel, anchor, iterations);\r
-\r
-    cv::Mat dst_gold;\r
-    cv::erode(img_gray, dst_gold, kernel, anchor, iterations);\r
-\r
-    cv::Size ksize = cv::Size(kernel.cols + iterations * (kernel.cols - 1), kernel.rows + iterations * (kernel.rows - 1));\r
-\r
-    EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0);\r
-}\r
-\r
-TEST_P(Erode, Color)\r
+TEST_P(Erode, Accuracy)\r
 {\r
-    cv::Mat img_rgba;\r
-    cv::cvtColor(img, img_rgba, CV_BGR2BGRA);\r
+    cv::Mat src = randomMat(size, type);\r
+    cv::Mat kernel = cv::Mat::ones(3, 3, CV_8U);\r
 \r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::erode(loadMat(img_rgba, useRoi), dst, kernel, anchor, iterations);\r
+    cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
+    cv::gpu::erode(loadMat(src, useRoi), dst, kernel, anchor, iterations);\r
 \r
     cv::Mat dst_gold;\r
-    cv::erode(img_rgba, dst_gold, kernel, anchor, iterations);\r
+    cv::erode(src, dst_gold, kernel, anchor, iterations);\r
 \r
     cv::Size ksize = cv::Size(kernel.cols + iterations * (kernel.cols - 1), kernel.rows + iterations * (kernel.rows - 1));\r
 \r
@@ -493,6 +383,8 @@ TEST_P(Erode, Color)
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Filter, Erode, testing::Combine(\r
     ALL_DEVICES,\r
+    DIFFERENT_SIZES,\r
+    testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)),\r
     testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))),\r
     testing::Values(Iterations(1), Iterations(2), Iterations(3)),\r
     WHOLE_SUBMAT));\r
@@ -500,58 +392,38 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Erode, testing::Combine(
 /////////////////////////////////////////////////////////////////////////////////////////////////\r
 // Dilate\r
 \r
-PARAM_TEST_CASE(Dilate, cv::gpu::DeviceInfo, Anchor, Iterations, UseRoi)\r
+PARAM_TEST_CASE(Dilate, cv::gpu::DeviceInfo, cv::Size, MatType, Anchor, Iterations, UseRoi)\r
 {\r
     cv::gpu::DeviceInfo devInfo;\r
+    cv::Size size;\r
+    int type;\r
     cv::Point anchor;\r
     int iterations;\r
     bool useRoi;\r
 \r
-    cv::Mat img;\r
-    cv::Mat kernel;\r
-\r
     virtual void SetUp()\r
     {\r
         devInfo = GET_PARAM(0);\r
-        anchor = GET_PARAM(1);\r
-        iterations = GET_PARAM(2);\r
-        useRoi = GET_PARAM(3);\r
+        size = GET_PARAM(1);\r
+        type = GET_PARAM(2);\r
+        anchor = GET_PARAM(3);\r
+        iterations = GET_PARAM(4);\r
+        useRoi = GET_PARAM(5);\r
 \r
         cv::gpu::setDevice(devInfo.deviceID());\r
-\r
-        img = readImage("stereobp/aloe-L.png");\r
-        ASSERT_FALSE(img.empty());\r
-\r
-        kernel = cv::Mat::ones(3, 3, CV_8U);\r
     }\r
 };\r
 \r
-TEST_P(Dilate, Gray)\r
-{\r
-    cv::Mat img_gray;\r
-    cv::cvtColor(img, img_gray, CV_BGR2GRAY);\r
-\r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::dilate(loadMat(img_gray, useRoi), dst, kernel, anchor, iterations);\r
-\r
-    cv::Mat dst_gold;\r
-    cv::dilate(img_gray, dst_gold, kernel, anchor, iterations);\r
-\r
-    cv::Size ksize = cv::Size(kernel.cols + iterations * (kernel.cols - 1), kernel.rows + iterations * (kernel.rows - 1));\r
-\r
-    EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0);\r
-}\r
-\r
-TEST_P(Dilate, Color)\r
+TEST_P(Dilate, Accuracy)\r
 {\r
-    cv::Mat img_rgba;\r
-    cv::cvtColor(img, img_rgba, CV_BGR2BGRA);\r
+    cv::Mat src = randomMat(size, type);\r
+    cv::Mat kernel = cv::Mat::ones(3, 3, CV_8U);\r
 \r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::dilate(loadMat(img_rgba, useRoi), dst, kernel, anchor, iterations);\r
+    cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
+    cv::gpu::dilate(loadMat(src, useRoi), dst, kernel, anchor, iterations);\r
 \r
     cv::Mat dst_gold;\r
-    cv::dilate(img_rgba, dst_gold, kernel, anchor, iterations);\r
+    cv::dilate(src, dst_gold, kernel, anchor, iterations);\r
 \r
     cv::Size ksize = cv::Size(kernel.cols + iterations * (kernel.cols - 1), kernel.rows + iterations * (kernel.rows - 1));\r
 \r
@@ -560,6 +432,8 @@ TEST_P(Dilate, Color)
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Filter, Dilate, testing::Combine(\r
     ALL_DEVICES,\r
+    DIFFERENT_SIZES,\r
+    testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)),\r
     testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))),\r
     testing::Values(Iterations(1), Iterations(2), Iterations(3)),\r
     WHOLE_SUBMAT));\r
@@ -570,60 +444,40 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Dilate, testing::Combine(
 CV_ENUM(MorphOp, cv::MORPH_OPEN, cv::MORPH_CLOSE, cv::MORPH_GRADIENT, cv::MORPH_TOPHAT, cv::MORPH_BLACKHAT)\r
 #define ALL_MORPH_OPS testing::Values(MorphOp(cv::MORPH_OPEN), MorphOp(cv::MORPH_CLOSE), MorphOp(cv::MORPH_GRADIENT), MorphOp(cv::MORPH_TOPHAT), MorphOp(cv::MORPH_BLACKHAT))\r
 \r
-PARAM_TEST_CASE(MorphEx, cv::gpu::DeviceInfo, MorphOp, Anchor, Iterations, UseRoi)\r
+PARAM_TEST_CASE(MorphEx, cv::gpu::DeviceInfo, cv::Size, MatType, MorphOp, Anchor, Iterations, UseRoi)\r
 {\r
     cv::gpu::DeviceInfo devInfo;\r
+    cv::Size size;\r
+    int type;\r
     int morphOp;\r
     cv::Point anchor;\r
     int iterations;\r
     bool useRoi;\r
 \r
-    cv::Mat img;\r
-    cv::Mat kernel;\r
-\r
     virtual void SetUp()\r
     {\r
         devInfo = GET_PARAM(0);\r
-        morphOp = GET_PARAM(1);\r
-        anchor = GET_PARAM(2);\r
-        iterations = GET_PARAM(3);\r
-        useRoi = GET_PARAM(4);\r
+        size = GET_PARAM(1);\r
+        type = GET_PARAM(2);\r
+        morphOp = GET_PARAM(3);\r
+        anchor = GET_PARAM(4);\r
+        iterations = GET_PARAM(5);\r
+        useRoi = GET_PARAM(6);\r
 \r
         cv::gpu::setDevice(devInfo.deviceID());\r
-\r
-        img = readImage("stereobp/aloe-L.png");\r
-        ASSERT_FALSE(img.empty());\r
-\r
-        kernel = cv::Mat::ones(3, 3, CV_8U);\r
     }\r
 };\r
 \r
-TEST_P(MorphEx, Gray)\r
-{\r
-    cv::Mat img_gray;\r
-    cv::cvtColor(img, img_gray, CV_BGR2GRAY);\r
-\r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::morphologyEx(loadMat(img_gray, useRoi), dst, morphOp, kernel, anchor, iterations);\r
-\r
-    cv::Mat dst_gold;\r
-    cv::morphologyEx(img_gray, dst_gold, morphOp, kernel, anchor, iterations);\r
-\r
-    cv::Size border = cv::Size(kernel.cols + (iterations + 1) * kernel.cols + 2, kernel.rows + (iterations + 1) * kernel.rows + 2);\r
-\r
-    EXPECT_MAT_NEAR(getInnerROI(dst_gold, border), getInnerROI(dst, border), 0.0);\r
-}\r
-\r
-TEST_P(MorphEx, Color)\r
+TEST_P(MorphEx, Accuracy)\r
 {\r
-    cv::Mat img_rgba;\r
-    cv::cvtColor(img, img_rgba, CV_BGR2BGRA);\r
+    cv::Mat src = randomMat(size, type);\r
+    cv::Mat kernel = cv::Mat::ones(3, 3, CV_8U);\r
 \r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::morphologyEx(loadMat(img_rgba, useRoi), dst, morphOp, kernel, anchor, iterations);\r
+    cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
+    cv::gpu::morphologyEx(loadMat(src, useRoi), dst, morphOp, kernel, anchor, iterations);\r
 \r
     cv::Mat dst_gold;\r
-    cv::morphologyEx(img_rgba, dst_gold, morphOp, kernel, anchor, iterations);\r
+    cv::morphologyEx(src, dst_gold, morphOp, kernel, anchor, iterations);\r
 \r
     cv::Size border = cv::Size(kernel.cols + (iterations + 1) * kernel.cols + 2, kernel.rows + (iterations + 1) * kernel.rows + 2);\r
 \r
@@ -632,6 +486,8 @@ TEST_P(MorphEx, Color)
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Filter, MorphEx, testing::Combine(\r
     ALL_DEVICES,\r
+    DIFFERENT_SIZES,\r
+    testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)),\r
     ALL_MORPH_OPS,\r
     testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))),\r
     testing::Values(Iterations(1), Iterations(2), Iterations(3)),\r
@@ -640,9 +496,11 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, MorphEx, testing::Combine(
 /////////////////////////////////////////////////////////////////////////////////////////////////\r
 // Filter2D\r
 \r
-PARAM_TEST_CASE(Filter2D, cv::gpu::DeviceInfo, KSize, Anchor, UseRoi)\r
+PARAM_TEST_CASE(Filter2D, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, Anchor, UseRoi)\r
 {\r
     cv::gpu::DeviceInfo devInfo;\r
+    cv::Size size;\r
+    int type;\r
     cv::Size ksize;\r
     cv::Point anchor;\r
     bool useRoi;\r
@@ -653,64 +511,37 @@ PARAM_TEST_CASE(Filter2D, cv::gpu::DeviceInfo, KSize, Anchor, UseRoi)
     virtual void SetUp()\r
     {\r
         devInfo = GET_PARAM(0);\r
-        ksize = GET_PARAM(1);\r
-        anchor = GET_PARAM(2);\r
-        useRoi = GET_PARAM(3);\r
+        size = GET_PARAM(1);\r
+        type = GET_PARAM(2);\r
+        ksize = GET_PARAM(3);\r
+        anchor = GET_PARAM(4);\r
+        useRoi = GET_PARAM(5);\r
 \r
         cv::gpu::setDevice(devInfo.deviceID());\r
-\r
-        img = readImage("stereobp/aloe-L.png");\r
-        ASSERT_FALSE(img.empty());\r
-\r
-        kernel = cv::Mat::ones(ksize.height, ksize.width, CV_32FC1);\r
     }\r
 };\r
 \r
-TEST_P(Filter2D, Gray)\r
-{\r
-    cv::Mat img_gray;\r
-    cv::cvtColor(img, img_gray, CV_BGR2GRAY);\r
-\r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::filter2D(loadMat(img_gray, useRoi), dst, -1, kernel, anchor);\r
-\r
-    cv::Mat dst_gold;\r
-    cv::filter2D(img_gray, dst_gold, -1, kernel, anchor, 0, cv::BORDER_CONSTANT);\r
-\r
-    EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0);\r
-}\r
-\r
-TEST_P(Filter2D, Color)\r
-{\r
-    cv::Mat img_rgba;\r
-    cv::cvtColor(img, img_rgba, CV_BGR2BGRA);\r
-\r
-    cv::gpu::GpuMat dst;\r
-    cv::gpu::filter2D(loadMat(img_rgba, useRoi), dst, -1, kernel, anchor);\r
-\r
-    cv::Mat dst_gold;\r
-    cv::filter2D(img_rgba, dst_gold, -1, kernel, anchor, 0, cv::BORDER_CONSTANT);\r
-\r
-    EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0);\r
-}\r
-\r
-TEST_P(Filter2D, Gray_32FC1)\r
+TEST_P(Filter2D, Accuracy)\r
 {\r
-    cv::Mat src;\r
-    cv::cvtColor(img, src, CV_BGR2GRAY);\r
-    src.convertTo(src, CV_32F, 1.0 / 255.0);\r
+    cv::Mat src = randomMat(size, type);\r
+    cv::Mat kernel = cv::Mat::ones(ksize.height, ksize.width, CV_32FC1);\r
 \r
-    cv::gpu::GpuMat dst;\r
+    cv::gpu::GpuMat dst = createMat(size, type, useRoi);\r
     cv::gpu::filter2D(loadMat(src, useRoi), dst, -1, kernel, anchor);\r
 \r
     cv::Mat dst_gold;\r
-    cv::filter2D(src, dst_gold, -1, kernel, anchor);\r
+    cv::filter2D(src, dst_gold, -1, kernel, anchor, 0, type == CV_32FC1 ? cv::BORDER_DEFAULT : cv::BORDER_CONSTANT);\r
 \r
-    EXPECT_MAT_NEAR(dst_gold, dst, 1e-3);\r
+    if (type == CV_32FC1)\r
+        EXPECT_MAT_NEAR(dst_gold, dst, 1e-1);\r
+    else\r
+        EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0);\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(GPU_Filter, Filter2D, testing::Combine(\r
     ALL_DEVICES,\r
+    DIFFERENT_SIZES,\r
+    testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)),\r
     testing::Values(KSize(3, 3), KSize(5, 5), KSize(7, 7), KSize(11, 11), KSize(13, 13), KSize(15, 15)),\r
     testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))),\r
     WHOLE_SUBMAT));\r