Chiebot-Mirror
/
opencv
8
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Enable runtime type definition in kernels

Browse Source
pull/811/head
Peng Xiao 12 years ago
parent fd1528795e
commit 1db20099a9
2 changed files with 101 additions and 79 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 28
      modules/ocl/src/brute_force_matcher.cpp
  2. 152
      modules/ocl/src/opencl/brute_force_match.cl

28
modules/ocl/src/brute_force_matcher.cpp
Unescape View File

@ -64,7 +64,14 @@ namespace cv
static const int OPT_SIZE = 100; static const int OPT_SIZE = 100;
static const char * T_ARR [] = {"uchar", "char", "ushort", "short", "int", "float", "double"}; static const char * T_ARR [] = {
"uchar",
"char",
"ushort",
"short",
"int",
"float -D T_FLOAT",
"double"};
template < int BLOCK_SIZE, int MAX_DESC_LEN/*, typename Mask*/ > template < int BLOCK_SIZE, int MAX_DESC_LEN/*, typename Mask*/ >
void matchUnrolledCached(const oclMat &query, const oclMat &train, const oclMat &/*mask*/, void matchUnrolledCached(const oclMat &query, const oclMat &train, const oclMat &/*mask*/,
@ -100,7 +107,7 @@ void matchUnrolledCached(const oclMat &query, const oclMat &train, const oclMat
std::string kernelName = "BruteForceMatch_UnrollMatch"; std::string kernelName = "BruteForceMatch_UnrollMatch";
openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth(), opt); openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt);
} }
} }
@ -126,7 +133,6 @@ void match(const oclMat &query, const oclMat &train, const oclMat &/*mask*/,
sprintf(opt, sprintf(opt,
"-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d", "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d",
T_ARR[query.depth()], distType, block_size); T_ARR[query.depth()], distType, block_size);
if(globalSize[0] != 0) if(globalSize[0] != 0)
{ {
args.push_back( make_pair( sizeof(cl_mem), (void *)&query.data )); args.push_back( make_pair( sizeof(cl_mem), (void *)&query.data ));
@ -143,7 +149,7 @@ void match(const oclMat &query, const oclMat &train, const oclMat &/*mask*/,
std::string kernelName = "BruteForceMatch_Match"; std::string kernelName = "BruteForceMatch_Match";
openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth(), opt); openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt);
} }
} }
@ -192,7 +198,7 @@ void matchUnrolledCached(const oclMat &query, const oclMat &train, float maxDist
std::string kernelName = "BruteForceMatch_RadiusUnrollMatch"; std::string kernelName = "BruteForceMatch_RadiusUnrollMatch";
openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth(), opt); openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt);
} }
} }
@ -234,7 +240,7 @@ void radius_match(const oclMat &query, const oclMat &train, float maxDistance, c
std::string kernelName = "BruteForceMatch_RadiusMatch"; std::string kernelName = "BruteForceMatch_RadiusMatch";
openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth(), opt); openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt);
} }
} }
@ -330,7 +336,7 @@ void knn_matchUnrolledCached(const oclMat &query, const oclMat &train, const ocl
std::string kernelName = "BruteForceMatch_knnUnrollMatch"; std::string kernelName = "BruteForceMatch_knnUnrollMatch";
openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth(), opt); openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt);
} }
} }
@ -366,7 +372,7 @@ void knn_match(const oclMat &query, const oclMat &train, const oclMat &/*mask*/,
std::string kernelName = "BruteForceMatch_knnMatch"; std::string kernelName = "BruteForceMatch_knnMatch";
openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth(), opt); openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt);
} }
} }
@ -403,7 +409,7 @@ void calcDistanceUnrolled(const oclMat &query, const oclMat &train, const oclMat
std::string kernelName = "BruteForceMatch_calcDistanceUnrolled"; std::string kernelName = "BruteForceMatch_calcDistanceUnrolled";
openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth(), opt); openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt);
} }
} }
@ -438,7 +444,7 @@ void calcDistance(const oclMat &query, const oclMat &train, const oclMat &/*mask
std::string kernelName = "BruteForceMatch_calcDistance"; std::string kernelName = "BruteForceMatch_calcDistance";
openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth(), opt); openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt);
} }
} }
@ -500,7 +506,7 @@ void findKnnMatch(int k, const oclMat &trainIdx, const oclMat &distance, const o
//args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols )); //args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols ));
//args.push_back( make_pair( sizeof(cl_int), (void *)&query.step )); //args.push_back( make_pair( sizeof(cl_int), (void *)&query.step ));
openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, trainIdx.depth(), -1); openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1);
} }
} }

152
modules/ocl/src/opencl/brute_force_match.cl
Unescape View File

@ -65,7 +65,7 @@
int bit1Count(int x) int bit1Count(int x)
{ {
int c = 0; int c = 0;
int ix = (int)x; int ix = x;
for (int i = 0 ; i < 32 ; i++) for (int i = 0 ; i < 32 ; i++)
{ {
c += ix & 0x1; c += ix & 0x1;
@ -74,42 +74,60 @@ int bit1Count(int x)
return c; return c;
} }
#if (DIST_TYPE == 0) // dirty fix for non-template support
#define DIST(x, y) fabs((x) - (y)) #if (DIST_TYPE == 0) // L1Dist
#elif (DIST_TYPE == 1) # ifdef T_FLOAT
# define DIST(x, y) fabs((x) - (y))
typedef float value_type;
typedef float result_type;
# else
# define DIST(x, y) abs((x) - (y))
typedef int value_type;
typedef int result_type;
# endif
#elif (DIST_TYPE == 1) // L2Dist
#define DIST(x, y) (((x) - (y)) * ((x) - (y))) #define DIST(x, y) (((x) - (y)) * ((x) - (y)))
#elif (DIST_TYPE == 2) typedef float value_type;
#define DIST(x, y) bit1Count((uint)(x) ^ (uint)(y)) typedef float result_type;
#endif #elif (DIST_TYPE == 2) // Hamming
#define DIST(x, y) bit1Count(((x) ^ (y))
typedef int value_type;
typedef int result_type;
#endif
float reduce_block(__local float *s_query, result_type reduce_block(
__local float *s_train, __local value_type *s_query,
int lidx, __local value_type *s_train,
int lidy int lidx,
) int lidy
)
{ {
float result = 0; result_type result = 0;
#pragma unroll #pragma unroll
for (int j = 0 ; j < BLOCK_SIZE ; j++) for (int j = 0 ; j < BLOCK_SIZE ; j++)
{ {
result += DIST(s_query[lidy * BLOCK_SIZE + j], s_train[j * BLOCK_SIZE + lidx]); result += DIST(
s_query[lidy * BLOCK_SIZE + j],
s_train[j * BLOCK_SIZE + lidx]);
} }
return result; return result;
} }
float reduce_multi_block(__local float *s_query, result_type reduce_multi_block(
__local float *s_train, __local value_type *s_query,
int block_index, __local value_type *s_train,
int lidx, int block_index,
int lidy int lidx,
) int lidy
)
{ {
float result = 0; result_type result = 0;
#pragma unroll #pragma unroll
for (int j = 0 ; j < BLOCK_SIZE ; j++) for (int j = 0 ; j < BLOCK_SIZE ; j++)
{ {
result += DIST(s_query[lidy * MAX_DESC_LEN + block_index * BLOCK_SIZE + j], s_train[j * BLOCK_SIZE + lidx]); result += DIST(
s_query[lidy * MAX_DESC_LEN + block_index * BLOCK_SIZE + j],
s_train[j * BLOCK_SIZE + lidx]);
} }
return result; return result;
} }
@ -117,9 +135,9 @@ float reduce_multi_block(__local float *s_query,
/* 2dim launch, global size: dim0 is (query rows + BLOCK_SIZE - 1) / BLOCK_SIZE * BLOCK_SIZE, dim1 is BLOCK_SIZE /* 2dim launch, global size: dim0 is (query rows + BLOCK_SIZE - 1) / BLOCK_SIZE * BLOCK_SIZE, dim1 is BLOCK_SIZE
local size: dim0 is BLOCK_SIZE, dim1 is BLOCK_SIZE. local size: dim0 is BLOCK_SIZE, dim1 is BLOCK_SIZE.
*/ */
__kernel void BruteForceMatch_UnrollMatch_D5( __kernel void BruteForceMatch_UnrollMatch(
__global float *query, __global T *query,
__global float *train, __global T *train,
//__global float *mask, //__global float *mask,
__global int *bestTrainIdx, __global int *bestTrainIdx,
__global float *bestDistance, __global float *bestDistance,
@ -131,13 +149,12 @@ __kernel void BruteForceMatch_UnrollMatch_D5(
int step int step
) )
{ {
const int lidx = get_local_id(0); const int lidx = get_local_id(0);
const int lidy = get_local_id(1); const int lidy = get_local_id(1);
const int groupidx = get_group_id(0); const int groupidx = get_group_id(0);
__local float *s_query = sharebuffer; __local value_type *s_query = sharebuffer;
__local float *s_train = sharebuffer + BLOCK_SIZE * MAX_DESC_LEN; __local value_type *s_train = sharebuffer + BLOCK_SIZE * MAX_DESC_LEN;
int queryIdx = groupidx * BLOCK_SIZE + lidy; int queryIdx = groupidx * BLOCK_SIZE + lidy;
// load the query into local memory. // load the query into local memory.
@ -155,7 +172,7 @@ __kernel void BruteForceMatch_UnrollMatch_D5(
volatile int imgIdx = 0; volatile int imgIdx = 0;
for (int t = 0, endt = (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE; t < endt; t++) for (int t = 0, endt = (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE; t < endt; t++)
{ {
float result = 0; result_type result = 0;
#pragma unroll #pragma unroll
for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE ; i++) for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE ; i++)
{ {
@ -211,9 +228,9 @@ __kernel void BruteForceMatch_UnrollMatch_D5(
} }
} }
__kernel void BruteForceMatch_Match_D5( __kernel void BruteForceMatch_Match(
__global float *query, __global T *query,
__global float *train, __global T *train,
//__global float *mask, //__global float *mask,
__global int *bestTrainIdx, __global int *bestTrainIdx,
__global float *bestDistance, __global float *bestDistance,
@ -234,14 +251,13 @@ __kernel void BruteForceMatch_Match_D5(
float myBestDistance = MAX_FLOAT; float myBestDistance = MAX_FLOAT;
int myBestTrainIdx = -1; int myBestTrainIdx = -1;
__local float *s_query = sharebuffer; __local value_type *s_query = sharebuffer;
__local float *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE; __local value_type *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
// loop // loop
for (int t = 0 ; t < (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE ; t++) for (int t = 0 ; t < (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE ; t++)
{ {
//Dist dist; result_type result = 0;
float result = 0;
for (int i = 0 ; i < (query_cols + BLOCK_SIZE - 1) / BLOCK_SIZE ; i++) for (int i = 0 ; i < (query_cols + BLOCK_SIZE - 1) / BLOCK_SIZE ; i++)
{ {
const int loadx = lidx + i * BLOCK_SIZE; const int loadx = lidx + i * BLOCK_SIZE;
@ -303,9 +319,9 @@ __kernel void BruteForceMatch_Match_D5(
} }
//radius_unrollmatch //radius_unrollmatch
__kernel void BruteForceMatch_RadiusUnrollMatch_D5( __kernel void BruteForceMatch_RadiusUnrollMatch(
__global float *query, __global T *query,
__global float *train, __global T *train,
float maxDistance, float maxDistance,
//__global float *mask, //__global float *mask,
__global int *bestTrainIdx, __global int *bestTrainIdx,
@ -329,10 +345,10 @@ __kernel void BruteForceMatch_RadiusUnrollMatch_D5(
const int queryIdx = groupidy * BLOCK_SIZE + lidy; const int queryIdx = groupidy * BLOCK_SIZE + lidy;
const int trainIdx = groupidx * BLOCK_SIZE + lidx; const int trainIdx = groupidx * BLOCK_SIZE + lidx;
__local float *s_query = sharebuffer; __local value_type *s_query = sharebuffer;
__local float *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE; __local value_type *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
float result = 0; result_type result = 0;
for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE ; ++i) for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE ; ++i)
{ {
//load a BLOCK_SIZE * BLOCK_SIZE block into local train. //load a BLOCK_SIZE * BLOCK_SIZE block into local train.
@ -363,9 +379,9 @@ __kernel void BruteForceMatch_RadiusUnrollMatch_D5(
} }
//radius_match //radius_match
__kernel void BruteForceMatch_RadiusMatch_D5( __kernel void BruteForceMatch_RadiusMatch(
__global float *query, __global T *query,
__global float *train, __global T *train,
float maxDistance, float maxDistance,
//__global float *mask, //__global float *mask,
__global int *bestTrainIdx, __global int *bestTrainIdx,
@ -389,10 +405,10 @@ __kernel void BruteForceMatch_RadiusMatch_D5(
const int queryIdx = groupidy * BLOCK_SIZE + lidy; const int queryIdx = groupidy * BLOCK_SIZE + lidy;
const int trainIdx = groupidx * BLOCK_SIZE + lidx; const int trainIdx = groupidx * BLOCK_SIZE + lidx;
__local float *s_query = sharebuffer; __local value_type *s_query = sharebuffer;
__local float *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE; __local value_type *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
float result = 0; result_type result = 0;
for (int i = 0 ; i < (query_cols + BLOCK_SIZE - 1) / BLOCK_SIZE ; ++i) for (int i = 0 ; i < (query_cols + BLOCK_SIZE - 1) / BLOCK_SIZE ; ++i)
{ {
//load a BLOCK_SIZE * BLOCK_SIZE block into local train. //load a BLOCK_SIZE * BLOCK_SIZE block into local train.
@ -423,9 +439,9 @@ __kernel void BruteForceMatch_RadiusMatch_D5(
} }
__kernel void BruteForceMatch_knnUnrollMatch_D5( __kernel void BruteForceMatch_knnUnrollMatch(
__global float *query, __global T *query,
__global float *train, __global T *train,
//__global float *mask, //__global float *mask,
__global int2 *bestTrainIdx, __global int2 *bestTrainIdx,
__global float2 *bestDistance, __global float2 *bestDistance,
@ -442,8 +458,8 @@ __kernel void BruteForceMatch_knnUnrollMatch_D5(
const int groupidx = get_group_id(0); const int groupidx = get_group_id(0);
const int queryIdx = groupidx * BLOCK_SIZE + lidy; const int queryIdx = groupidx * BLOCK_SIZE + lidy;
local float *s_query = sharebuffer; local value_type *s_query = sharebuffer;
local float *s_train = sharebuffer + BLOCK_SIZE * MAX_DESC_LEN; local value_type *s_train = sharebuffer + BLOCK_SIZE * MAX_DESC_LEN;
// load the query into local memory. // load the query into local memory.
for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE; i ++) for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE; i ++)
@ -461,7 +477,7 @@ __kernel void BruteForceMatch_knnUnrollMatch_D5(
volatile int imgIdx = 0; volatile int imgIdx = 0;
for (int t = 0 ; t < (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE ; t++) for (int t = 0 ; t < (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE ; t++)
{ {
float result = 0; result_type result = 0;
for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE ; i++) for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE ; i++)
{ {
const int loadX = lidx + i * BLOCK_SIZE; const int loadX = lidx + i * BLOCK_SIZE;
@ -569,9 +585,9 @@ __kernel void BruteForceMatch_knnUnrollMatch_D5(
} }
} }
__kernel void BruteForceMatch_knnMatch_D5( __kernel void BruteForceMatch_knnMatch(
__global float *query, __global T *query,
__global float *train, __global T *train,
//__global float *mask, //__global float *mask,
__global int2 *bestTrainIdx, __global int2 *bestTrainIdx,
__global float2 *bestDistance, __global float2 *bestDistance,
@ -588,8 +604,8 @@ __kernel void BruteForceMatch_knnMatch_D5(
const int groupidx = get_group_id(0); const int groupidx = get_group_id(0);
const int queryIdx = groupidx * BLOCK_SIZE + lidy; const int queryIdx = groupidx * BLOCK_SIZE + lidy;
local float *s_query = sharebuffer; local value_type *s_query = sharebuffer;
local float *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE; local value_type *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
float myBestDistance1 = MAX_FLOAT; float myBestDistance1 = MAX_FLOAT;
float myBestDistance2 = MAX_FLOAT; float myBestDistance2 = MAX_FLOAT;
@ -599,7 +615,7 @@ __kernel void BruteForceMatch_knnMatch_D5(
//loop //loop
for (int t = 0 ; t < (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE ; t++) for (int t = 0 ; t < (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE ; t++)
{ {
float result = 0.0f; result_type result = 0.0f;
for (int i = 0 ; i < (query_cols + BLOCK_SIZE -1) / BLOCK_SIZE ; i++) for (int i = 0 ; i < (query_cols + BLOCK_SIZE -1) / BLOCK_SIZE ; i++)
{ {
const int loadx = lidx + i * BLOCK_SIZE; const int loadx = lidx + i * BLOCK_SIZE;
@ -712,9 +728,9 @@ __kernel void BruteForceMatch_knnMatch_D5(
} }
} }
kernel void BruteForceMatch_calcDistanceUnrolled_D5( kernel void BruteForceMatch_calcDistanceUnrolled(
__global float *query, __global T *query,
__global float *train, __global T *train,
//__global float *mask, //__global float *mask,
__global float *allDist, __global float *allDist,
__local float *sharebuffer, __local float *sharebuffer,
@ -727,9 +743,9 @@ kernel void BruteForceMatch_calcDistanceUnrolled_D5(
/* Todo */ /* Todo */
} }
kernel void BruteForceMatch_calcDistance_D5( kernel void BruteForceMatch_calcDistance(
__global float *query, __global T *query,
__global float *train, __global T *train,
//__global float *mask, //__global float *mask,
__global float *allDist, __global float *allDist,
__local float *sharebuffer, __local float *sharebuffer,
@ -742,7 +758,7 @@ kernel void BruteForceMatch_calcDistance_D5(
/* Todo */ /* Todo */
} }
kernel void BruteForceMatch_findBestMatch_D5( kernel void BruteForceMatch_findBestMatch(
__global float *allDist, __global float *allDist,
__global int *bestTrainIdx, __global int *bestTrainIdx,
__global float *bestDistance, __global float *bestDistance,

Write Preview
Loading…
Cancel
Save