From 1bfe39f485ba0b5e00c2d9d64a2e3ed61e1f67d4 Mon Sep 17 00:00:00 2001 From: Jin Ma Date: Sun, 22 Sep 2013 10:23:54 +0800 Subject: [PATCH] Added knearest neighbor of OpenCL version. It includes the accuracy/performance test and the implementation of KNN. --- modules/ocl/perf/perf_ml.cpp | 109 +++++++++++++++++ modules/ocl/src/knearest.cpp | 163 +++++++++++++++++++++++++ modules/ocl/src/opencl/knearest.cl | 186 +++++++++++++++++++++++++++++ modules/ocl/test/test_ml.cpp | 124 +++++++++++++++++++ 4 files changed, 582 insertions(+) create mode 100644 modules/ocl/perf/perf_ml.cpp create mode 100644 modules/ocl/src/knearest.cpp create mode 100644 modules/ocl/src/opencl/knearest.cl create mode 100644 modules/ocl/test/test_ml.cpp diff --git a/modules/ocl/perf/perf_ml.cpp b/modules/ocl/perf/perf_ml.cpp new file mode 100644 index 000000000..fac471ed4 --- /dev/null +++ b/modules/ocl/perf/perf_ml.cpp @@ -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 KNNParamType; +typedef TestBaseWithParam 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 diff --git a/modules/ocl/src/knearest.cpp b/modules/ocl/src/knearest.cpp new file mode 100644 index 000000000..4f78e85ea --- /dev/null +++ b/modules/ocl/src/knearest.cpp @@ -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(pos) = t1[j]; + } + + Point pos_label(CvKNearest::var_count, i); + samples_mat.at(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 > 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 diff --git a/modules/ocl/src/opencl/knearest.cl b/modules/ocl/src/opencl/knearest.cl new file mode 100644 index 000000000..47af57a7e --- /dev/null +++ b/modules/ocl/src/opencl/knearest.cl @@ -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!*/ +} diff --git a/modules/ocl/test/test_ml.cpp b/modules/ocl/test/test_ml.cpp new file mode 100644 index 000000000..834fc4e37 --- /dev/null +++ b/modules/ocl/test/test_ml.cpp @@ -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 \ No newline at end of file