[Piglit] [PATCH 2/3] cl: Add support for OpenCV unit tests v3
Tom Stellard
tom at stellard.net
Wed Jan 22 10:19:28 PST 2014
From: Tom Stellard <thomas.stellard at amd.com>
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 <thomas.stellard at amd.com>
+#
+
+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, newlife20080214 at gmail.com
-// Jia Haipeng, jiahaipeng95 at gmail.com
-// Xu Pang, pangxu010 at 163.com
-// Wenju He, wenju at 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 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
More information about the Piglit
mailing list