//! supports source images of 8UC1 type only
CV_EXPORTS void sqrIntegral(const GpuMat& src, GpuMat& sqsum, Stream& stream = Stream::Null());
-//! computes vertical sum, supports only CV_32FC1 images
-CV_EXPORTS void columnSum(const GpuMat& src, GpuMat& sum);
-
//! computes the standard deviation of integral images
//! supports only CV_32SC1 source type and CV_32FC1 sqr type
//! output will have CV_32FC1 type
}
//////////////////////////////////////////////////////////////////////
-// ColumnSum
-
-PERF_TEST_P(Sz, ImgProc_ColumnSum,
- GPU_TYPICAL_MAT_SIZES)
-{
- const cv::Size size = GetParam();
-
- cv::Mat src(size, CV_32FC1);
- declare.in(src, WARMUP_RNG);
-
- if (PERF_RUN_GPU())
- {
- const cv::gpu::GpuMat d_src(src);
- cv::gpu::GpuMat dst;
-
- TEST_CYCLE() cv::gpu::columnSum(d_src, dst);
-
- GPU_SANITY_CHECK(dst);
- }
- else
- {
- FAIL_NO_CPU();
- }
-}
-
-//////////////////////////////////////////////////////////////////////
// Canny
DEF_PARAM_TEST(Image_AppertureSz_L2gradient, string, int, bool);
cudaSafeCall(cudaDeviceSynchronize());
}
- ////////////////////////////// Column Sum //////////////////////////////////////
-
- __global__ void column_sumKernel_32F(int cols, int rows, const PtrStepb src, const PtrStepb dst)
- {
- int x = blockIdx.x * blockDim.x + threadIdx.x;
-
- if (x < cols)
- {
- const unsigned char* src_data = src.data + x * sizeof(float);
- unsigned char* dst_data = dst.data + x * sizeof(float);
-
- float sum = 0.f;
- for (int y = 0; y < rows; ++y)
- {
- sum += *(const float*)src_data;
- *(float*)dst_data = sum;
- src_data += src.step;
- dst_data += dst.step;
- }
- }
- }
-
-
- void columnSum_32F(const PtrStepSzb src, const PtrStepSzb dst)
- {
- dim3 threads(256);
- dim3 grid(divUp(src.cols, threads.x));
-
- column_sumKernel_32F<<<grid, threads>>>(src.cols, src.rows, src, dst);
- cudaSafeCall( cudaGetLastError() );
-
- cudaSafeCall( cudaDeviceSynchronize() );
- }
-
-
//////////////////////////////////////////////////////////////////////////
// mulSpectrums
void cv::gpu::integral(const GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); }
void cv::gpu::integralBuffered(const GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); }
void cv::gpu::sqrIntegral(const GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); }
-void cv::gpu::columnSum(const GpuMat&, GpuMat&) { throw_no_cuda(); }
void cv::gpu::rectStdDev(const GpuMat&, const GpuMat&, GpuMat&, const Rect&, Stream&) { throw_no_cuda(); }
void cv::gpu::evenLevels(GpuMat&, int, int, int) { throw_no_cuda(); }
void cv::gpu::histEven(const GpuMat&, GpuMat&, int, int, int, Stream&) { throw_no_cuda(); }
}
//////////////////////////////////////////////////////////////////////////////
-// columnSum
-
-namespace cv { namespace gpu { namespace cudev
-{
- namespace imgproc
- {
- void columnSum_32F(const PtrStepSzb src, const PtrStepSzb dst);
- }
-}}}
-
-void cv::gpu::columnSum(const GpuMat& src, GpuMat& dst)
-{
- using namespace ::cv::gpu::cudev::imgproc;
-
- CV_Assert(src.type() == CV_32F);
-
- dst.create(src.size(), CV_32F);
-
- cudev::imgproc::columnSum_32F(src, dst);
-}
+// rectStdDev
void cv::gpu::rectStdDev(const GpuMat& src, const GpuMat& sqr, GpuMat& dst, const Rect& rect, Stream& s)
{
DIFFERENT_SIZES,
testing::Values(0.0, 40.0)));
-////////////////////////////////////////////////////////////////////////
-// ColumnSum
-
-PARAM_TEST_CASE(ColumnSum, cv::gpu::DeviceInfo, cv::Size)
-{
- cv::gpu::DeviceInfo devInfo;
- cv::Size size;
-
- virtual void SetUp()
- {
- devInfo = GET_PARAM(0);
- size = GET_PARAM(1);
-
- cv::gpu::setDevice(devInfo.deviceID());
- }
-};
-
-GPU_TEST_P(ColumnSum, Accuracy)
-{
- cv::Mat src = randomMat(size, CV_32FC1);
-
- cv::gpu::GpuMat d_dst;
- cv::gpu::columnSum(loadMat(src), d_dst);
-
- cv::Mat dst(d_dst);
-
- for (int j = 0; j < src.cols; ++j)
- {
- float gold = src.at<float>(0, j);
- float res = dst.at<float>(0, j);
- ASSERT_NEAR(res, gold, 1e-5);
- }
-
- for (int i = 1; i < src.rows; ++i)
- {
- for (int j = 0; j < src.cols; ++j)
- {
- float gold = src.at<float>(i, j) += src.at<float>(i - 1, j);
- float res = dst.at<float>(i, j);
- ASSERT_NEAR(res, gold, 1e-5);
- }
- }
-}
-
-INSTANTIATE_TEST_CASE_P(GPU_ImgProc, ColumnSum, testing::Combine(
- ALL_DEVICES,
- DIFFERENT_SIZES));
-
////////////////////////////////////////////////////////
// Canny
void cv::gpu::reduce(const GpuMat&, GpuMat&, int, int, int, Stream&) { throw_no_cuda(); }
#else
-#include "opencv2/core/utility.hpp"
namespace
{