Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

Added knearest neighbor of OpenCL version.

Browse Source 
It includes the accuracy/performance test and the implementation of KNN.
pull/1489/head
Jin Ma 11 years ago
parent 06c33df307
commit 1bfe39f485
4 changed files with 582 additions and 0 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 109
      modules/ocl/perf/perf_ml.cpp
  2. 163
      modules/ocl/src/knearest.cpp
  3. 186
      modules/ocl/src/opencl/knearest.cl
  4. 124
      modules/ocl/test/test_ml.cpp

109
modules/ocl/perf/perf_ml.cpp
Unescape View File

@ -0,0 +1,109 @@
/*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
// Xiaopeng Fu, fuxiaopeng2222@163.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 "perf_precomp.hpp"
using namespace perf;
using namespace std;
using namespace cv::ocl;
using namespace cv;
using std::tr1::tuple;
using std::tr1::get;
////////////////////////////////// K-NEAREST NEIGHBOR ////////////////////////////////////
static void genData(Mat& trainData, Size size, Mat& trainLabel = Mat().setTo(Scalar::all(0)), int nClasses = 0)
{
trainData.create(size, CV_32FC1);
randu(trainData, 1.0, 100.0);
if(nClasses != 0)
{
trainLabel.create(size.height, 1, CV_8UC1);
randu(trainLabel, 0, nClasses - 1);
trainLabel.convertTo(trainLabel, CV_32FC1);
}
}
typedef tuple<int> KNNParamType;
typedef TestBaseWithParam<KNNParamType> KNNFixture;
PERF_TEST_P(KNNFixture, KNN,
testing::Values(1000, 2000, 4000))
{
KNNParamType params = GetParam();
const int rows = get<0>(params);
int columns = 100;
int k = rows/250;
Mat trainData, trainLabels;
Size size(columns, rows);
genData(trainData, size, trainLabels, 3);
Mat testData;
genData(testData, size);
Mat best_label;
if(RUN_PLAIN_IMPL)
{
TEST_CYCLE()
{
CvKNearest knn_cpu;
knn_cpu.train(trainData, trainLabels);
knn_cpu.find_nearest(testData, k, &best_label);
}
}else if(RUN_OCL_IMPL)
{
cv::ocl::oclMat best_label_ocl;
cv::ocl::oclMat testdata;
testdata.upload(testData);
OCL_TEST_CYCLE()
{
cv::ocl::KNearestNeighbour knn_ocl;
knn_ocl.train(trainData, trainLabels);
knn_ocl.find_nearest(testdata, k, best_label_ocl);
}
best_label_ocl.download(best_label);
}else
OCL_PERF_ELSE
SANITY_CHECK(best_label);
}

163
modules/ocl/src/knearest.cpp
Unescape View File

@ -0,0 +1,163 @@
/*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();
}
KNearestNeighbour::KNearestNeighbour(const Mat& train_data, const Mat& responses,
const Mat& sample_idx, bool is_regression, int max_k)
{
max_k = max_k;
CvKNearest::train(train_data, responses, sample_idx, is_regression, max_k);
}
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
Unescape 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!*/
}

124
modules/ocl/test/test_ml.cpp
Unescape View File

@ -0,0 +1,124 @@
///////////////////////////////////////////////////////////////////////////////////////
//
// 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
// Xiaopeng Fu, fuxiaopeng2222@163.com
// Erping Pang, pang_er_ping@163.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 "test_precomp.hpp"
#ifdef HAVE_OPENCL
using namespace cv;
using namespace cv::ocl;
using namespace cvtest;
using namespace testing;
///////K-NEAREST NEIGHBOR//////////////////////////
static void genTrainData(Mat& trainData, int trainDataRow, int trainDataCol,
Mat& trainLabel = Mat().setTo(Scalar::all(0)), int nClasses = 0)
{
cv::RNG &rng = TS::ptr()->get_rng();
cv::Size size(trainDataCol, trainDataRow);
trainData = randomMat(rng, size, CV_32FC1, 1.0, 1000.0, false);
if(nClasses != 0)
{
cv::Size size1(trainDataRow, 1);
trainLabel = randomMat(rng, size1, CV_8UC1, 0, nClasses - 1, false);
trainLabel.convertTo(trainLabel, CV_32FC1);
}
}
PARAM_TEST_CASE(KNN, int, Size, int, bool)
{
int k;
int trainDataCol;
int testDataRow;
int nClass;
bool regression;
virtual void SetUp()
{
k = GET_PARAM(0);
nClass = GET_PARAM(2);
trainDataCol = GET_PARAM(1).width;
testDataRow = GET_PARAM(1).height;
regression = GET_PARAM(3);
}
};
TEST_P(KNN, Accuracy)
{
Mat trainData, trainLabels;
const int trainDataRow = 500;
genTrainData(trainData, trainDataRow, trainDataCol, trainLabels, nClass);
Mat testData, testLabels;
genTrainData(testData, testDataRow, trainDataCol);
KNearestNeighbour knn_ocl;
CvKNearest knn_cpu;
Mat best_label_cpu;
oclMat best_label_ocl;
/*ocl k-Nearest_Neighbor start*/
oclMat trainData_ocl;
trainData_ocl.upload(trainData);
Mat simpleIdx;
knn_ocl.train(trainData, trainLabels, simpleIdx, regression);
oclMat testdata;
testdata.upload(testData);
knn_ocl.find_nearest(testdata, k, best_label_ocl);
/*ocl k-Nearest_Neighbor end*/
/*cpu k-Nearest_Neighbor start*/
knn_cpu.train(trainData, trainLabels, simpleIdx, regression);
knn_cpu.find_nearest(testData, k, &best_label_cpu);
/*cpu k-Nearest_Neighbor end*/
if(regression)
{
EXPECT_MAT_SIMILAR(Mat(best_label_ocl), best_label_cpu, 1e-5);
}
else
{
EXPECT_MAT_NEAR(Mat(best_label_ocl), best_label_cpu, 0.0);
}
}
INSTANTIATE_TEST_CASE_P(OCL_ML, KNN, Combine(Values(6, 5), Values(Size(200, 400), Size(300, 600)),
Values(4, 3), Values(false, true)));
#endif // HAVE_OPENCL
Write Preview
Loading…
Cancel
Save