implemented brute force convolve for small kernel sizes

[profile/ivi/opencv.git] / modules / gpu / src / cuda / imgproc.cu

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