From 9801d07a4697ea27e7cdb53fb8be7e1146a5eedd Mon Sep 17 00:00:00 2001
From: Alexey Spizhevoy <no@email>
Date: Wed, 8 Dec 2010 16:51:12 +0000
Subject: [PATCH] added test for gpu:columnSum

---
 modules/gpu/src/cuda/imgproc.cu | 12 ++++----
 tests/gpu/src/imgproc_gpu.cpp   | 51 +++++++++++++++++++++++++++++++++
 2 files changed, 58 insertions(+), 5 deletions(-)

diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu
index 94cfb40c7d..143e8d09f1 100644
--- a/modules/gpu/src/cuda/imgproc.cu
+++ b/modules/gpu/src/cuda/imgproc.cu
@@ -723,16 +723,18 @@ namespace cv { namespace gpu { namespace imgproc
     {
         int x = blockIdx.x * blockDim.x + threadIdx.x;
 
-        const float* src_data = (const float*)src.data + x;
-        float* dst_data = (float*)dst.data + 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 += src_data[y];
-                dst_data[y] = sum;
+                sum += *(const float*)src_data;
+                *(float*)dst_data = sum;
+                src_data += src.step;
+                dst_data += dst.step;
             }
         }
     }
diff --git a/tests/gpu/src/imgproc_gpu.cpp b/tests/gpu/src/imgproc_gpu.cpp
index f9fe87876b..53c93e52df 100644
--- a/tests/gpu/src/imgproc_gpu.cpp
+++ b/tests/gpu/src/imgproc_gpu.cpp
@@ -775,6 +775,56 @@ struct CV_GpuCornerMinEigenValTest: CvTest
     }
 };
 
+struct CV_GpuColumnSumTest: CvTest 
+{
+    CV_GpuColumnSumTest(): CvTest("GPU-ColumnSumTest", "columnSum") {}
+
+    void run(int)
+    {
+        try
+        {
+            int n = 375;
+            int m = 1072;
+            Mat src(n, m, CV_32F);
+            RNG rng;
+            rng.fill(src, RNG::UNIFORM, Scalar(0), Scalar(1));
+            Mat dst_gold, dst2_gold;
+
+            integral(src, dst_gold, dst2_gold);
+
+            GpuMat dsrc(src);
+            GpuMat buf;
+            GpuMat dst;
+            columnSum(dsrc, buf);
+            transpose(buf, dst);
+            columnSum(dst, buf);
+            transpose(buf, dst);
+            
+            Mat dst_ = dst;
+            for (int i = 0; i < dst_.rows; ++i)
+            {
+                const double* dst_gold_data = (const double*)dst_gold.ptr(i + 1);
+                for (int j = 0; j < dst_.cols; ++j)
+                {
+                    float a = (float)dst_gold_data[j + 1];
+                    float b = dst_.at<float>(i, j);
+                    if (fabs(a - b) > 0.5f)
+                    {
+                        ts->printf(CvTS::CONSOLE, "%d %d %f %f\n", i, j, a, b);
+                        ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);
+                        return;
+                    }
+                }
+            }
+        }
+        catch (const Exception& e)
+        {
+            if (!check_and_treat_gpu_exception(e, ts)) throw;
+            return;
+        }
+    }
+};
+
 /////////////////////////////////////////////////////////////////////////////
 /////////////////// tests registration  /////////////////////////////////////
 /////////////////////////////////////////////////////////////////////////////
@@ -794,4 +844,5 @@ CV_GpuCvtColorTest CV_GpuCvtColor_test;
 CV_GpuHistogramsTest CV_GpuHistograms_test;
 CV_GpuCornerHarrisTest CV_GpuCornerHarris_test;
 CV_GpuCornerMinEigenValTest CV_GpuCornerMinEigenVal_test;
+CV_GpuColumnSumTest CV_GpuColumnSum_test;