added gpu transpose and integral based on NPP Staging.

added mask support to SURF_GPU.

modules/gpu/include/opencv2/gpu/gpu.hpp

@@ -364,7 +364,7 @@ namespace cv
         ////////////////////////////// Arithmetics ///////////////////////////////////
 
         //! transposes the matrix
-        //! supports CV_8UC1, CV_8SC1, CV_8UC4, CV_8SC4, CV_16UC2, CV_16SC2, CV_32SC1, CV_32FC1 type
+        //! supports matrix with element size = 1, 4 and 8 bytes (CV_8UC1, CV_8UC4, CV_16UC2, CV_32FC1, etc)
         CV_EXPORTS void transpose(const GpuMat& src1, GpuMat& dst);
 
         //! reverses the order of the rows, columns or both in a matrix
@@ -594,6 +594,11 @@ namespace cv
         //! supports CV_8UC1, CV_8UC4, CV_32SC1 and CV_32FC1 types
         CV_EXPORTS void copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value = Scalar());
 
+        //! computes the integral image
+        //! sum will have CV_32S type, but will contain unsigned int values
+        //! supports only CV_8UC1 source type
+        CV_EXPORTS void integral(const GpuMat& src, GpuMat& sum);
+
         //! computes the integral image and integral for the squared image
         //! sum will have CV_32S type, sqsum - CV32F type
         //! supports only CV_8UC1 source type
@@ -1433,27 +1438,28 @@ namespace cv
             static void downloadDescriptors(const GpuMat& descriptorsGPU, vector<float>& descriptors);
             
             //! finds the keypoints using fast hessian detector used in SURF
-            //! supports CV_8UC1 (0..255) and CV_32FC1 (0..1) images
+            //! supports CV_8UC1 images
             //! keypoints will have 1 row and type CV_32FC(6)
-            //! keypoints.at<float6>(1, i) contains i'th keypoint
+            //! keypoints.at<float[6]>(1, i) contains i'th keypoint
             //! format: (x, y, size, response, angle, octave)
-            void operator()(const GpuMat& img, GpuMat& keypoints);
+            void operator()(const GpuMat& img, const GpuMat& mask, GpuMat& keypoints);
             //! finds the keypoints and computes their descriptors. 
             //! Optionally it can compute descriptors for the user-provided keypoints and recompute keypoints direction
-            void operator()(const GpuMat& img, GpuMat& keypoints, GpuMat& descriptors, 
+            void operator()(const GpuMat& img, const GpuMat& mask, GpuMat& keypoints, GpuMat& descriptors, 
                 bool useProvidedKeypoints = false, bool calcOrientation = true);
         
-            void operator()(const GpuMat& img, std::vector<KeyPoint>& keypoints);
+            void operator()(const GpuMat& img, const GpuMat& mask, std::vector<KeyPoint>& keypoints);
-            void operator()(const GpuMat& img, std::vector<KeyPoint>& keypoints, GpuMat& descriptors, 
+            void operator()(const GpuMat& img, const GpuMat& mask, std::vector<KeyPoint>& keypoints, GpuMat& descriptors, 
                 bool useProvidedKeypoints = false, bool calcOrientation = true);
             
-            void operator()(const GpuMat& img, std::vector<KeyPoint>& keypoints, std::vector<float>& descriptors, 
+            void operator()(const GpuMat& img, const GpuMat& mask, std::vector<KeyPoint>& keypoints, std::vector<float>& descriptors, 
                 bool useProvidedKeypoints = false, bool calcOrientation = true);
 
-            GpuMat img_float;
-            GpuMat img_float_tr;
-
             GpuMat sum;
+            GpuMat sumf;
+
+            GpuMat mask1;
+            GpuMat maskSum;
 
             GpuMat hessianBuffer;
             GpuMat maxPosBuffer;

modules/gpu/src/arithm.cpp

@@ -71,19 +71,13 @@ void cv::gpu::polarToCart(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool,
 ////////////////////////////////////////////////////////////////////////
 // transpose
 
-namespace cv { namespace gpu { namespace mathfunc
-{
-    void transpose_gpu(const DevMem2Di& src, const DevMem2Di& dst);
-}}}
-
 void cv::gpu::transpose(const GpuMat& src, GpuMat& dst)
 {
-    CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8SC1 || src.type() == CV_8UC4 || src.type() == CV_8SC4 
+    CV_Assert(src.elemSize() == 1 || src.elemSize() == 4 || src.elemSize() == 8);
-        || src.type() == CV_16UC2 || src.type() == CV_16SC2 || src.type() == CV_32SC1 || src.type() == CV_32FC1);
 
     dst.create( src.cols, src.rows, src.type() );
 
-    if (src.type() == CV_8UC1 || src.type() == CV_8SC1)
+    if (src.elemSize() == 1)
     {
         NppiSize sz;
         sz.width  = src.cols;
@@ -91,9 +85,23 @@ void cv::gpu::transpose(const GpuMat& src, GpuMat& dst)
 
         nppSafeCall( nppiTranspose_8u_C1R(src.ptr<Npp8u>(), src.step, dst.ptr<Npp8u>(), dst.step, sz) );
     }
-    else
+    else if (src.elemSize() == 4)
     {
-        mathfunc::transpose_gpu(src, dst);
+        NppStSize32u sz;
+        sz.width  = src.cols;
+        sz.height = src.rows;
+
+        nppSafeCall( nppiStTranspose_32u_C1R(const_cast<NppSt32u*>(src.ptr<NppSt32u>()), src.step, 
+            dst.ptr<NppSt32u>(), dst.step, sz) );
+    }
+    else // if (src.elemSize() == 8)
+    {
+        NppStSize32u sz;
+        sz.width  = src.cols;
+        sz.height = src.rows;
+
+        nppSafeCall( nppiStTranspose_64u_C1R(const_cast<NppSt64u*>(src.ptr<NppSt64u>()), src.step, 
+            dst.ptr<NppSt64u>(), dst.step, sz) );
     }
 }
 

modules/gpu/src/cuda/mathfunc.cu

@@ -214,44 +214,6 @@ namespace cv { namespace gpu { namespace mathfunc
 
         callers[mag.data == 0](mag, angle, x, y, angleInDegrees, stream);
     }
-
-
-//////////////////////////////////////////////////////////////////////////////////////////////////////////
-// transpose
-
-    __global__ void transpose(const DevMem2Di src, PtrStepi dst)
-    {
-    	__shared__ int s_mem[16 * 17];
-
-    	int x = blockIdx.x * blockDim.x + threadIdx.x;
-    	int y = blockIdx.y * blockDim.y + threadIdx.y;
-	    int smem_idx = threadIdx.y * blockDim.x + threadIdx.x + threadIdx.y;
-
-	    if (y < src.rows && x < src.cols)
-	    {
-            s_mem[smem_idx] = src.ptr(y)[x];
-	    }
-	    __syncthreads();
-
-	    smem_idx = threadIdx.x * blockDim.x + threadIdx.y + threadIdx.x;
-
-	    x = blockIdx.y * blockDim.x + threadIdx.x;
-	    y = blockIdx.x * blockDim.y + threadIdx.y;
-
-	    if (y < src.cols && x < src.rows)
-	    {
-		    dst.ptr(y)[x] = s_mem[smem_idx];
-	    }
-    }
-
-    void transpose_gpu(const DevMem2Di& src, const DevMem2Di& dst)
-    {
-	    dim3 threads(16, 16, 1);
-	    dim3 grid(divUp(src.cols, 16), divUp(src.rows, 16), 1);
-
-	    transpose<<<grid, threads>>>(src, dst);
-        cudaSafeCall( cudaThreadSynchronize() );
-    }
 }}}
 
 

modules/gpu/src/cuda/surf.cu

@@ -259,7 +259,36 @@ namespace cv { namespace gpu { namespace surf
 
     ////////////////////////////////////////////////////////////////////////
     // NONMAX
+    
+    texture<int, 2, cudaReadModeElementType> maskSumTex(0, cudaFilterModePoint, cudaAddressModeClamp);
 
+    struct WithOutMask
+    {
+        static __device__ bool check(float, float, float)
+        {
+            return true;
+        }
+    };
+    struct WithMask
+    {
+        static __device__ bool check(float x, float y, float fscale)
+        {
+	        float half_width = fscale / 2;
+    
+	        float result = 0.f;
+
+            result += tex2D(maskSumTex, x - half_width, y - half_width);
+            result -= tex2D(maskSumTex, x + half_width, y - half_width);
+            result -= tex2D(maskSumTex, x - half_width, y + half_width);
+            result += tex2D(maskSumTex, x + half_width, y + half_width);
+	
+	        result /= (fscale * fscale);
+
+            return (result >= 0.5f);
+        }
+    };
+
+    template <typename Mask>
     __global__ void nonmaxonly(PtrStepf hessianBuffer, int4* maxPosBuffer, unsigned int* maxCounter)
     {        
         #if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 110
@@ -287,7 +316,12 @@ namespace cv { namespace gpu { namespace surf
 
         float val = fh_vals[localLin];
 
-        if (inBounds2 && val >= c_threshold)
+	    // Compute the lookup location of the mask center
+        float x = hidx_x * c_step + c_border;
+        float y = hidx_y * c_step + c_border;
+        float fscale = calcScale(hidx_z);
+
+        if (inBounds2 && val >= c_threshold && Mask::check(x, y, fscale))
         {
             // Check to see if we have a max (in its 26 neighbours)
             int zoff = blockDim.x * blockDim.y;
@@ -337,7 +371,7 @@ namespace cv { namespace gpu { namespace surf
     }
 
     void nonmaxonly_gpu(PtrStepf hessianBuffer, int4* maxPosBuffer, unsigned int& maxCounter, 
-        int nIntervals, int x_size, int y_size)
+        int nIntervals, int x_size, int y_size, bool use_mask)
     {
         dim3 threads;
         threads.x = 16;
@@ -353,7 +387,10 @@ namespace cv { namespace gpu { namespace surf
 
         DeviceReference<unsigned int> maxCounterWrapper(maxCounter);
 
-        nonmaxonly<<<grid, threads, smem_size>>>(hessianBuffer, maxPosBuffer, maxCounterWrapper);
+        if (use_mask)
+            nonmaxonly<WithMask><<<grid, threads, smem_size>>>(hessianBuffer, maxPosBuffer, maxCounterWrapper);
+        else
+            nonmaxonly<WithOutMask><<<grid, threads, smem_size>>>(hessianBuffer, maxPosBuffer, maxCounterWrapper);
 
         cudaSafeCall( cudaThreadSynchronize() );
     }

modules/gpu/src/imgproc_gpu.cpp

@@ -60,6 +60,7 @@ void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, const S
 void cv::gpu::warpAffine(const GpuMat&, GpuMat&, const Mat&, Size, int) { throw_nogpu(); }
 void cv::gpu::warpPerspective(const GpuMat&, GpuMat&, const Mat&, Size, int) { throw_nogpu(); }
 void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int) { throw_nogpu(); }
+void cv::gpu::integral(const GpuMat&, GpuMat&) { throw_nogpu(); }
 void cv::gpu::integral(const GpuMat&, GpuMat&, GpuMat&) { throw_nogpu(); }
 void cv::gpu::columnSum(const GpuMat&, GpuMat&) { throw_nogpu(); }
 void cv::gpu::rectStdDev(const GpuMat&, const GpuMat&, GpuMat&, const Rect&) { throw_nogpu(); }
@@ -547,6 +548,26 @@ void cv::gpu::rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, d
 ////////////////////////////////////////////////////////////////////////
 // integral
 
+void cv::gpu::integral(const GpuMat& src, GpuMat& sum)
+{
+    CV_Assert(src.type() == CV_8UC1);
+
+    sum.create(src.rows + 1, src.cols + 1, CV_32S);
+    
+    NppStSize32u roiSize;
+    roiSize.width = src.cols;
+    roiSize.height = src.rows;
+
+    NppSt32u bufSize;
+
+    nppSafeCall( nppiStIntegralGetSize_8u32u(roiSize, &bufSize) );
+
+    GpuMat buffer(1, bufSize, CV_8UC1);
+
+    nppSafeCall( nppiStIntegral_8u32u_C1R(const_cast<NppSt8u*>(src.ptr<NppSt8u>()), src.step, 
+        sum.ptr<NppSt32u>(), sum.step, roiSize, buffer.ptr<NppSt8u>(), bufSize) );
+}
+
 void cv::gpu::integral(const GpuMat& src, GpuMat& sum, GpuMat& sqsum)
 {
     CV_Assert(src.type() == CV_8UC1);

modules/gpu/src/surf.cpp

@@ -52,11 +52,11 @@ int cv::gpu::SURF_GPU::descriptorSize() const { throw_nogpu(); return 0;}
 void cv::gpu::SURF_GPU::uploadKeypoints(const vector<KeyPoint>&, GpuMat&) { throw_nogpu(); }
 void cv::gpu::SURF_GPU::downloadKeypoints(const GpuMat&, vector<KeyPoint>&) { throw_nogpu(); }
 void cv::gpu::SURF_GPU::downloadDescriptors(const GpuMat&, vector<float>&) { throw_nogpu(); }
-void cv::gpu::SURF_GPU::operator()(const GpuMat&, GpuMat&) { throw_nogpu(); }
+void cv::gpu::SURF_GPU::operator()(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }
-void cv::gpu::SURF_GPU::operator()(const GpuMat&, GpuMat&, GpuMat&, bool, bool) { throw_nogpu(); }
+void cv::gpu::SURF_GPU::operator()(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, bool) { throw_nogpu(); }
-void cv::gpu::SURF_GPU::operator()(const GpuMat&, vector<KeyPoint>&) { throw_nogpu(); }
+void cv::gpu::SURF_GPU::operator()(const GpuMat&, const GpuMat&, vector<KeyPoint>&) { throw_nogpu(); }
-void cv::gpu::SURF_GPU::operator()(const GpuMat&, vector<KeyPoint>&, GpuMat&, bool, bool) { throw_nogpu(); }
+void cv::gpu::SURF_GPU::operator()(const GpuMat&, const GpuMat&, vector<KeyPoint>&, GpuMat&, bool, bool) { throw_nogpu(); }
-void cv::gpu::SURF_GPU::operator()(const GpuMat&, vector<KeyPoint>&, vector<float>&, bool, bool) { throw_nogpu(); }
+void cv::gpu::SURF_GPU::operator()(const GpuMat&, const GpuMat&, vector<KeyPoint>&, vector<float>&, bool, bool) { throw_nogpu(); }
 
 #else /* !defined (HAVE_CUDA) */
 
@@ -65,7 +65,7 @@ namespace cv { namespace gpu { namespace surf
     void fasthessian_gpu(PtrStepf hessianBuffer, int nIntervals, int x_size, int y_size);
     
     void nonmaxonly_gpu(PtrStepf hessianBuffer, int4* maxPosBuffer, unsigned int& maxCounter, 
-        int nIntervals, int x_size, int y_size);
+        int nIntervals, int x_size, int y_size, bool use_mask);
     
     void fh_interp_extremum_gpu(PtrStepf hessianBuffer, const int4* maxPosBuffer, unsigned int maxCounter, 
         KeyPoint_GPU* featuresBuffer, unsigned int& featureCounter);
@@ -82,12 +82,12 @@ namespace
     class SURF_GPU_Invoker : private SURFParams_GPU
     {
     public:
-        SURF_GPU_Invoker(SURF_GPU& surf, const GpuMat& img) : 
+        SURF_GPU_Invoker(SURF_GPU& surf, const GpuMat& img, const GpuMat& mask) : 
             SURFParams_GPU(surf),
 
-            img_float(surf.img_float), img_float_tr(surf.img_float_tr),
+            sum(surf.sum), sumf(surf.sumf),
 
-            sum(surf.sum), 
+            mask1(surf.mask1), maskSum(surf.maskSum),
 
             hessianBuffer(surf.hessianBuffer), 
             maxPosBuffer(surf.maxPosBuffer), 
@@ -95,11 +95,15 @@ namespace
 
             img_cols(img.cols), img_rows(img.rows),
 
+            use_mask(!mask.empty()),
+
             mask_width(0), mask_height(0),
 
             featureCounter(0), maxCounter(0)
         {
-            CV_Assert((img.type() == CV_8UC1 || img.type() == CV_32FC1) && nOctaves > 0 && nIntervals > 2);
+            CV_Assert(img.type() == CV_8UC1);
+            CV_Assert(mask.empty() || (mask.size() == img.size() && mask.type() == CV_8UC1));
+            CV_Assert(nOctaves > 0 && nIntervals > 2);
             CV_Assert(hasAtomicsSupport(getDevice()));
 
             max_features = static_cast<int>(img.size().area() * featuresRatio);
@@ -139,22 +143,25 @@ namespace
             
             hessianBuffer.create(height0 * nIntervals, width0, CV_32F);
 
-            if (img.type() == CV_32FC1)
+            integral(img, sum);
-                img_float = img;
+            sum.convertTo(sumf, CV_32F, 1.0 / 255.0);
-            else
-                img.convertTo(img_float, CV_32F, 1.0 / 255.0);
-
-            transpose(img_float, img_float_tr);
-            columnSum(img_float_tr, img_float_tr);
-            transpose(img_float_tr, sum);
-            columnSum(sum, sum);
             
-            bindTexture("cv::gpu::surf::sumTex", (DevMem2Df)sum);
+            bindTexture("cv::gpu::surf::sumTex", (DevMem2Df)sumf);
+
+            if (!mask.empty())
+		    {
+                min(mask, 1.0, mask1);
+                integral(mask1, maskSum);
+            
+                bindTexture("cv::gpu::surf::maskSumTex", (DevMem2Di)maskSum);
+		    }
         }
 
         ~SURF_GPU_Invoker()
         {
             unbindTexture("cv::gpu::surf::sumTex");
+            if (use_mask)
+                unbindTexture("cv::gpu::surf::maskSumTex");
         }
 
         void detectKeypoints(GpuMat& keypoints)
@@ -185,7 +192,7 @@ namespace
                 // Reset the candidate count.
                 maxCounter = 0;
 
-                nonmaxonly_gpu(hessianBuffer, maxPosBuffer.ptr<int4>(), maxCounter, nIntervals, x_size, y_size); 
+                nonmaxonly_gpu(hessianBuffer, maxPosBuffer.ptr<int4>(), maxCounter, nIntervals, x_size, y_size, use_mask); 
                 
                 maxCounter = std::min(maxCounter, static_cast<unsigned int>(max_candidates));
 
@@ -214,16 +221,19 @@ namespace
         }
 
     private:
-        GpuMat& img_float;
-        GpuMat& img_float_tr;
-
         GpuMat& sum;
+        GpuMat& sumf;
+
+        GpuMat& mask1;
+        GpuMat& maskSum;
 
         GpuMat& hessianBuffer;
         GpuMat& maxPosBuffer;
         GpuMat& featuresBuffer;
 
         int img_cols, img_rows;
+
+        bool use_mask;
         
         float mask_width, mask_height;
 
@@ -298,19 +308,19 @@ void cv::gpu::SURF_GPU::downloadDescriptors(const GpuMat& descriptorsGPU, vector
     descriptorsGPU.download(descriptorsCPU);
 }
 
-void cv::gpu::SURF_GPU::operator()(const GpuMat& img, GpuMat& keypoints)
+void cv::gpu::SURF_GPU::operator()(const GpuMat& img, const GpuMat& mask, GpuMat& keypoints)
 {
-    SURF_GPU_Invoker surf(*this, img);
+    SURF_GPU_Invoker surf(*this, img, mask);
 
     surf.detectKeypoints(keypoints);
 
     surf.findOrientation(keypoints);
 }
 
-void cv::gpu::SURF_GPU::operator()(const GpuMat& img, GpuMat& keypoints, GpuMat& descriptors, 
+void cv::gpu::SURF_GPU::operator()(const GpuMat& img, const GpuMat& mask, GpuMat& keypoints, GpuMat& descriptors, 
                                    bool useProvidedKeypoints, bool calcOrientation)
 {
-    SURF_GPU_Invoker surf(*this, img);
+    SURF_GPU_Invoker surf(*this, img, mask);
     
     if (!useProvidedKeypoints)
         surf.detectKeypoints(keypoints);
@@ -321,34 +331,34 @@ void cv::gpu::SURF_GPU::operator()(const GpuMat& img, GpuMat& keypoints, GpuMat&
     surf.computeDescriptors(keypoints, descriptors, descriptorSize());
 }
 
-void cv::gpu::SURF_GPU::operator()(const GpuMat& img, vector<KeyPoint>& keypoints)
+void cv::gpu::SURF_GPU::operator()(const GpuMat& img, const GpuMat& mask, vector<KeyPoint>& keypoints)
 {
     GpuMat keypointsGPU;
 
-    (*this)(img, keypointsGPU);
+    (*this)(img, mask, keypointsGPU);
 
     downloadKeypoints(keypointsGPU, keypoints);
 }
 
-void cv::gpu::SURF_GPU::operator()(const GpuMat& img, vector<KeyPoint>& keypoints, GpuMat& descriptors, 
+void cv::gpu::SURF_GPU::operator()(const GpuMat& img, const GpuMat& mask, vector<KeyPoint>& keypoints, 
-                                   bool useProvidedKeypoints, bool calcOrientation)
+    GpuMat& descriptors, bool useProvidedKeypoints, bool calcOrientation)
 {
     GpuMat keypointsGPU;
 
     if (useProvidedKeypoints)
         uploadKeypoints(keypoints, keypointsGPU);    
 
-    (*this)(img, keypointsGPU, descriptors, useProvidedKeypoints, calcOrientation);
+    (*this)(img, mask, keypointsGPU, descriptors, useProvidedKeypoints, calcOrientation);
 
     downloadKeypoints(keypointsGPU, keypoints);
 }
 
-void cv::gpu::SURF_GPU::operator()(const GpuMat& img, vector<KeyPoint>& keypoints, vector<float>& descriptors, 
+void cv::gpu::SURF_GPU::operator()(const GpuMat& img, const GpuMat& mask, vector<KeyPoint>& keypoints, 
-                                   bool useProvidedKeypoints, bool calcOrientation)
+    vector<float>& descriptors, bool useProvidedKeypoints, bool calcOrientation)
 {
     GpuMat descriptorsGPU;
 
-    (*this)(img, keypoints, descriptorsGPU, useProvidedKeypoints, calcOrientation);
+    (*this)(img, mask, keypoints, descriptorsGPU, useProvidedKeypoints, calcOrientation);
 
     downloadDescriptors(descriptorsGPU, descriptors);
 }

tests/gpu/src/imgproc_gpu.cpp

@@ -384,29 +384,14 @@ struct CV_GpuNppImageIntegralTest : public CV_GpuImageProcTest
             return CvTS::OK;
         }
 
-        Mat cpusum, cpusqsum;
+        Mat cpusum;
-        cv::integral(img, cpusum, cpusqsum, CV_32S);
+        cv::integral(img, cpusum, CV_32S);
 
         GpuMat gpu1(img);
-        GpuMat gpusum, gpusqsum;
+        GpuMat gpusum;
-        cv::gpu::integral(gpu1, gpusum, gpusqsum);
+        cv::gpu::integral(gpu1, gpusum);
 
-        gpusqsum.convertTo(gpusqsum, CV_64F);
+        return CheckNorm(cpusum, gpusum) == CvTS::OK ? CvTS::OK : CvTS::FAIL_GENERIC;
-
-        int test_res = CvTS::OK;
-
-        if (CheckNorm(cpusum, gpusum) != CvTS::OK)
-        {
-            ts->printf(CvTS::LOG, "\nSum failed\n");
-            test_res = CvTS::FAIL_GENERIC;
-        }
-        if (CheckNorm(cpusqsum, gpusqsum) != CvTS::OK)
-        {
-            ts->printf(CvTS::LOG, "\nSquared sum failed\n");
-            test_res = CvTS::FAIL_GENERIC;
-        }
-
-        return test_res;
     }
 };