On Tue, Jan 21, 2014 at 11:17:43AM -0800, Dylan Baker wrote: > On Monday, January 20, 2014 09:16:54 AM Tom Stellard wrote: > > On Fri, Jan 17, 2014 at 03:04:46PM -0800, Dylan Baker wrote: > > > I think you meant to ping the piglit list not mesa dev :) > > > > > > In general this looks a lot better, I've got a few little things for you > > > > > > On Friday, January 17, 2014 02:30:13 PM Tom Stellard wrote: > > > > 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 > > > > > > > > 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, [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]; > > > > -} > > > > 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 <[email protected]> > > > > +# > > > > + > > > > +# 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 <[email protected]> > > > > +# > > > > + > > > > +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 > > > > > > the imports on the last line are wrong, they should be all caps, and > > > OPENCV_TEST_OCL_BIN does not match what you have in ocl_config.py.in. > > > There > > > are a few other cases of this > > > > > > Also, you can remove the following imports you're not using: os, sys, > > > string, TestProfile, testBinDir, Exectest > > > > > > > + > > > > +opencv_test_ocl = path.join(opencv_test_ocl_bin, 'opencv_test_ocl') > > > > > > here's one of those cases > > > > > > > + > > > > +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)) > > > > > > here's another of those cases > > > > > > > + ExecTest.__init__(self, options) > > > > + > > > > + > > > > +def add_opencv_tests(profile, individual = False): > > > > > > > + if opencv_test_ocl_bin == '': > > > and here > > > > > > > + 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 > > > > > > you only need one level of indent here beneath profile > > > > > > profile.tests['{}/{}'.format(full_test_name,test_name)] = \ > > > > > > OpenCVTest('{}{}'.format(group_name > > > > > > > > ,test_name)) > > > > > > The only other thing is that this file should probably be in framework, > > > not > > > tests. the only python that should be in tests are profiles, this is more > > > test classes. > > > > I moved opencv.py to framework, but now I can't import the ocl_config > > module. I would like to fix this by implementing your initial suggestion of > > adding a piglit config file and using python's config parser. Was your > > idea to add a piglit.config.in file in the main directory and then a > > piglitconfig class in framework for parsing and storing the values? > > > > -Tom > > It seems really strange to me that you can't import ocl_config. Can you push > your current version to a branch or send me patches? >
Here is the current version of the patch that I'm having trouble with. -Tom > for the config file I'm not sure that we really need to have cmake configure > it for us using the .in method. I'm not against it, just not sure that it's > the necessary. > > I had started working on a config parser for piglit, but never finished. I > rebased it and pushed it to github so you could see what I was thinking: > https://github.com/crymson/piglit/compare/config_parser > > > > > with those changes you have my r-b: > > > Reviewed-by: Dylan Baker <[email protected]>
>From 0126595fea071a7e32c4744269a8cee817f5f0c1 Mon Sep 17 00:00:00 2001 From: Tom Stellard <[email protected]> Date: Wed, 15 Jan 2014 11:28:44 -0800 Subject: [PATCH] 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. 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 --- CMakeLists.txt | 9 ++- framework/opencv.py | 70 +++++++++++++++++++ tests/all_cl.py | 4 ++ tests/cl/program/execute/opencv-merge-hist.cl | 98 --------------------------- tests/ocl_config.py.in | 31 +++++++++ 5 files changed, 113 insertions(+), 99 deletions(-) create mode 100644 framework/opencv.py delete mode 100644 tests/cl/program/execute/opencv-merge-hist.cl create mode 100644 tests/ocl_config.py.in diff --git a/CMakeLists.txt b/CMakeLists.txt index 47c8b10..48a761a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -107,7 +107,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") @@ -386,6 +389,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/framework/opencv.py b/framework/opencv.py new file mode 100644 index 0000000..cbdd7fd --- /dev/null +++ b/framework/opencv.py @@ -0,0 +1,70 @@ +#!/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 tests.ocl_config import OPENCV_TEST_OCL_BINDIR, OPENCV_TEST_OCL_WORKDIR + +opencv_test_ocl = path.join(OPENCV_TEST_OCL_BINDIR, '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_BINDIR == '': + 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('{}*'.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)) 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]; -} 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 <[email protected]> +# + +# 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}' -- 1.8.1.4
_______________________________________________ Piglit mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/piglit
