gftt.cu 5.4 KB
Newer Older
wester committed
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49
/*M///////////////////////////////////////////////////////////////////////////////////////
//
//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
//  By downloading, copying, installing or using the software you agree to this license.
//  If you do not agree to this license, do not download, install,
//  copy or use the software.
//
//
//                           License Agreement
//                For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
//   * Redistribution's of source code must retain the above copyright notice,
//     this list of conditions and the following disclaimer.
//
//   * Redistribution's in binary form must reproduce the above copyright notice,
//     this list of conditions and the following disclaimer in the documentation
//     and/or other materials provided with the distribution.
//
//   * The name of the copyright holders may not be used to endorse or promote products
//     derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/

#if !defined CUDA_DISABLER

#include <thrust/device_ptr.h>
#include <thrust/sort.h>

#include "opencv2/core/cuda/common.hpp"
#include "opencv2/core/cuda/utility.hpp"
a  
Kai Westerkamp committed
50

wester committed
51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93
namespace cv { namespace cuda { namespace device
{
    namespace gfft
    {
        texture<float, cudaTextureType2D, cudaReadModeElementType> eigTex(0, cudaFilterModePoint, cudaAddressModeClamp);

        __device__ int g_counter = 0;

        template <class Mask> __global__ void findCorners(float threshold, const Mask mask, float2* corners, int max_count, int rows, int cols)
        {
            const int j = blockIdx.x * blockDim.x + threadIdx.x;
            const int i = blockIdx.y * blockDim.y + threadIdx.y;

            if (i > 0 && i < rows - 1 && j > 0 && j < cols - 1 && mask(i, j))
            {
                float val = tex2D(eigTex, j, i);

                if (val > threshold)
                {
                    float maxVal = val;

                    maxVal = ::fmax(tex2D(eigTex, j - 1, i - 1), maxVal);
                    maxVal = ::fmax(tex2D(eigTex, j    , i - 1), maxVal);
                    maxVal = ::fmax(tex2D(eigTex, j + 1, i - 1), maxVal);

                    maxVal = ::fmax(tex2D(eigTex, j - 1, i), maxVal);
                    maxVal = ::fmax(tex2D(eigTex, j + 1, i), maxVal);

                    maxVal = ::fmax(tex2D(eigTex, j - 1, i + 1), maxVal);
                    maxVal = ::fmax(tex2D(eigTex, j    , i + 1), maxVal);
                    maxVal = ::fmax(tex2D(eigTex, j + 1, i + 1), maxVal);

                    if (val == maxVal)
                    {
                        const int ind = ::atomicAdd(&g_counter, 1);

                        if (ind < max_count)
                            corners[ind] = make_float2(j, i);
                    }
                }
            }
        }

a  
Kai Westerkamp committed
94
        int findCorners_gpu(PtrStepSzf eig, float threshold, PtrStepSzb mask, float2* corners, int max_count)
wester committed
95 96 97 98
        {
            void* counter_ptr;
            cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, g_counter) );

a  
Kai Westerkamp committed
99
            cudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(int)) );
wester committed
100 101 102 103 104 105 106

            bindTexture(&eigTex, eig);

            dim3 block(16, 16);
            dim3 grid(divUp(eig.cols, block.x), divUp(eig.rows, block.y));

            if (mask.data)
a  
Kai Westerkamp committed
107
                findCorners<<<grid, block>>>(threshold, SingleMask(mask), corners, max_count, eig.rows, eig.cols);
wester committed
108
            else
a  
Kai Westerkamp committed
109
                findCorners<<<grid, block>>>(threshold, WithOutMask(), corners, max_count, eig.rows, eig.cols);
wester committed
110 111 112

            cudaSafeCall( cudaGetLastError() );

a  
Kai Westerkamp committed
113 114
            cudaSafeCall( cudaDeviceSynchronize() );

wester committed
115
            int count;
a  
Kai Westerkamp committed
116 117
            cudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) );

wester committed
118 119 120 121 122 123 124 125 126 127 128 129 130
            return std::min(count, max_count);
        }

        class EigGreater
        {
        public:
            __device__ __forceinline__ bool operator()(float2 a, float2 b) const
            {
                return tex2D(eigTex, a.x, a.y) > tex2D(eigTex, b.x, b.y);
            }
        };


a  
Kai Westerkamp committed
131
        void sortCorners_gpu(PtrStepSzf eig, float2* corners, int count)
wester committed
132 133 134 135
        {
            bindTexture(&eigTex, eig);

            thrust::device_ptr<float2> ptr(corners);
a  
Kai Westerkamp committed
136

wester committed
137 138 139 140 141 142 143
            thrust::sort(ptr, ptr + count, EigGreater());
        }
    } // namespace optical_flow
}}}


#endif /* CUDA_DISABLER */