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

minor fix

Browse Source
This commit is contained in:
Vladislav Vinogradov 2011-10-12 11:36:10 +00:00
parent da3a60a5a7
commit 631d4b483a
3 changed files with 15 additions and 15 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
modules/gpu/src/cuda/bf_knnmatch.cu
View File

@ -742,7 +742,7 @@ namespace cv { namespace gpu { namespace bf_knnmatch
{ {
matchUnrolledCached<16, 128, Dist>(query, train, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> > (distance), stream); matchUnrolledCached<16, 128, Dist>(query, train, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> > (distance), stream);
} }
else if (query.cols <= 256) /*else if (query.cols <= 256)
{ {
matchUnrolled<16, 256, Dist>(query, train, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> > (distance), stream); matchUnrolled<16, 256, Dist>(query, train, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> > (distance), stream);
} }
@ -753,7 +753,7 @@ namespace cv { namespace gpu { namespace bf_knnmatch
else if (query.cols <= 1024) else if (query.cols <= 1024)
{ {
matchUnrolled<16, 1024, Dist>(query, train, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> > (distance), stream); matchUnrolled<16, 1024, Dist>(query, train, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> > (distance), stream);
} }*/
else else
{ {
match<16, Dist>(query, train, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> > (distance), stream); match<16, Dist>(query, train, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<float2> > (distance), stream);
@ -773,7 +773,7 @@ namespace cv { namespace gpu { namespace bf_knnmatch
{ {
matchUnrolledCached<16, 128, Dist>(query, trains, n, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<int2> >(imgIdx), static_cast< DevMem2D_<float2> > (distance), stream); matchUnrolledCached<16, 128, Dist>(query, trains, n, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<int2> >(imgIdx), static_cast< DevMem2D_<float2> > (distance), stream);
} }
else if (query.cols <= 256) /*else if (query.cols <= 256)
{ {
matchUnrolled<16, 256, Dist>(query, trains, n, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<int2> >(imgIdx), static_cast< DevMem2D_<float2> > (distance), stream); matchUnrolled<16, 256, Dist>(query, trains, n, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<int2> >(imgIdx), static_cast< DevMem2D_<float2> > (distance), stream);
} }
@ -784,7 +784,7 @@ namespace cv { namespace gpu { namespace bf_knnmatch
else if (query.cols <= 1024) else if (query.cols <= 1024)
{ {
matchUnrolled<16, 1024, Dist>(query, trains, n, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<int2> >(imgIdx), static_cast< DevMem2D_<float2> > (distance), stream); matchUnrolled<16, 1024, Dist>(query, trains, n, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<int2> >(imgIdx), static_cast< DevMem2D_<float2> > (distance), stream);
} }*/
else else
{ {
match<16, Dist>(query, trains, n, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<int2> >(imgIdx), static_cast< DevMem2D_<float2> > (distance), stream); match<16, Dist>(query, trains, n, mask, static_cast< DevMem2D_<int2> >(trainIdx), static_cast< DevMem2D_<int2> >(imgIdx), static_cast< DevMem2D_<float2> > (distance), stream);
@ -937,7 +937,7 @@ namespace cv { namespace gpu { namespace bf_knnmatch
{ {
calcDistanceUnrolled<16, 128, Dist>(query, train, mask, allDist, stream); calcDistanceUnrolled<16, 128, Dist>(query, train, mask, allDist, stream);
} }
else if (query.cols <= 256) /*else if (query.cols <= 256)
{ {
calcDistanceUnrolled<16, 256, Dist>(query, train, mask, allDist, stream); calcDistanceUnrolled<16, 256, Dist>(query, train, mask, allDist, stream);
} }
@ -948,7 +948,7 @@ namespace cv { namespace gpu { namespace bf_knnmatch
else if (query.cols <= 1024) else if (query.cols <= 1024)
{ {
calcDistanceUnrolled<16, 1024, Dist>(query, train, mask, allDist, stream); calcDistanceUnrolled<16, 1024, Dist>(query, train, mask, allDist, stream);
} }*/
else else
{ {
calcDistance<16, Dist>(query, train, mask, allDist, stream); calcDistance<16, Dist>(query, train, mask, allDist, stream);

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

@ -536,7 +536,7 @@ namespace cv { namespace gpu { namespace bf_match
const dim3 block(BLOCK_SIZE, BLOCK_SIZE); const dim3 block(BLOCK_SIZE, BLOCK_SIZE);
const dim3 grid(divUp(query.rows, BLOCK_SIZE)); const dim3 grid(divUp(query.rows, BLOCK_SIZE));
const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); const size_t smemSize = (3 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int);
match<BLOCK_SIZE, Dist><<<grid, block, smemSize, stream>>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data); match<BLOCK_SIZE, Dist><<<grid, block, smemSize, stream>>>(query, trains, n, mask, trainIdx.data, imgIdx.data, distance.data);
cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaGetLastError() );
@ -561,7 +561,7 @@ namespace cv { namespace gpu { namespace bf_match
{ {
matchUnrolledCached<16, 128, Dist>(query, train, mask, trainIdx, distance, stream); matchUnrolledCached<16, 128, Dist>(query, train, mask, trainIdx, distance, stream);
} }
else if (query.cols <= 256) /*else if (query.cols <= 256)
{ {
matchUnrolled<16, 256, Dist>(query, train, mask, trainIdx, distance, stream); matchUnrolled<16, 256, Dist>(query, train, mask, trainIdx, distance, stream);
} }
@ -572,7 +572,7 @@ namespace cv { namespace gpu { namespace bf_match
else if (query.cols <= 1024) else if (query.cols <= 1024)
{ {
matchUnrolled<16, 1024, Dist>(query, train, mask, trainIdx, distance, stream); matchUnrolled<16, 1024, Dist>(query, train, mask, trainIdx, distance, stream);
} }*/
else else
{ {
match<16, Dist>(query, train, mask, trainIdx, distance, stream); match<16, Dist>(query, train, mask, trainIdx, distance, stream);
@ -592,7 +592,7 @@ namespace cv { namespace gpu { namespace bf_match
{ {
matchUnrolledCached<16, 128, Dist>(query, trains, n, mask, trainIdx, imgIdx, distance, stream); matchUnrolledCached<16, 128, Dist>(query, trains, n, mask, trainIdx, imgIdx, distance, stream);
} }
else if (query.cols <= 256) /*else if (query.cols <= 256)
{ {
matchUnrolled<16, 256, Dist>(query, trains, n, mask, trainIdx, imgIdx, distance, stream); matchUnrolled<16, 256, Dist>(query, trains, n, mask, trainIdx, imgIdx, distance, stream);
} }
@ -603,7 +603,7 @@ namespace cv { namespace gpu { namespace bf_match
else if (query.cols <= 1024) else if (query.cols <= 1024)
{ {
matchUnrolled<16, 1024, Dist>(query, trains, n, mask, trainIdx, imgIdx, distance, stream); matchUnrolled<16, 1024, Dist>(query, trains, n, mask, trainIdx, imgIdx, distance, stream);
} }*/
else else
{ {
match<16, Dist>(query, trains, n, mask, trainIdx, imgIdx, distance, stream); match<16, Dist>(query, trains, n, mask, trainIdx, imgIdx, distance, stream);

8
modules/gpu/src/cuda/bf_radius_match.cu
View File

@ -283,7 +283,7 @@ namespace cv { namespace gpu { namespace bf_radius_match
{ {
matchUnrolled<16, 128, Dist>(query, train, maxDistance, mask, trainIdx, distance, nMatches, stream); matchUnrolled<16, 128, Dist>(query, train, maxDistance, mask, trainIdx, distance, nMatches, stream);
} }
else if (query.cols <= 256) /*else if (query.cols <= 256)
{ {
matchUnrolled<16, 256, Dist>(query, train, maxDistance, mask, trainIdx, distance, nMatches, stream); matchUnrolled<16, 256, Dist>(query, train, maxDistance, mask, trainIdx, distance, nMatches, stream);
} }
@ -294,7 +294,7 @@ namespace cv { namespace gpu { namespace bf_radius_match
else if (query.cols <= 1024) else if (query.cols <= 1024)
{ {
matchUnrolled<16, 1024, Dist>(query, train, maxDistance, mask, trainIdx, distance, nMatches, stream); matchUnrolled<16, 1024, Dist>(query, train, maxDistance, mask, trainIdx, distance, nMatches, stream);
} }*/
else else
{ {
match<16, Dist>(query, train, maxDistance, mask, trainIdx, distance, nMatches, stream); match<16, Dist>(query, train, maxDistance, mask, trainIdx, distance, nMatches, stream);
@ -314,7 +314,7 @@ namespace cv { namespace gpu { namespace bf_radius_match
{ {
matchUnrolled<16, 128, Dist>(query, trains, n, maxDistance, masks, trainIdx, imgIdx, distance, nMatches, stream); matchUnrolled<16, 128, Dist>(query, trains, n, maxDistance, masks, trainIdx, imgIdx, distance, nMatches, stream);
} }
else if (query.cols <= 256) /*else if (query.cols <= 256)
{ {
matchUnrolled<16, 256, Dist>(query, trains, n, maxDistance, masks, trainIdx, imgIdx, distance, nMatches, stream); matchUnrolled<16, 256, Dist>(query, trains, n, maxDistance, masks, trainIdx, imgIdx, distance, nMatches, stream);
} }
@ -325,7 +325,7 @@ namespace cv { namespace gpu { namespace bf_radius_match
else if (query.cols <= 1024) else if (query.cols <= 1024)
{ {
matchUnrolled<16, 1024, Dist>(query, trains, n, maxDistance, masks, trainIdx, imgIdx, distance, nMatches, stream); matchUnrolled<16, 1024, Dist>(query, trains, n, maxDistance, masks, trainIdx, imgIdx, distance, nMatches, stream);
} }*/
else else
{ {
match<16, Dist>(query, trains, n, maxDistance, masks, trainIdx, imgIdx, distance, nMatches, stream); match<16, Dist>(query, trains, n, maxDistance, masks, trainIdx, imgIdx, distance, nMatches, stream);