Commit 4a13528a authored by Tom Stellard's avatar Tom Stellard

cl: Add support for OpenCV unit tests v3

This enables piglit to run and interpret the results from OpenCV's
gtest based opencv_test_ocl program.

You can enable the OpenCV tests by adding the path to the
opencv_test_ocl program to your piglit.conf file.

You can also optionally specify the path to OpenCV's work
directory if you buid OpenCV in a non-standard way.  For example:

[opencv]
opencv_test_ocl_bindir=/home/user/opencv/build/bin
opencv_workdir=/home/user/opencv/samples/c/

v2:
  - Python code cleanups

v3:
  - More cleanups
  - Move opencv.py into framework
  - Use ConfigParser
Reviewed-by: Dylan Baker's avatarDylan Baker <baker.dylan.c@gmail.com>
parent 06b92174
#!/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@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))
;[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/
;
......@@ -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)
/*!
[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@gmail.com
// Jia Haipeng, jiahaipeng95@gmail.com
// Xu Pang, pangxu010@163.com
// Wenju He, wenju@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];
}
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment