removed columnSum function (it is a duplicate for reduce)

modules/gpu/include/opencv2/gpu.hpp

@@ -184,9 +184,6 @@ CV_EXPORTS void integralBuffered(const GpuMat& src, GpuMat& sum, GpuMat& buffer,
 //! 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

modules/gpu/perf/perf_imgproc.cpp

@@ -631,32 +631,6 @@ PERF_TEST_P(Sz_ClipLimit, ImgProc_CLAHE,
     }
 }
 
-//////////////////////////////////////////////////////////////////////
-// 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
 

modules/gpu/src/cuda/imgproc.cu

@@ -582,41 +582,6 @@ namespace cv { namespace gpu { namespace cudev
                 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
 

modules/gpu/src/imgproc.cpp

@@ -59,7 +59,6 @@ void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int,
 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(); }
@@ -630,26 +629,7 @@ void cv::gpu::sqrIntegral(const GpuMat& src, GpuMat& sqsum, Stream& s)
 }
 
 //////////////////////////////////////////////////////////////////////////////
-// 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)
 {

modules/gpu/test/test_imgproc.cpp

@@ -261,54 +261,6 @@ INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CLAHE, testing::Combine(
     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
 

modules/gpuarithm/src/matrix_reductions.cpp

@@ -71,7 +71,6 @@ int cv::gpu::countNonZero(const GpuMat&, GpuMat&) { throw_no_cuda(); return 0; }
 void cv::gpu::reduce(const GpuMat&, GpuMat&, int, int, int, Stream&) { throw_no_cuda(); }
 
 #else
-#include "opencv2/core/utility.hpp"
 
 namespace
 {