\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
\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
}\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
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
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
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