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

Merge pull request #1489 from melody-rain:2_4_knearest_neighbor

Browse Source
This commit is contained in:
Alexander Smorkalov
2013-09-23 13:34:38 +04:00
committed by OpenCV Buildbot
parent 0cd3d1f4d6 b88f5302e2
commit d6d8e96003
8 changed files with 605 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
modules/ocl/src/initialization.cpp
View File

@@ -1033,6 +1033,14 @@ namespace cv
return impl->maxComputeUnits;
}
unsigned long queryLocalMemInfo()
{
Info::Impl* impl = Context::getContext()->impl;
cl_ulong local_memory_size = 0;
clGetDeviceInfo(impl->devices[impl->devnum], CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), (void*)&local_memory_size, 0);
return local_memory_size;
}
void* Context::oclContext()
{
return impl->oclcontext;

157
modules/ocl/src/knearest.cpp Normal file
View File

@@ -0,0 +1,157 @@
/*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) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
// Jin Ma, jin@multicorewareinc.com
//
// 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 oclMaterials 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*/
#include "precomp.hpp"
using namespace cv;
using namespace cv::ocl;
namespace cv
{
namespace ocl
{
extern const char* knearest;//knearest
}
}
KNearestNeighbour::KNearestNeighbour()
{
clear();
}
KNearestNeighbour::~KNearestNeighbour()
{
clear();
samples_ocl.release();
}
void KNearestNeighbour::clear()
{
CvKNearest::clear();
}
bool KNearestNeighbour::train(const Mat& trainData, Mat& labels, Mat& sampleIdx,
bool isRegression, int _max_k, bool updateBase)
{
max_k = _max_k;
bool cv_knn_train = CvKNearest::train(trainData, labels, sampleIdx, isRegression, max_k, updateBase);
CvVectors* s = CvKNearest::samples;
cv::Mat samples_mat(s->count, CvKNearest::var_count + 1, s->type);
float* s1 = (float*)(s + 1);
for(int i = 0; i < s->count; i++)
{
float* t1 = s->data.fl[i];
for(int j = 0; j < CvKNearest::var_count; j++)
{
Point pos(j, i);
samples_mat.at<float>(pos) = t1[j];
}
Point pos_label(CvKNearest::var_count, i);
samples_mat.at<float>(pos_label) = s1[i];
}
samples_ocl = samples_mat;
return cv_knn_train;
}
void KNearestNeighbour::find_nearest(const oclMat& samples, int k, oclMat& lables)
{
CV_Assert(!samples_ocl.empty());
lables.create(samples.rows, 1, CV_32FC1);
CV_Assert(samples.cols == CvKNearest::var_count);
CV_Assert(samples.type() == CV_32FC1);
CV_Assert(k >= 1 && k <= max_k);
int k1 = KNearest::get_sample_count();
k1 = MIN( k1, k );
String kernel_name = "knn_find_nearest";
cl_ulong local_memory_size = queryLocalMemInfo();
int nThreads = local_memory_size / (2 * k * 4);
if(nThreads >= 256)
nThreads = 256;
int smem_size = nThreads * k * 4 * 2;
size_t local_thread[] = {1, nThreads, 1};
size_t global_thread[] = {1, samples.rows, 1};
char build_option[50];
if(!Context::getContext()->supportsFeature(Context::CL_DOUBLE))
{
sprintf(build_option, " ");
}else
sprintf(build_option, "-D DOUBLE_SUPPORT");
std::vector< std::pair<size_t, const void*> > args;
int samples_ocl_step = samples_ocl.step/samples_ocl.elemSize();
int samples_step = samples.step/samples.elemSize();
int lables_step = lables.step/lables.elemSize();
int _regression = 0;
if(CvKNearest::regression)
_regression = 1;
args.push_back(make_pair(sizeof(cl_mem), (void*)&samples.data));
args.push_back(make_pair(sizeof(cl_int), (void*)&samples.rows));
args.push_back(make_pair(sizeof(cl_int), (void*)&samples.cols));
args.push_back(make_pair(sizeof(cl_int), (void*)&samples_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&k));
args.push_back(make_pair(sizeof(cl_mem), (void*)&samples_ocl.data));
args.push_back(make_pair(sizeof(cl_int), (void*)&samples_ocl.rows));
args.push_back(make_pair(sizeof(cl_int), (void*)&samples_ocl_step));
args.push_back(make_pair(sizeof(cl_mem), (void*)&lables.data));
args.push_back(make_pair(sizeof(cl_int), (void*)&lables_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&_regression));
args.push_back(make_pair(sizeof(cl_int), (void*)&k1));
args.push_back(make_pair(sizeof(cl_int), (void*)&samples_ocl.cols));
args.push_back(make_pair(sizeof(cl_int), (void*)&nThreads));
args.push_back(make_pair(smem_size, (void*)NULL));
openCLExecuteKernel(Context::getContext(), &knearest, kernel_name, global_thread, local_thread, args, -1, -1, build_option);
}

186
modules/ocl/src/opencl/knearest.cl Normal file
View File

@@ -0,0 +1,186 @@
/*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) 2010-2012, Multicoreware, Inc., all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
// Jin Ma, jin@multicorewareinc.com
//
// 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 oclMaterials 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 (DOUBLE_SUPPORT)
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#define TYPE double
#else
#define TYPE float
#endif
#define CV_SWAP(a,b,t) ((t) = (a), (a) = (b), (b) = (t))
///////////////////////////////////// find_nearest //////////////////////////////////////
__kernel void knn_find_nearest(__global float* sample, int sample_row, int sample_col, int sample_step,
int k, __global float* samples_ocl, int sample_ocl_row, int sample_ocl_step,
__global float* _results, int _results_step, int _regression, int K1,
int sample_ocl_col, int nThreads, __local float* nr)
{
int k1 = 0;
int k2 = 0;
bool regression = false;
if(_regression)
regression = true;
TYPE inv_scale;
#ifdef DOUBLE_SUPPORT
inv_scale = 1.0/K1;
#else
inv_scale = 1.0f/K1;
#endif
int y = get_global_id(1);
int j, j1;
int threadY = (y % nThreads);
__local float* dd = nr + nThreads * k;
if(y >= sample_row)
{
return;
}
for(j = 0; j < sample_ocl_row; j++)
{
TYPE sum;
#ifdef DOUBLE_SUPPORT
sum = 0.0;
#else
sum = 0.0f;
#endif
float si;
int t, ii, ii1;
for(t = 0; t < sample_col - 16; t += 16)
{
float16 t0 = vload16(0, sample + y * sample_step + t) - vload16(0, samples_ocl + j * sample_ocl_step + t);
t0 *= t0;
sum += t0.s0 + t0.s1 + t0.s2 + t0.s3 + t0.s4 + t0.s5 + t0.s6 + t0.s7 +
t0.s8 + t0.s9 + t0.sa + t0.sb + t0.sc + t0.sd + t0.se + t0.sf;
}
for(; t < sample_col; t++)
{
#ifdef DOUBLE_SUPPORT
double t0 = sample[y * sample_step + t] - samples_ocl[j * sample_ocl_step + t];
#else
float t0 = sample[y * sample_step + t] - samples_ocl[j * sample_ocl_step + t];
#endif
sum = sum + t0 * t0;
}
si = (float)sum;
for(ii = k1 - 1; ii >= 0; ii--)
{
if(as_int(si) > as_int(dd[ii * nThreads + threadY]))
break;
}
if(ii < k - 1)
{
for(ii1 = k2 - 1; ii1 > ii; ii1--)
{
dd[(ii1 + 1) * nThreads + threadY] = dd[ii1 * nThreads + threadY];
nr[(ii1 + 1) * nThreads + threadY] = nr[ii1 * nThreads + threadY];
}
dd[(ii + 1) * nThreads + threadY] = si;
nr[(ii + 1) * nThreads + threadY] = samples_ocl[sample_col + j * sample_ocl_step];
}
k1 = (k1 + 1) < k ? (k1 + 1) : k;
k2 = k1 < (k - 1) ? k1 : (k - 1);
}
/*! find_nearest_neighbor done!*/
/*! write_results start!*/
switch (regression)
{
case true:
{
TYPE s;
#ifdef DOUBLE_SUPPORT
s = 0.0;
#else
s = 0.0f;
#endif
for(j = 0; j < K1; j++)
s += nr[j * nThreads + threadY];
_results[y * _results_step] = (float)(s * inv_scale);
}
break;
case false:
{
int prev_start = 0, best_count = 0, cur_count;
float best_val;
for(j = K1 - 1; j > 0; j--)
{
bool swap_f1 = false;
for(j1 = 0; j1 < j; j1++)
{
if(nr[j1 * nThreads + threadY] > nr[(j1 + 1) * nThreads + threadY])
{
int t;
CV_SWAP(nr[j1 * nThreads + threadY], nr[(j1 + 1) * nThreads + threadY], t);
swap_f1 = true;
}
}
if(!swap_f1)
break;
}
best_val = 0;
for(j = 1; j <= K1; j++)
if(j == K1 || nr[j * nThreads + threadY] != nr[(j - 1) * nThreads + threadY])
{
cur_count = j - prev_start;
if(best_count < cur_count)
{
best_count = cur_count;
best_val = nr[(j - 1) * nThreads + threadY];
}
prev_start = j;
}
_results[y * _results_step] = best_val;
}
break;
}
///*! write_results done!*/
}