implemented brute force convolve for small kernel sizes

diff --git a/modules/gpu/perf/perf_imgproc.cpp b/modules/gpu/perf/perf_imgproc.cpp
index 81dd559..3f72641 100644 (file)
--- a/modules/gpu/perf/perf_imgproc.cpp
+++ b/modules/gpu/perf/perf_imgproc.cpp
@@ -737,7 +737,7 @@ PERF_TEST_P(DevInfo_Size, dft, testing::Combine(testing::ValuesIn(devices()),
 \r
 PERF_TEST_P(DevInfo_Int_Int, convolve, testing::Combine(testing::ValuesIn(devices()),\r
                                                      testing::Values(512, 1024, 1536, 2048, 2560, 3072, 3584),\r
-                                                     testing::Values(27, 32, 64)))\r
+                                                     testing::Values(3, 9, 27, 32, 64)))\r
 {\r
     DeviceInfo devInfo = std::tr1::get<0>(GetParam());\r
     int image_size = std::tr1::get<1>(GetParam());\r
@@ -745,13 +745,12 @@ PERF_TEST_P(DevInfo_Int_Int, convolve, testing::Combine(testing::ValuesIn(device
 \r
     setDevice(devInfo.deviceID());\r
 \r
-    Mat image_host(image_size, image_size, CV_32FC1);\r
-    Mat templ_host(templ_size, templ_size, CV_32FC1);\r
+    GpuMat image = createContinuous(image_size, image_size, CV_32FC1);\r
+    GpuMat templ = createContinuous(templ_size, templ_size, CV_32FC1);\r
 \r
-    declare.in(image_host, templ_host, WARMUP_RNG);\r
+    image.setTo(Scalar(1.0));\r
+    templ.setTo(Scalar(1.0));\r
 \r
-    GpuMat image(image_host);\r
-    GpuMat templ(templ_host);\r
     GpuMat dst;\r
     ConvolveBuf buf;\r
 \r

diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu
index 79b5e10..d2ce30e 100644 (file)
--- a/modules/gpu/src/cuda/imgproc.cu
+++ b/modules/gpu/src/cuda/imgproc.cu
@@ -951,6 +951,84 @@ namespace cv { namespace gpu { namespace imgproc
     }\r
 \r
 \r
+    //////////////////////////////////////////////////////////////////////////\r
+    // convolve\r
+\r
+    #define CONVOLVE_MAX_KERNEL_SIZE 17\r
+\r
+    __constant__ float c_convolveKernel[CONVOLVE_MAX_KERNEL_SIZE * CONVOLVE_MAX_KERNEL_SIZE];\r
+\r
+    __global__ void convolve(const DevMem2Df src, PtrStepf dst, int kWidth, int kHeight)\r
+    {\r
+        __shared__ float smem[16 + 2 * 8][16 + 2 * 8];\r
+\r
+        const int x = blockIdx.x * blockDim.x + threadIdx.x;\r
+        const int y = blockIdx.y * blockDim.y + threadIdx.y;\r
+\r
+        // x | x 0 | 0\r
+        // -----------\r
+        // x | x 0 | 0\r
+        // 0 | 0 0 | 0\r
+        // -----------\r
+        // 0 | 0 0 | 0\r
+        smem[threadIdx.y][threadIdx.x] = src.ptr(min(max(y - 8, 0), src.rows - 1))[min(max(x - 8, 0), src.cols - 1)];\r
+\r
+        // 0 | 0 x | x\r
+        // -----------\r
+        // 0 | 0 x | x\r
+        // 0 | 0 0 | 0\r
+        // -----------\r
+        // 0 | 0 0 | 0\r
+        smem[threadIdx.y][threadIdx.x + 16] = src.ptr(min(max(y - 8, 0), src.rows - 1))[min(x + 8, src.cols - 1)];\r
+\r
+        // 0 | 0 0 | 0\r
+        // -----------\r
+        // 0 | 0 0 | 0\r
+        // x | x 0 | 0\r
+        // -----------\r
+        // x | x 0 | 0\r
+        smem[threadIdx.y + 16][threadIdx.x] = src.ptr(min(y + 8, src.rows - 1))[min(max(x - 8, 0), src.cols - 1)];\r
+\r
+        // 0 | 0 0 | 0\r
+        // -----------\r
+        // 0 | 0 0 | 0\r
+        // 0 | 0 x | x\r
+        // -----------\r
+        // 0 | 0 x | x\r
+        smem[threadIdx.y + 16][threadIdx.x + 16] = src.ptr(min(y + 8, src.rows - 1))[min(x + 8, src.cols - 1)];\r
+\r
+        __syncthreads();\r
+\r
+        if (x < src.cols && y < src.rows)\r
+        {\r
+            float res = 0;\r
+\r
+            for (int i = 0; i < kHeight; ++i)\r
+            {\r
+                for (int j = 0; j < kWidth; ++j)\r
+                {\r
+                    res += smem[threadIdx.y + 8 - kHeight / 2 + i][threadIdx.x + 8 - kWidth / 2 + j] * c_convolveKernel[i * kWidth + j];\r
+                }\r
+            }\r
+\r
+            dst.ptr(y)[x] = res;\r
+        }\r
+    }\r
+\r
+    void convolve_gpu(const DevMem2Df& src, const PtrStepf& dst, int kWidth, int kHeight, float* kernel)\r
+    {\r
+        cudaSafeCall(cudaMemcpyToSymbol(c_convolveKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice) );\r
+\r
+        const dim3 block(16, 16);\r
+        const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));\r
+\r
+        convolve<<<grid, block>>>(src, dst, kWidth, kHeight);\r
+        cudaSafeCall(cudaGetLastError());\r
+\r
+        cudaSafeCall(cudaDeviceSynchronize());\r
+    }\r
+\r
+\r
 }}}\r
 \r
 \r

diff --git a/modules/gpu/src/imgproc.cpp b/modules/gpu/src/imgproc.cpp
index 47b0998..35878c4 100644 (file)
--- a/modules/gpu/src/imgproc.cpp
+++ b/modules/gpu/src/imgproc.cpp
@@ -1576,6 +1576,10 @@ void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result,
     convolve(image, templ, result, ccorr, buf);\r
 }\r
 \r
+namespace cv { namespace gpu { namespace imgproc\r
+{\r
+    void convolve_gpu(const DevMem2Df& src, const PtrStepf& dst, int kWidth, int kHeight, float* kernel);\r
+}}}\r
 \r
 void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result, \r
                        bool ccorr, ConvolveBuf& buf)\r
@@ -1586,6 +1590,24 @@ void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result,
     CV_Assert(image.type() == CV_32F);\r
     CV_Assert(templ.type() == CV_32F);\r
 \r
+    if (templ.cols < 13 && templ.rows < 13)\r
+    {\r
+        result.create(image.size(), CV_32F);\r
+        GpuMat contKernel;\r
+\r
+        if (templ.isContinuous())\r
+            contKernel = templ;\r
+        else\r
+        {\r
+            contKernel = createContinuous(templ.size(), templ.type());\r
+            templ.copyTo(contKernel);\r
+        }\r
+\r
+        imgproc::convolve_gpu(image, result, templ.cols, templ.rows, contKernel.ptr<float>());\r
+\r
+        return;\r
+    }\r
+\r
     buf.create(image.size(), templ.size());\r
     result.create(buf.result_size, CV_32F);\r
 \r

diff --git a/modules/gpu/test/test_imgproc.cpp b/modules/gpu/test/test_imgproc.cpp
index 30dbb9b..1ff6cf2 100644 (file)
--- a/modules/gpu/test/test_imgproc.cpp
+++ b/modules/gpu/test/test_imgproc.cpp
@@ -4221,4 +4221,59 @@ INSTANTIATE_TEST_CASE_P(ImgProc, Canny, testing::Combine(
                         testing::Values(3, 5),\r
                         testing::Values(false, true)));\r
 \r
+////////////////////////////////////////////////////////\r
+// convolve\r
+\r
+struct Convolve: testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int> >\r
+{    \r
+    cv::gpu::DeviceInfo devInfo;\r
+    int ksize;\r
+    \r
+    cv::Size size;    \r
+    cv::Mat src;\r
+    cv::Mat kernel;\r
+    \r
+    cv::Mat dst_gold;\r
+\r
+    virtual void SetUp()\r
+    {\r
+        devInfo = std::tr1::get<0>(GetParam());\r
+        ksize = std::tr1::get<1>(GetParam());\r
+\r
+        cv::gpu::setDevice(devInfo.deviceID());\r
+        \r
+        cv::RNG& rng = cvtest::TS::ptr()->get_rng();\r
+\r
+        size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200));\r
+\r
+        src = cvtest::randomMat(rng, size, CV_32FC1, 0.0, 255.0, false);\r
+        kernel = cvtest::randomMat(rng, cv::Size(ksize, ksize), CV_32FC1, 0.0, 1.0, false);\r
+        \r
+        cv::filter2D(src, dst_gold, CV_32F, kernel, cv::Point(-1, -1), 0, cv::BORDER_REPLICATE);\r
+    }\r
+};\r
+\r
+TEST_P(Convolve, Accuracy)\r
+{\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_PARAM(ksize);\r
+    \r
+    cv::Mat dst;\r
+\r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat d_dst;\r
+\r
+        cv::gpu::convolve(cv::gpu::GpuMat(src), cv::gpu::GpuMat(kernel), d_dst);\r
+        \r
+        d_dst.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst, dst_gold, 1e-2);\r
+}\r
+\r
+\r
+INSTANTIATE_TEST_CASE_P(ImgProc, Convolve, testing::Combine(\r
+                        testing::ValuesIn(devices()), \r
+                        testing::Values(3, 5, 7, 9, 11)));\r
+\r
 #endif // HAVE_CUDA\r