From: Tom Stellard <[email protected]> This enables piglit to run and interpret the results from OpenCV's gtest based opencv_test_ocl program.
This patch adds two new CMake configuration variables: OPENCL_OpenCVTestOCL_BINDIR: You can use this variable to enable the OpenCV tests by setting it to the full path of the bin directory in your opencv build tree (e.g. /home/user/opencv/build/bin). OPENCL_OpenCVTestOCL_WORKDIR: (Optional)If you don't want to use OpenCV's default work directory, you can use this variable to specify a different one. v2: - Python code cleanups v3: - More cleanups - Move opencv.py into framework - Use ConfigParser --- framework/opencv.py | 75 ++++++++++++++++++++ piglit.conf.example | 5 ++ tests/all_cl.py | 4 ++ tests/cl/program/execute/opencv-merge-hist.cl | 98 --------------------------- 4 files changed, 84 insertions(+), 98 deletions(-) create mode 100644 framework/opencv.py create mode 100644 piglit.conf.example delete mode 100644 tests/cl/program/execute/opencv-merge-hist.cl diff --git a/framework/opencv.py b/framework/opencv.py new file mode 100644 index 0000000..2dc98bb --- /dev/null +++ b/framework/opencv.py @@ -0,0 +1,75 @@ +#!/usr/bin/env python +# +# Copyright 2014 Advanced Micro Devices, Inc. +# +# Permission is hereby granted, free of charge, to any person obtaining a +# copy of this software and associated documentation files (the "Software"), +# to deal in the Software without restriction, including without limitation +# the rights to use, copy, modify, merge, publish, distribute, sublicense, +# and/or sell copies of the Software, and to permit persons to whom the +# Software is furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice (including the next +# paragraph) shall be included in all copies or substantial portions of the +# Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +# THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +# SOFTWARE. +# +# Authors: Tom Stellard <[email protected]> +# + +import re +import subprocess +from os import path +from framework.gtest import GTest +from framework.core import PIGLIT_CONFIG + + +class OpenCVTest(GTest): + def __init__(self, test_prog, testname): + options = [test_prog, '--gtest_filter=' + testname, '--gtest_color=no'] + if PIGLIT_CONFIG.has_option('opencv', 'workdir'): + options.append('-w {}'.format(PIGLIT_CONFIG.get('opencv', 'workdir'))) + GTest.__init__(self, options) + + +def add_opencv_tests(profile, individual = False): + if not PIGLIT_CONFIG.has_option('opencv', 'opencv_test_ocl_bindir'): + return + + opencv_test_ocl = path.join(PIGLIT_CONFIG.get('opencv', + 'opencv_test_ocl_bindir'), 'opencv_test_ocl') + if not path.isfile(opencv_test_ocl): + print 'Warning: ', opencv_test_ocl, 'does not exist.\nSkipping OpenCV tests...' + return + tests = subprocess.check_output([opencv_test_ocl, '--gtest_list_tests']) + test_list = tests.splitlines() + group_name = '' + full_test_name = '' + for line in test_list: + #Test groups names start at the beginning of the line and end with '.' + m = re.match('([^.]+\.)$', line) + if m: + group_name = m.group(1) + group_desc = group_name[:-1] + full_test_name = 'opencv/{}'.format(group_desc) + if not individual: + profile.tests[full_test_name] = OpenCVTest(opencv_test_ocl, + '{}*'.format(group_name)) + continue + + if not individual: + continue + + # Test names are indent by 2 spaces + m = re.match(' (.+)', line) + if m: + test_name = m.group(1) + profile.tests['{}/{}'.format(full_test_name,test_name)] = \ + OpenCVTest(opencv_test_ocl, '{}{}'.format(group_name ,test_name)) diff --git a/piglit.conf.example b/piglit.conf.example new file mode 100644 index 0000000..72ee1ef --- /dev/null +++ b/piglit.conf.example @@ -0,0 +1,5 @@ +;[opencv] +; Set the opencv_test_ocl_bindir variable to run the OpenCV OpenCL tests. +;opencv_test_ocl_bindir=/home/user/opencv/build/bin +;opencv_workdir=/home/user/opencv/samples/c/ +; diff --git a/tests/all_cl.py b/tests/all_cl.py index a7d7106..d45036a 100644 --- a/tests/all_cl.py +++ b/tests/all_cl.py @@ -7,6 +7,8 @@ __all__ = ['profile'] import os import os.path as path +from framework.opencv import add_opencv_tests + from framework.core import Group, TestProfile from framework.exectest import PlainExecTest @@ -122,3 +124,5 @@ add_program_test_dir(program_execute_builtin, 'generated_tests/cl/builtin/relati program_execute_store = Group() program["Execute"]["Store"] = program_execute_store add_program_test_dir(program_execute_store, 'generated_tests/cl/store') + +add_opencv_tests(profile) diff --git a/tests/cl/program/execute/opencv-merge-hist.cl b/tests/cl/program/execute/opencv-merge-hist.cl deleted file mode 100644 index 8dccf24..0000000 --- a/tests/cl/program/execute/opencv-merge-hist.cl +++ /dev/null @@ -1,98 +0,0 @@ -/*! -[config] -name: OpenCV merge-hist -kernel_name: merge_hist -dimensions: 1 -global_size: 65536 1 1 -local_size: 256 1 1 - -# This test checks that the loop: -# for(int stride = HISTOGRAM256_WORK_GROUP_SIZE /2; stride > 0; stride >>= 1) -# in this kernel is unrolled correctly or not unrolled at all. The Clang -# OpenCL frontend was unrolling this in way that created illegal uses of the -# barrier() builtin which resulted in GPU hangs on r600g. -[test] -arg_in: 0 buffer int[65536] repeat 0 -arg_out: 1 buffer int[256] repeat 0 -arg_in: 2 int 256 - -!*/ - -// The kernel in this test is taken from the opencv library (opencv.org) -// File: modules/ocl/src/opencl/imgproc_histogram.cl -// -// License Agreement -// For Open Source Computer Vision Library -// -// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// 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 -// Niko Li, [email protected] -// Jia Haipeng, [email protected] -// Xu Pang, [email protected] -// Wenju He, [email protected] -// 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 GpuMaterials 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. -// -// - - -#define PARTIAL_HISTOGRAM256_COUNT (256) -#define HISTOGRAM256_BIN_COUNT (256) - -#define HISTOGRAM256_WORK_GROUP_SIZE (256) -#define HISTOGRAM256_LOCAL_MEM_SIZE (HISTOGRAM256_BIN_COUNT) - -#define NBANKS (16) -#define NBANKS_BIT (4) - - -__kernel __attribute__((reqd_work_group_size(256,1,1)))void merge_hist(__global int* buf, - __global int* hist, - int src_step) -{ - int lx = get_local_id(0); - int gx = get_group_id(0); - - int sum = 0; - - for(int i = lx; i < PARTIAL_HISTOGRAM256_COUNT; i += HISTOGRAM256_WORK_GROUP_SIZE) - sum += buf[ mad24(i, src_step, gx)]; - - __local int data[HISTOGRAM256_WORK_GROUP_SIZE]; - data[lx] = sum; - - for(int stride = HISTOGRAM256_WORK_GROUP_SIZE /2; stride > 0; stride >>= 1) - { - barrier(CLK_LOCAL_MEM_FENCE); - if(lx < stride) - data[lx] += data[lx + stride]; - } - - if(lx == 0) - hist[gx] = data[0]; -} -- 1.8.1.4 _______________________________________________ Piglit mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/piglit
