It includes the accuracy/performance test and the implementation of KNN.
--- /dev/null
+/*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);
+}
\ No newline at end of file
--- /dev/null
+/*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);
+}
\ No newline at end of file
--- /dev/null
+/*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!*/
+}
--- /dev/null
+///////////////////////////////////////////////////////////////////////////////////////
+//
+// 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
\ No newline at end of file
projects / profile / ivi / opencv.git / commitdiff