[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