[Mesa-dev] [PATCH] cl: Add support for OpenCV unit tests v2
Tom Stellard
tom at stellard.net
Fri Jan 17 14:30:13 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
XXX: opencv
---
CMakeLists.txt | 9 ++-
tests/all_cl.py | 4 ++
tests/cl/program/execute/opencv-merge-hist.cl | 98 ---------------------------
tests/ocl_config.py.in | 31 +++++++++
tests/opencv.py | 76 +++++++++++++++++++++
5 files changed, 119 insertions(+), 99 deletions(-)
delete mode 100644 tests/cl/program/execute/opencv-merge-hist.cl
create mode 100644 tests/ocl_config.py.in
create mode 100644 tests/opencv.py
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 1befffb..37eac65 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -92,7 +92,10 @@ if(PIGLIT_BUILD_GLES3_TESTS AND NOT PIGLIT_USE_WAFFLE)
endif(PIGLIT_BUILD_GLES3_TESTS AND NOT PIGLIT_USE_WAFFLE)
if(PIGLIT_BUILD_CL_TESTS)
- find_package(OpenCL REQUIRED)
+ find_package(OpenCL REQUIRED)
+ set(OPENCL_OpenCVTestOCL_BINDIR "" CACHE STRING "Directory with the opencv_ocl_test program" )
+ set(OPENCL_OpenCVTestOCL_WORKDIR "" CACHE STRING
+ "Directory with the opencv test resources (This is not needed if your OpenCV build directory was $(opencv_src_dir)/build )" )
endif(PIGLIT_BUILD_CL_TESTS)
IF(${CMAKE_SYSTEM_NAME} MATCHES "Linux")
@@ -356,6 +359,10 @@ configure_file(
"${piglit_SOURCE_DIR}/tests/util/config.h.in"
"${piglit_BINARY_DIR}/tests/util/config.h"
)
+configure_file(
+ "${piglit_SOURCE_DIR}/tests/ocl_config.py.in"
+ "${piglit_BINARY_DIR}/tests/ocl_config.py"
+)
include(cmake/piglit_util.cmake)
include(cmake/piglit_glapi.cmake)
diff --git a/tests/all_cl.py b/tests/all_cl.py
index a7d7106..b9c112a 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 tests.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];
-}
diff --git a/tests/ocl_config.py.in b/tests/ocl_config.py.in
new file mode 100644
index 0000000..7ba515c
--- /dev/null
+++ b/tests/ocl_config.py.in
@@ -0,0 +1,31 @@
+#!/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>
+#
+
+# TODO: Move this into a global piglit config file and use python's
+# ConfigParser
+
+OPENCV_TEST_OCL_BINDIR = '${OPENCL_OpenCVTestOCL_BINDIR}'
+OPENCV_TEST_OCL_WORKDIR = '${OPENCL_OpenCVTestOCL_WORKDIR}'
diff --git a/tests/opencv.py b/tests/opencv.py
new file mode 100644
index 0000000..746a25b
--- /dev/null
+++ b/tests/opencv.py
@@ -0,0 +1,76 @@
+#!/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 os
+import sys
+import re
+import string
+import subprocess
+from os import path
+from framework.exectest import ExecTest
+from framework.core import TestProfile, testBinDir, Group
+from framework.gtest import GTest
+from tests.ocl_config import opencv_test_ocl_bin, opencv_test_ocl_workdir
+
+opencv_test_ocl = path.join(opencv_test_ocl_bin, 'opencv_test_ocl')
+
+class OpenCVTest(GTest):
+ def __init__(self, testname):
+ options = [opencv_test_ocl, '--gtest_filter=' + testname, '--gtest_color=no']
+ if opencv_test_ocl_workdir != '':
+ options.append('-w {}'.format(opencv_test_ocl_workdir))
+ ExecTest.__init__(self, options)
+
+
+def add_opencv_tests(profile, individual = False):
+ if opencv_test_ocl_bin == '':
+ return
+
+ opencv = Group()
+ 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('{}*'.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('{}{}'.format(group_name ,test_name))
--
1.8.1.4
More information about the mesa-dev
mailing list