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*/ >
void matchUnrolledCached(const oclMat &query, const oclMat &train, const oclMat &/*mask*/,
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);
}
}
sprintf(opt,
"-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d",
T_ARR[query.depth()], distType, block_size);
-
if(globalSize[0] != 0)
{
args.push_back( make_pair( sizeof(cl_mem), (void *)&query.data ));
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);
}
}
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);
}
}
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);
}
}
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);
}
}
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);
}
}
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);
}
}
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);
}
}
//args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols ));
//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);
}
}
int bit1Count(int x)
{
int c = 0;
- int ix = (int)x;
+ int ix = x;
for (int i = 0 ; i < 32 ; i++)
{
c += ix & 0x1;
return c;
}
-#if (DIST_TYPE == 0)
-#define DIST(x, y) fabs((x) - (y))
-#elif (DIST_TYPE == 1)
+// dirty fix for non-template support
+#if (DIST_TYPE == 0) // L1Dist
+# 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)))
-#elif (DIST_TYPE == 2)
-#define DIST(x, y) bit1Count((uint)(x) ^ (uint)(y))
-#endif
-
+typedef float value_type;
+typedef float result_type;
+#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,
- __local float *s_train,
- int lidx,
- int lidy
- )
+result_type reduce_block(
+ __local value_type *s_query,
+ __local value_type *s_train,
+ int lidx,
+ int lidy
+ )
{
- float result = 0;
+ result_type result = 0;
#pragma unroll
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;
}
-float reduce_multi_block(__local float *s_query,
- __local float *s_train,
- int block_index,
- int lidx,
- int lidy
- )
+result_type reduce_multi_block(
+ __local value_type *s_query,
+ __local value_type *s_train,
+ int block_index,
+ int lidx,
+ int lidy
+ )
{
- float result = 0;
+ result_type result = 0;
#pragma unroll
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;
}
/* 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.
*/
-__kernel void BruteForceMatch_UnrollMatch_D5(
- __global float *query,
- __global float *train,
+__kernel void BruteForceMatch_UnrollMatch(
+ __global T *query,
+ __global T *train,
//__global float *mask,
__global int *bestTrainIdx,
__global float *bestDistance,
int step
)
{
-
const int lidx = get_local_id(0);
const int lidy = get_local_id(1);
const int groupidx = get_group_id(0);
- __local float *s_query = sharebuffer;
- __local float *s_train = sharebuffer + BLOCK_SIZE * MAX_DESC_LEN;
+ __local value_type *s_query = sharebuffer;
+ __local value_type *s_train = sharebuffer + BLOCK_SIZE * MAX_DESC_LEN;
int queryIdx = groupidx * BLOCK_SIZE + lidy;
// load the query into local memory.
volatile int imgIdx = 0;
for (int t = 0, endt = (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE; t < endt; t++)
{
- float result = 0;
+ result_type result = 0;
#pragma unroll
for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE ; i++)
{
}
}
-__kernel void BruteForceMatch_Match_D5(
- __global float *query,
- __global float *train,
+__kernel void BruteForceMatch_Match(
+ __global T *query,
+ __global T *train,
//__global float *mask,
__global int *bestTrainIdx,
__global float *bestDistance,
float myBestDistance = MAX_FLOAT;
int myBestTrainIdx = -1;
- __local float *s_query = sharebuffer;
- __local float *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
+ __local value_type *s_query = sharebuffer;
+ __local value_type *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
// loop
for (int t = 0 ; t < (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE ; t++)
{
- //Dist dist;
- float result = 0;
+ result_type result = 0;
for (int i = 0 ; i < (query_cols + BLOCK_SIZE - 1) / BLOCK_SIZE ; i++)
{
const int loadx = lidx + i * BLOCK_SIZE;
}
//radius_unrollmatch
-__kernel void BruteForceMatch_RadiusUnrollMatch_D5(
- __global float *query,
- __global float *train,
+__kernel void BruteForceMatch_RadiusUnrollMatch(
+ __global T *query,
+ __global T *train,
float maxDistance,
//__global float *mask,
__global int *bestTrainIdx,
const int queryIdx = groupidy * BLOCK_SIZE + lidy;
const int trainIdx = groupidx * BLOCK_SIZE + lidx;
- __local float *s_query = sharebuffer;
- __local float *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
+ __local value_type *s_query = sharebuffer;
+ __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)
{
//load a BLOCK_SIZE * BLOCK_SIZE block into local train.
}
//radius_match
-__kernel void BruteForceMatch_RadiusMatch_D5(
- __global float *query,
- __global float *train,
+__kernel void BruteForceMatch_RadiusMatch(
+ __global T *query,
+ __global T *train,
float maxDistance,
//__global float *mask,
__global int *bestTrainIdx,
const int queryIdx = groupidy * BLOCK_SIZE + lidy;
const int trainIdx = groupidx * BLOCK_SIZE + lidx;
- __local float *s_query = sharebuffer;
- __local float *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
+ __local value_type *s_query = sharebuffer;
+ __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)
{
//load a BLOCK_SIZE * BLOCK_SIZE block into local train.
}
-__kernel void BruteForceMatch_knnUnrollMatch_D5(
- __global float *query,
- __global float *train,
+__kernel void BruteForceMatch_knnUnrollMatch(
+ __global T *query,
+ __global T *train,
//__global float *mask,
__global int2 *bestTrainIdx,
__global float2 *bestDistance,
const int groupidx = get_group_id(0);
const int queryIdx = groupidx * BLOCK_SIZE + lidy;
- local float *s_query = sharebuffer;
- local float *s_train = sharebuffer + BLOCK_SIZE * MAX_DESC_LEN;
+ local value_type *s_query = sharebuffer;
+ local value_type *s_train = sharebuffer + BLOCK_SIZE * MAX_DESC_LEN;
// load the query into local memory.
for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE; i ++)
volatile int imgIdx = 0;
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++)
{
const int loadX = lidx + i * BLOCK_SIZE;
}
}
-__kernel void BruteForceMatch_knnMatch_D5(
- __global float *query,
- __global float *train,
+__kernel void BruteForceMatch_knnMatch(
+ __global T *query,
+ __global T *train,
//__global float *mask,
__global int2 *bestTrainIdx,
__global float2 *bestDistance,
const int groupidx = get_group_id(0);
const int queryIdx = groupidx * BLOCK_SIZE + lidy;
- local float *s_query = sharebuffer;
- local float *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
+ local value_type *s_query = sharebuffer;
+ local value_type *s_train = sharebuffer + BLOCK_SIZE * BLOCK_SIZE;
float myBestDistance1 = MAX_FLOAT;
float myBestDistance2 = MAX_FLOAT;
//loop
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++)
{
const int loadx = lidx + i * BLOCK_SIZE;
}
}
-kernel void BruteForceMatch_calcDistanceUnrolled_D5(
- __global float *query,
- __global float *train,
+kernel void BruteForceMatch_calcDistanceUnrolled(
+ __global T *query,
+ __global T *train,
//__global float *mask,
__global float *allDist,
__local float *sharebuffer,
/* Todo */
}
-kernel void BruteForceMatch_calcDistance_D5(
- __global float *query,
- __global float *train,
+kernel void BruteForceMatch_calcDistance(
+ __global T *query,
+ __global T *train,
//__global float *mask,
__global float *allDist,
__local float *sharebuffer,
/* Todo */
}
-kernel void BruteForceMatch_findBestMatch_D5(
+kernel void BruteForceMatch_findBestMatch(
__global float *allDist,
__global int *bestTrainIdx,
__global float *bestDistance,
projects / platform / upstream / opencv.git / commitdiff