generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Remerged the 2.4.0 branch

Browse Source
This commit is contained in:
Andrey Kamaev
2012-05-02 20:20:14 +00:00
parent 77717e1abc
commit 3a1f85d4e8
41 changed files with 2452 additions and 4857 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
modules/gpu/src/cuda/bf_match.cu
View File

@@ -97,7 +97,7 @@ namespace cv { namespace gpu { namespace device
}
template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>
__device__ void loopUnrolledCached(int queryIdx, const DevMem2D_<T>& query, int imgIdx, const DevMem2D_<T>& train, const Mask& mask,
__device__ void loopUnrolledCached(int queryIdx, const DevMem2D_<T>& query,volatile int imgIdx, const DevMem2D_<T>& train, const Mask& mask,
typename Dist::value_type* s_query, typename Dist::value_type* s_train,
float& bestDistance, int& bestTrainIdx, int& bestImgIdx)
{
@@ -253,7 +253,7 @@ namespace cv { namespace gpu { namespace device
// Match Unrolled
template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>
__device__ void loopUnrolled(int queryIdx, const DevMem2D_<T>& query, int imgIdx, const DevMem2D_<T>& train, const Mask& mask,
__device__ void loopUnrolled(int queryIdx, const DevMem2D_<T>& query,volatile int imgIdx, const DevMem2D_<T>& train, const Mask& mask,
typename Dist::value_type* s_query, typename Dist::value_type* s_train,
float& bestDistance, int& bestTrainIdx, int& bestImgIdx)
{
@@ -409,7 +409,7 @@ namespace cv { namespace gpu { namespace device
// Match
template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>
__device__ void loop(int queryIdx, const DevMem2D_<T>& query, int imgIdx, const DevMem2D_<T>& train, const Mask& mask,
__device__ void loop(int queryIdx, const DevMem2D_<T>& query, volatile int imgIdx, const DevMem2D_<T>& train, const Mask& mask,
typename Dist::value_type* s_query, typename Dist::value_type* s_train,
float& bestDistance, int& bestTrainIdx, int& bestImgIdx)
{

9
modules/gpu/src/pyrlk.cpp
View File

@@ -126,18 +126,19 @@ void cv::gpu::PyrLKOpticalFlow::buildImagePyramid(const GpuMat& img0, vector<Gpu
namespace
{
void calcPatchSize(cv::Size winSize, int cn, dim3& block, dim3& patch)
void calcPatchSize(cv::Size winSize, int cn, dim3& block, dim3& patch, bool isDeviceArch11)
{
winSize.width *= cn;
if (winSize.width > 32 && winSize.width > 2 * winSize.height)
{
block.x = 32;
block.x = isDeviceArch11 ? 16 : 32;
block.y = 8;
}
else
{
block.x = block.y = 16;
block.x = 16;
block.y = isDeviceArch11 ? 8 : 16;
}
patch.x = (winSize.width + block.x - 1) / block.x;
@@ -166,7 +167,7 @@ void cv::gpu::PyrLKOpticalFlow::sparse(const GpuMat& prevImg, const GpuMat& next
const int cn = prevImg.channels();
dim3 block, patch;
calcPatchSize(winSize, cn, block, patch);
calcPatchSize(winSize, cn, block, patch, isDeviceArch11_);
CV_Assert(derivLambda >= 0);
CV_Assert(maxLevel >= 0 && winSize.width > 2 && winSize.height > 2);