From: Alexander Alekhin Date: Fri, 6 Dec 2013 14:41:23 +0000 (+0400) Subject: core/ocl: update dynamic runtime X-Git-Tag: accepted/tizen/6.0/unified/20201030.111113~3635^2~4 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=5ecf7597428bbb8a7dfad98c309ef3f92d7d210f;p=platform%2Fupstream%2Fopencv.git core/ocl: update dynamic runtime --- diff --git a/modules/core/CMakeLists.txt b/modules/core/CMakeLists.txt index f0603ac..2e57d2e 100644 --- a/modules/core/CMakeLists.txt +++ b/modules/core/CMakeLists.txt @@ -1,5 +1,5 @@ set(the_description "The Core Functionality") -ocv_add_module(core PRIVATE_REQUIRED ${ZLIB_LIBRARIES} OPTIONAL opencv_cudev) +ocv_add_module(core PRIVATE_REQUIRED ${ZLIB_LIBRARIES} "${OPENCL_LIBRARIES}" OPTIONAL opencv_cudev) ocv_module_include_directories(${ZLIB_INCLUDE_DIRS}) if(HAVE_WINRT) diff --git a/modules/core/include/opencv2/core/opencl/runtime/opencl_clamdblas.hpp b/modules/core/include/opencv2/core/opencl/runtime/opencl_clamdblas.hpp new file mode 100644 index 0000000..b19f911 --- /dev/null +++ b/modules/core/include/opencv2/core/opencl/runtime/opencl_clamdblas.hpp @@ -0,0 +1,59 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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 materials 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 OpenCV Foundation 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. +// +//M*/ + +#ifndef __OPENCV_CORE_OCL_RUNTIME_CLAMDBLAS_HPP__ +#define __OPENCV_CORE_OCL_RUNTIME_CLAMDBLAS_HPP__ + +#ifdef HAVE_CLAMDBLAS + +#ifndef CL_RUNTIME_EXPORT +#if (defined(BUILD_SHARED_LIBS) || defined(OPENCV_CORE_SHARED)) && (defined WIN32 || defined _WIN32 || defined WINCE) +#define CL_RUNTIME_EXPORT __declspec(dllimport) +#else +#define CL_RUNTIME_EXPORT +#endif +#endif + +#include "autogenerated/opencl_clamdblas.hpp" + +#endif // HAVE_CLAMDBLAS + +#endif // __OPENCV_CORE_OCL_RUNTIME_CLAMDBLAS_HPP__ diff --git a/modules/core/include/opencv2/core/opencl/runtime/opencl_clamdfft.hpp b/modules/core/include/opencv2/core/opencl/runtime/opencl_clamdfft.hpp new file mode 100644 index 0000000..0af9307 --- /dev/null +++ b/modules/core/include/opencv2/core/opencl/runtime/opencl_clamdfft.hpp @@ -0,0 +1,59 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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 materials 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 OpenCV Foundation 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. +// +//M*/ + +#ifndef __OPENCV_CORE_OCL_RUNTIME_CLAMDFFT_HPP__ +#define __OPENCV_CORE_OCL_RUNTIME_CLAMDFFT_HPP__ + +#ifdef HAVE_CLAMDFFT + +#ifndef CL_RUNTIME_EXPORT +#if (defined(BUILD_SHARED_LIBS) || defined(OPENCV_CORE_SHARED)) && (defined WIN32 || defined _WIN32 || defined WINCE) +#define CL_RUNTIME_EXPORT __declspec(dllimport) +#else +#define CL_RUNTIME_EXPORT +#endif +#endif + +#include "autogenerated/opencl_clamdfft.hpp" + +#endif // HAVE_CLAMDFFT + +#endif // __OPENCV_CORE_OCL_RUNTIME_CLAMDFFT_HPP__ diff --git a/modules/core/include/opencv2/core/opencl/runtime/opencl_core.hpp b/modules/core/include/opencv2/core/opencl/runtime/opencl_core.hpp new file mode 100644 index 0000000..b19563c --- /dev/null +++ b/modules/core/include/opencv2/core/opencl/runtime/opencl_core.hpp @@ -0,0 +1,83 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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 materials 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 OpenCV Foundation 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. +// +//M*/ + +#ifndef __OPENCV_CORE_OCL_RUNTIME_OPENCL_CORE_HPP__ +#define __OPENCV_CORE_OCL_RUNTIME_OPENCL_CORE_HPP__ + +#ifdef HAVE_OPENCL + +#if defined(HAVE_OPENCL_STATIC) + +#if defined __APPLE__ +#include +#else +#include +#endif + +#else // HAVE_OPENCL_STATIC + +#ifndef CL_RUNTIME_EXPORT +#if (defined(BUILD_SHARED_LIBS) || defined(OPENCV_CORE_SHARED)) && (defined WIN32 || defined _WIN32 || defined WINCE) +#define CL_RUNTIME_EXPORT __declspec(dllimport) +#else +#define CL_RUNTIME_EXPORT +#endif +#endif + +#include "autogenerated/opencl_core.hpp" + +#endif // HAVE_OPENCL_STATIC + +#ifndef CL_DEVICE_DOUBLE_FP_CONFIG +#define CL_DEVICE_DOUBLE_FP_CONFIG 0x1032 +#endif + +#ifndef CL_DEVICE_HALF_FP_CONFIG +#define CL_DEVICE_HALF_FP_CONFIG 0x1033 +#endif + +#ifndef CL_VERSION_1_2 +#define CV_REQUIRE_OPENCL_1_2_ERROR CV_ErrorNoReturn(cv::Error::OpenCLApiCallError, "OpenCV compiled without OpenCL v1.2 support, so we can't use functionality from OpenCL v1.2") +#endif + +#endif // HAVE_OPENCL + +#endif // __OPENCV_CORE_OCL_RUNTIME_OPENCL_CORE_HPP__ diff --git a/modules/core/include/opencv2/core/opencl/runtime/opencl_core_wrappers.hpp b/modules/core/include/opencv2/core/opencl/runtime/opencl_core_wrappers.hpp new file mode 100644 index 0000000..912429f --- /dev/null +++ b/modules/core/include/opencv2/core/opencl/runtime/opencl_core_wrappers.hpp @@ -0,0 +1,47 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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 materials 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 OpenCV Foundation 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. +// +//M*/ + +#ifndef __OPENCV_CORE_OCL_RUNTIME_OPENCL_WRAPPERS_HPP__ +#define __OPENCV_CORE_OCL_RUNTIME_OPENCL_WRAPPERS_HPP__ + +#include "autogenerated/opencl_core_wrappers.hpp" + +#endif // __OPENCV_CORE_OCL_RUNTIME_OPENCL_WRAPPERS_HPP__ diff --git a/modules/core/src/opencl/runtime/generator/common.py b/modules/core/src/opencl/runtime/generator/common.py index 19b2145..ed0face 100644 --- a/modules/core/src/opencl/runtime/generator/common.py +++ b/modules/core/src/opencl/runtime/generator/common.py @@ -65,6 +65,7 @@ def getParameters(i, tokens): return params def postProcessParameters(fns): + fns.sort(key=lambda x: x['name']) for fn in fns: fn['params_full'] = list(fn['params']) for i in range(len(fn['params'])): @@ -79,6 +80,35 @@ def postProcessParameters(fns): del parts[-1] fn['params'][i] = ' '.join(parts) +def readFunctionFilter(fns, fileName): + try: + f = open(fileName, "r") + except: + print "ERROR: Can't open filter file: %s" % fileName + return 0 + + count = 0 + while f: + line = f.readline() + if not line: + break + assert isinstance(line, str) + if line.startswith('#') or line.startswith('//'): + continue + line = line.replace('\n', '') + if len(line) == 0: + continue + found = False + for fn in fns: + if fn['name'] == line: + found = True + fn['enabled'] = True + if not found: + sys.exit("FATAL ERROR: Unknown function: %s" % line) + count = count + 1 + f.close() + return count + # # Generator helpers # @@ -92,7 +122,7 @@ def outputToString(f): assert res is None sys.stdout = old_stdout result = str_stdout.getvalue() - result = re.sub(r'([^\n ]) [ ]+', r'\1 ', result) # don't remove spaces at start of line + result = re.sub(r'([^\n /]) [ ]+', r'\1 ', result) # don't remove spaces at start of line result = re.sub(r' ,', ',', result) result = re.sub(r' \*', '*', result) result = re.sub(r'\( ', '(', result) @@ -101,29 +131,48 @@ def outputToString(f): return wrapped @outputToString -def generateEnums(fns, prefix='OPENCL_FN'): +def generateFilterNames(fns): + for fn in fns: + print '%s%s' % ('' if fn.has_key('enabled') else '//', fn['name']) + print '#total %d' % len(fns) + +@outputToString +def generateStructDefinitions(fns, lprefix='opencl_fn', enumprefix='OPENCL_FN'): print '// generated by %s' % os.path.basename(sys.argv[0]) - print 'enum %s_ID {' % prefix first = True for fn in fns: - print ' %s_%s%s,' % (prefix, fn['name'], ' = 0' if first else '') + commentStr = '' if fn.has_key('enabled') else '//' + print commentStr + ('%s%s (%s *%s)(%s) =\n%s %s%d<%s_%s, %s%s>::switch_fn;' % \ + ((' '.join(fn['modifiers'] + ' ') if len(fn['modifiers']) > 0 else ''), + ' '.join(fn['ret']), ' '.join(fn['calling']), fn['name'], ', '.join(fn['params']), \ + commentStr, lprefix, len(fn['params']), enumprefix, fn['name'], ' '.join(fn['ret']), ('' if len(fn['params']) == 0 else ', ' + ', '.join(fn['params'])))) + print commentStr + ('static const struct DynamicFnEntry %s_definition = { "%s", (void**)&%s};' % (fn['name'], fn['name'], fn['name'])) + print first = False - print '};' @outputToString -def generateNames(fns, prefix='opencl_fn'): +def generateListOfDefinitions(fns, name='opencl_fn_list'): print '// generated by %s' % os.path.basename(sys.argv[0]) - print 'const char* %s_names[] = {' % prefix + print 'static const struct DynamicFnEntry* %s[] = {' % (name) for fn in fns: - print ' "%s",' % (fn['name']) + commentStr = '' if fn.has_key('enabled') else '//' + if fn.has_key('enabled'): + print ' &%s_definition,' % (fn['name']) + else: + print ' NULL/*&%s_definition*/,' % (fn['name']) + first = False + print ' ADDITIONAL_FN_DEFINITIONS // macro for custom functions' print '};' @outputToString -def generatePtrs(fns, prefix='opencl_fn'): +def generateEnums(fns, prefix='OPENCL_FN'): print '// generated by %s' % os.path.basename(sys.argv[0]) - print 'void* %s_ptrs[] = {' % prefix + print 'enum %s_ID {' % prefix + i = -1 for fn in fns: - print ' &%s,' % (fn['name']) + i = i + 1 + commentStr = '' if fn.has_key('enabled') else '//' + print commentStr + (' %s_%s = %d,' % (prefix, fn['name'], i)) print '};' @outputToString @@ -137,28 +186,16 @@ def generateRemapDynamic(fns): print '// generated by %s' % os.path.basename(sys.argv[0]) for fn in fns: print '#undef %s' % (fn['name']) - print '#define %s %s_pfn' % (fn['name'], fn['name']) - -@outputToString -def generateParamsCfg(fns): - for fn in fns: - print '%s %d' % (fn['name'], len(fn['params'])) + commentStr = '' if fn.has_key('enabled') else '//' + print commentStr + ('#define %s %s_pfn' % (fn['name'], fn['name'])) @outputToString def generateFnDeclaration(fns): print '// generated by %s' % os.path.basename(sys.argv[0]) for fn in fns: - print 'extern CL_RUNTIME_EXPORT %s %s (%s *%s)(%s);' % (' '.join(fn['modifiers']), ' '.join(fn['ret']), ' '.join(fn['calling']), - fn['name'], ', '.join(fn['params'] if not fn.has_key('params_full') else fn['params_full'])) - -@outputToString -def generateFnDefinition(fns, lprefix='opencl_fn', uprefix='OPENCL_FN'): - print '// generated by %s' % os.path.basename(sys.argv[0]) - for fn in fns: - print '%s%s (%s *%s)(%s) = %s%d<%s_%s, %s%s>::switch_fn;' % \ - ((' '.join(fn['modifiers'] + ' ') if len(fn['modifiers']) > 0 else ''), - ' '.join(fn['ret']), ' '.join(fn['calling']), fn['name'], ', '.join(fn['params']), \ - lprefix, len(fn['params']), uprefix, fn['name'], ' '.join(fn['ret']), ('' if len(fn['params']) == 0 else ', ' + ', '.join(fn['params']))) + commentStr = '' if fn.has_key('enabled') else '//' + print commentStr + ('extern CL_RUNTIME_EXPORT %s %s (%s *%s)(%s);' % (' '.join(fn['modifiers']), ' '.join(fn['ret']), ' '.join(fn['calling']), + fn['name'], ', '.join(fn['params'] if not fn.has_key('params_full') else fn['params_full']))) @outputToString def generateTemplates(sz, lprefix, switch_name, calling_convention=''): @@ -186,8 +223,9 @@ def generateTemplates(sz, lprefix, switch_name, calling_convention=''): def generateInlineWrappers(fns): print '// generated by %s' % os.path.basename(sys.argv[0]) for fn in fns: + commentStr = '' if fn.has_key('enabled') else '//' print '#undef %s' % (fn['name']) - print '#define %s %s_fn' % (fn['name'], fn['name']) + print commentStr + ('#define %s %s_fn' % (fn['name'], fn['name'])) params = [] call_params = [] for i in range(0, len(fn['params'])): @@ -200,11 +238,11 @@ def generateInlineWrappers(fns): call_params.append('p%d' % (i)) if len(fn['ret']) == 1 and fn['ret'][0] == 'void': - print 'inline void %s(%s) { %s_pfn(%s); }' \ - % (fn['name'], ', '.join(params), fn['name'], ', '.join(call_params)) + print commentStr + ('inline void %s(%s) { %s_pfn(%s); }' \ + % (fn['name'], ', '.join(params), fn['name'], ', '.join(call_params))) else: - print 'inline %s %s(%s) { return %s_pfn(%s); }' \ - % (' '.join(fn['ret']), fn['name'], ', '.join(params), fn['name'], ', '.join(call_params)) + print commentStr + ('inline %s %s(%s) { return %s_pfn(%s); }' \ + % (' '.join(fn['ret']), fn['name'], ', '.join(params), fn['name'], ', '.join(call_params))) def ProcessTemplate(inputFile, ctx, noteLine='//\n// AUTOGENERATED, DO NOT EDIT\n//'): f = open(inputFile, "r") diff --git a/modules/core/src/opencl/runtime/generator/filter/opencl_clamdblas_functions.list b/modules/core/src/opencl/runtime/generator/filter/opencl_clamdblas_functions.list new file mode 100644 index 0000000..2b1f076 --- /dev/null +++ b/modules/core/src/opencl/runtime/generator/filter/opencl_clamdblas_functions.list @@ -0,0 +1,176 @@ +//clAmdBlasAddScratchImage +//clAmdBlasCaxpy +//clAmdBlasCcopy +//clAmdBlasCdotc +//clAmdBlasCdotu +//clAmdBlasCgbmv +//clAmdBlasCgemm +clAmdBlasCgemmEx +//clAmdBlasCgemv +//clAmdBlasCgemvEx +//clAmdBlasCgerc +//clAmdBlasCgeru +//clAmdBlasChbmv +//clAmdBlasChemm +//clAmdBlasChemv +//clAmdBlasCher +//clAmdBlasCher2 +//clAmdBlasCher2k +//clAmdBlasCherk +//clAmdBlasChpmv +//clAmdBlasChpr +//clAmdBlasChpr2 +//clAmdBlasCrotg +//clAmdBlasCscal +//clAmdBlasCsrot +//clAmdBlasCsscal +//clAmdBlasCswap +//clAmdBlasCsymm +//clAmdBlasCsyr2k +//clAmdBlasCsyr2kEx +//clAmdBlasCsyrk +//clAmdBlasCsyrkEx +//clAmdBlasCtbmv +//clAmdBlasCtbsv +//clAmdBlasCtpmv +//clAmdBlasCtpsv +//clAmdBlasCtrmm +//clAmdBlasCtrmmEx +//clAmdBlasCtrmv +//clAmdBlasCtrsm +//clAmdBlasCtrsmEx +//clAmdBlasCtrsv +//clAmdBlasDasum +//clAmdBlasDaxpy +//clAmdBlasDcopy +//clAmdBlasDdot +//clAmdBlasDgbmv +//clAmdBlasDgemm +clAmdBlasDgemmEx +//clAmdBlasDgemv +//clAmdBlasDgemvEx +//clAmdBlasDger +//clAmdBlasDnrm2 +//clAmdBlasDrot +//clAmdBlasDrotg +//clAmdBlasDrotm +//clAmdBlasDrotmg +//clAmdBlasDsbmv +//clAmdBlasDscal +//clAmdBlasDspmv +//clAmdBlasDspr +//clAmdBlasDspr2 +//clAmdBlasDswap +//clAmdBlasDsymm +//clAmdBlasDsymv +//clAmdBlasDsymvEx +//clAmdBlasDsyr +//clAmdBlasDsyr2 +//clAmdBlasDsyr2k +//clAmdBlasDsyr2kEx +//clAmdBlasDsyrk +//clAmdBlasDsyrkEx +//clAmdBlasDtbmv +//clAmdBlasDtbsv +//clAmdBlasDtpmv +//clAmdBlasDtpsv +//clAmdBlasDtrmm +//clAmdBlasDtrmmEx +//clAmdBlasDtrmv +//clAmdBlasDtrsm +//clAmdBlasDtrsmEx +//clAmdBlasDtrsv +//clAmdBlasDzasum +//clAmdBlasDznrm2 +//clAmdBlasGetVersion +//clAmdBlasRemoveScratchImage +//clAmdBlasSasum +//clAmdBlasSaxpy +//clAmdBlasScasum +//clAmdBlasScnrm2 +//clAmdBlasScopy +//clAmdBlasSdot +clAmdBlasSetup +//clAmdBlasSgbmv +//clAmdBlasSgemm +clAmdBlasSgemmEx +//clAmdBlasSgemv +//clAmdBlasSgemvEx +//clAmdBlasSger +//clAmdBlasSnrm2 +//clAmdBlasSrot +//clAmdBlasSrotg +//clAmdBlasSrotm +//clAmdBlasSrotmg +//clAmdBlasSsbmv +//clAmdBlasSscal +//clAmdBlasSspmv +//clAmdBlasSspr +//clAmdBlasSspr2 +//clAmdBlasSswap +//clAmdBlasSsymm +//clAmdBlasSsymv +//clAmdBlasSsymvEx +//clAmdBlasSsyr +//clAmdBlasSsyr2 +//clAmdBlasSsyr2k +//clAmdBlasSsyr2kEx +//clAmdBlasSsyrk +//clAmdBlasSsyrkEx +//clAmdBlasStbmv +//clAmdBlasStbsv +//clAmdBlasStpmv +//clAmdBlasStpsv +//clAmdBlasStrmm +//clAmdBlasStrmmEx +//clAmdBlasStrmv +//clAmdBlasStrsm +//clAmdBlasStrsmEx +//clAmdBlasStrsv +clAmdBlasTeardown +//clAmdBlasZaxpy +//clAmdBlasZcopy +//clAmdBlasZdotc +//clAmdBlasZdotu +//clAmdBlasZdrot +//clAmdBlasZdscal +//clAmdBlasZgbmv +//clAmdBlasZgemm +clAmdBlasZgemmEx +//clAmdBlasZgemv +//clAmdBlasZgemvEx +//clAmdBlasZgerc +//clAmdBlasZgeru +//clAmdBlasZhbmv +//clAmdBlasZhemm +//clAmdBlasZhemv +//clAmdBlasZher +//clAmdBlasZher2 +//clAmdBlasZher2k +//clAmdBlasZherk +//clAmdBlasZhpmv +//clAmdBlasZhpr +//clAmdBlasZhpr2 +//clAmdBlasZrotg +//clAmdBlasZscal +//clAmdBlasZswap +//clAmdBlasZsymm +//clAmdBlasZsyr2k +//clAmdBlasZsyr2kEx +//clAmdBlasZsyrk +//clAmdBlasZsyrkEx +//clAmdBlasZtbmv +//clAmdBlasZtbsv +//clAmdBlasZtpmv +//clAmdBlasZtpsv +//clAmdBlasZtrmm +//clAmdBlasZtrmmEx +//clAmdBlasZtrmv +//clAmdBlasZtrsm +//clAmdBlasZtrsmEx +//clAmdBlasZtrsv +//clAmdBlasiCamax +//clAmdBlasiDamax +//clAmdBlasiSamax +//clAmdBlasiZamax +#total 175 diff --git a/modules/core/src/opencl/runtime/generator/filter/opencl_clamdfft_functions.list b/modules/core/src/opencl/runtime/generator/filter/opencl_clamdfft_functions.list new file mode 100644 index 0000000..7f407ff --- /dev/null +++ b/modules/core/src/opencl/runtime/generator/filter/opencl_clamdfft_functions.list @@ -0,0 +1,33 @@ +clAmdFftBakePlan +//clAmdFftCopyPlan +clAmdFftCreateDefaultPlan +clAmdFftDestroyPlan +clAmdFftEnqueueTransform +//clAmdFftGetLayout +//clAmdFftGetPlanBatchSize +//clAmdFftGetPlanContext +//clAmdFftGetPlanDim +//clAmdFftGetPlanDistance +//clAmdFftGetPlanInStride +//clAmdFftGetPlanLength +//clAmdFftGetPlanOutStride +//clAmdFftGetPlanPrecision +//clAmdFftGetPlanScale +//clAmdFftGetPlanTransposeResult +//clAmdFftGetResultLocation +clAmdFftGetTmpBufSize +//clAmdFftGetVersion +clAmdFftSetLayout +clAmdFftSetPlanBatchSize +//clAmdFftSetPlanDim +clAmdFftSetPlanDistance +clAmdFftSetPlanInStride +//clAmdFftSetPlanLength +clAmdFftSetPlanOutStride +//clAmdFftSetPlanPrecision +clAmdFftSetPlanScale +//clAmdFftSetPlanTransposeResult +clAmdFftSetResultLocation +clAmdFftSetup +clAmdFftTeardown +#total 32 diff --git a/modules/core/src/opencl/runtime/generator/filter/opencl_core_functions.list b/modules/core/src/opencl/runtime/generator/filter/opencl_core_functions.list new file mode 100644 index 0000000..2dc5185 --- /dev/null +++ b/modules/core/src/opencl/runtime/generator/filter/opencl_core_functions.list @@ -0,0 +1,89 @@ +clBuildProgram +clCompileProgram +clCreateBuffer +clCreateCommandQueue +clCreateContext +clCreateContextFromType +clCreateImage +clCreateImage2D +clCreateImage3D +clCreateKernel +clCreateKernelsInProgram +clCreateProgramWithBinary +clCreateProgramWithBuiltInKernels +clCreateProgramWithSource +clCreateSampler +clCreateSubBuffer +clCreateSubDevices +clCreateUserEvent +clEnqueueBarrier +clEnqueueBarrierWithWaitList +clEnqueueCopyBuffer +clEnqueueCopyBufferRect +clEnqueueCopyBufferToImage +clEnqueueCopyImage +clEnqueueCopyImageToBuffer +clEnqueueFillBuffer +clEnqueueFillImage +clEnqueueMapBuffer +clEnqueueMapImage +clEnqueueMarker +clEnqueueMarkerWithWaitList +clEnqueueMigrateMemObjects +clEnqueueNDRangeKernel +clEnqueueNativeKernel +clEnqueueReadBuffer +clEnqueueReadBufferRect +clEnqueueReadImage +clEnqueueTask +clEnqueueUnmapMemObject +clEnqueueWaitForEvents +clEnqueueWriteBuffer +clEnqueueWriteBufferRect +clEnqueueWriteImage +clFinish +clFlush +clGetCommandQueueInfo +clGetContextInfo +clGetDeviceIDs +clGetDeviceInfo +clGetEventInfo +clGetEventProfilingInfo +clGetExtensionFunctionAddress +clGetExtensionFunctionAddressForPlatform +clGetImageInfo +clGetKernelArgInfo +clGetKernelInfo +clGetKernelWorkGroupInfo +clGetMemObjectInfo +clGetPlatformIDs +clGetPlatformInfo +clGetProgramBuildInfo +clGetProgramInfo +clGetSamplerInfo +clGetSupportedImageFormats +clLinkProgram +clReleaseCommandQueue +clReleaseContext +clReleaseDevice +clReleaseEvent +clReleaseKernel +clReleaseMemObject +clReleaseProgram +clReleaseSampler +clRetainCommandQueue +clRetainContext +clRetainDevice +clRetainEvent +clRetainKernel +clRetainMemObject +clRetainProgram +clRetainSampler +clSetEventCallback +clSetKernelArg +clSetMemObjectDestructorCallback +clSetUserEventStatus +clUnloadCompiler +clUnloadPlatformCompiler +clWaitForEvents +#total 88 diff --git a/modules/core/src/opencl/runtime/generator/generate.sh b/modules/core/src/opencl/runtime/generator/generate.sh index bc2d8d0..8649e99 100644 --- a/modules/core/src/opencl/runtime/generator/generate.sh +++ b/modules/core/src/opencl/runtime/generator/generate.sh @@ -1,6 +1,6 @@ #!/bin/bash -e echo "Generate files for CL runtime..." -cat sources/cl.h | python parser_cl.py cl_runtime_opencl -cat sources/clAmdBlas.h | python parser_clamdblas.py -cat sources/clAmdFft.h | python parser_clamdfft.py +python parser_cl.py opencl_core < sources/cl.h +python parser_clamdblas.py < sources/clAmdBlas.h +python parser_clamdfft.py < sources/clAmdFft.h echo "Generate files for CL runtime... Done" diff --git a/modules/core/src/opencl/runtime/generator/parser_cl.py b/modules/core/src/opencl/runtime/generator/parser_cl.py index 0ffbe75..87eeb27 100644 --- a/modules/core/src/opencl/runtime/generator/parser_cl.py +++ b/modules/core/src/opencl/runtime/generator/parser_cl.py @@ -8,9 +8,9 @@ from common import remove_comments, getTokens, getParameters, postProcessParamet try: if len(sys.argv) > 1: - outfile = open('../../../include/opencv2/ocl/cl_runtime/' + sys.argv[1] + '.hpp', "w") - outfile_impl = open('../' + sys.argv[1] + '_impl.hpp', "w") - outfile_wrappers = open('../../../include/opencv2/ocl/cl_runtime/' + sys.argv[1] + '_wrappers.hpp', "w") + outfile = open('../../../../include/opencv2/core/opencl/runtime/autogenerated/' + sys.argv[1] + '.hpp', 'wb') + outfile_impl = open('../autogenerated/' + sys.argv[1] + '_impl.hpp', 'wb') + outfile_wrappers = open('../../../../include/opencv2/core/opencl/runtime/autogenerated/' + sys.argv[1] + '_wrappers.hpp', 'wb') if len(sys.argv) > 2: f = open(sys.argv[2], "r") else: @@ -95,24 +95,31 @@ pprint(fns) from common import * +filterFileName='./filter/opencl_core_functions.list' +numEnabled = readFunctionFilter(fns, filterFileName) + +functionsFilter = generateFilterNames(fns) +filter_file = open(filterFileName, 'wb') +filter_file.write(functionsFilter) + ctx = {} ctx['CL_REMAP_ORIGIN'] = generateRemapOrigin(fns) ctx['CL_REMAP_DYNAMIC'] = generateRemapDynamic(fns) ctx['CL_FN_DECLARATIONS'] = generateFnDeclaration(fns) sys.stdout = outfile -ProcessTemplate('template/cl_runtime_opencl.hpp.in', ctx) +ProcessTemplate('template/opencl_core.hpp.in', ctx) ctx['CL_FN_INLINE_WRAPPERS'] = generateInlineWrappers(fns) sys.stdout = outfile_wrappers -ProcessTemplate('template/cl_runtime_opencl_wrappers.hpp.in', ctx) +ProcessTemplate('template/opencl_core_wrappers.hpp.in', ctx) +ctx['CL_FN_ENTRY_DEFINITIONS'] = generateStructDefinitions(fns) +ctx['CL_FN_ENTRY_LIST'] = generateListOfDefinitions(fns) ctx['CL_FN_ENUMS'] = generateEnums(fns) -ctx['CL_FN_NAMES'] = generateNames(fns) -ctx['CL_FN_DEFINITIONS'] = generateFnDefinition(fns) -ctx['CL_FN_PTRS'] = generatePtrs(fns) ctx['CL_FN_SWITCH'] = generateTemplates(15, 'opencl_fn', 'opencl_check_fn', 'CL_API_CALL') +ctx['CL_NUMBER_OF_ENABLED_FUNCTIONS'] = '// number of enabled functions: %d' % (numEnabled) sys.stdout = outfile_impl -ProcessTemplate('template/cl_runtime_impl_opencl.hpp.in', ctx) +ProcessTemplate('template/opencl_core_impl.hpp.in', ctx) diff --git a/modules/core/src/opencl/runtime/generator/parser_clamdblas.py b/modules/core/src/opencl/runtime/generator/parser_clamdblas.py index 52e62c5..2b473a9 100644 --- a/modules/core/src/opencl/runtime/generator/parser_clamdblas.py +++ b/modules/core/src/opencl/runtime/generator/parser_clamdblas.py @@ -89,19 +89,26 @@ pprint(fns) from common import * +filterFileName='./filter/opencl_clamdblas_functions.list' +numEnabled = readFunctionFilter(fns, filterFileName) + +functionsFilter = generateFilterNames(fns) +filter_file = open(filterFileName, 'wb') +filter_file.write(functionsFilter) + ctx = {} ctx['CLAMDBLAS_REMAP_ORIGIN'] = generateRemapOrigin(fns) ctx['CLAMDBLAS_REMAP_DYNAMIC'] = generateRemapDynamic(fns) ctx['CLAMDBLAS_FN_DECLARATIONS'] = generateFnDeclaration(fns) -sys.stdout = open('../../../include/opencv2/ocl/cl_runtime/clamdblas_runtime.hpp', 'w') -ProcessTemplate('template/clamdblas_runtime.hpp.in', ctx) +sys.stdout = open('../../../../include/opencv2/core/opencl/runtime/autogenerated/opencl_clamdblas.hpp', 'wb') +ProcessTemplate('template/opencl_clamdblas.hpp.in', ctx) -ctx['CL_FN_ENUMS'] = generateEnums(fns, 'OPENCLAMDBLAS_FN') -ctx['CL_FN_NAMES'] = generateNames(fns, 'openclamdblas_fn') -ctx['CL_FN_DEFINITIONS'] = generateFnDefinition(fns, 'openclamdblas_fn', 'OPENCLAMDBLAS_FN') -ctx['CL_FN_PTRS'] = generatePtrs(fns, 'openclamdblas_fn') +ctx['CL_FN_ENUMS'] = generateEnums(fns, 'OPENCLAMDBLAS_FN', ) ctx['CL_FN_SWITCH'] = generateTemplates(23, 'openclamdblas_fn', 'openclamdblas_check_fn', '') +ctx['CL_FN_ENTRY_DEFINITIONS'] = generateStructDefinitions(fns, 'openclamdblas_fn', 'OPENCLAMDBLAS_FN') +ctx['CL_FN_ENTRY_LIST'] = generateListOfDefinitions(fns, 'openclamdblas_fn') +ctx['CL_NUMBER_OF_ENABLED_FUNCTIONS'] = '// number of enabled functions: %d' % (numEnabled) -sys.stdout = open('../clamdblas_runtime.cpp', 'w') -ProcessTemplate('template/clamdblas_runtime.cpp.in', ctx) +sys.stdout = open('../autogenerated/opencl_clamdblas_impl.hpp', 'wb') +ProcessTemplate('template/opencl_clamdblas_impl.hpp.in', ctx) diff --git a/modules/core/src/opencl/runtime/generator/parser_clamdfft.py b/modules/core/src/opencl/runtime/generator/parser_clamdfft.py index 35b78ca..accc20b 100644 --- a/modules/core/src/opencl/runtime/generator/parser_clamdfft.py +++ b/modules/core/src/opencl/runtime/generator/parser_clamdfft.py @@ -86,19 +86,26 @@ pprint(fns) from common import * +filterFileName='./filter/opencl_clamdfft_functions.list' +numEnabled = readFunctionFilter(fns, filterFileName) + +functionsFilter = generateFilterNames(fns) +filter_file = open(filterFileName, 'wb') +filter_file.write(functionsFilter) + ctx = {} ctx['CLAMDFFT_REMAP_ORIGIN'] = generateRemapOrigin(fns) ctx['CLAMDFFT_REMAP_DYNAMIC'] = generateRemapDynamic(fns) ctx['CLAMDFFT_FN_DECLARATIONS'] = generateFnDeclaration(fns) -sys.stdout = open('../../../include/opencv2/ocl/cl_runtime/clamdfft_runtime.hpp', 'w') -ProcessTemplate('template/clamdfft_runtime.hpp.in', ctx) +sys.stdout = open('../../../../include/opencv2/core/opencl/runtime/autogenerated/opencl_clamdfft.hpp', 'wb') +ProcessTemplate('template/opencl_clamdfft.hpp.in', ctx) ctx['CL_FN_ENUMS'] = generateEnums(fns, 'OPENCLAMDFFT_FN') -ctx['CL_FN_NAMES'] = generateNames(fns, 'openclamdfft_fn') -ctx['CL_FN_DEFINITIONS'] = generateFnDefinition(fns, 'openclamdfft_fn', 'OPENCLAMDFFT_FN') -ctx['CL_FN_PTRS'] = generatePtrs(fns, 'openclamdfft_fn') ctx['CL_FN_SWITCH'] = generateTemplates(23, 'openclamdfft_fn', 'openclamdfft_check_fn', '') +ctx['CL_FN_ENTRY_DEFINITIONS'] = generateStructDefinitions(fns, 'openclamdfft_fn', 'OPENCLAMDFFT_FN') +ctx['CL_FN_ENTRY_LIST'] = generateListOfDefinitions(fns, 'openclamdfft_fn') +ctx['CL_NUMBER_OF_ENABLED_FUNCTIONS'] = '// number of enabled functions: %d' % (numEnabled) -sys.stdout = open('../clamdfft_runtime.cpp', 'w') -ProcessTemplate('template/clamdfft_runtime.cpp.in', ctx) +sys.stdout = open('../autogenerated/opencl_clamdfft_impl.hpp', 'wb') +ProcessTemplate('template/opencl_clamdfft_impl.hpp.in', ctx) diff --git a/modules/core/src/opencl/runtime/generator/template/opencl_clamdblas.hpp.in b/modules/core/src/opencl/runtime/generator/template/opencl_clamdblas.hpp.in new file mode 100644 index 0000000..251c085 --- /dev/null +++ b/modules/core/src/opencl/runtime/generator/template/opencl_clamdblas.hpp.in @@ -0,0 +1,11 @@ +#ifndef __OPENCV_CORE_OCL_RUNTIME_CLAMDBLAS_HPP__ +#error "Invalid usage" +#endif + +@CLAMDBLAS_REMAP_ORIGIN@ + +#include + +@CLAMDBLAS_REMAP_DYNAMIC@ + +@CLAMDBLAS_FN_DECLARATIONS@ diff --git a/modules/core/src/opencl/runtime/generator/template/opencl_clamdblas_impl.hpp.in b/modules/core/src/opencl/runtime/generator/template/opencl_clamdblas_impl.hpp.in new file mode 100644 index 0000000..11c834f --- /dev/null +++ b/modules/core/src/opencl/runtime/generator/template/opencl_clamdblas_impl.hpp.in @@ -0,0 +1,15 @@ +#ifndef ADDITIONAL_FN_DEFINITIONS +#define ADDITIONAL_FN_DEFINITIONS +#endif + +@CL_FN_ENUMS@ + +namespace { +@CL_FN_SWITCH@ +} + +@CL_FN_ENTRY_DEFINITIONS@ + +@CL_FN_ENTRY_LIST@ + +@CL_NUMBER_OF_ENABLED_FUNCTIONS@ diff --git a/modules/core/src/opencl/runtime/generator/template/opencl_clamdfft.hpp.in b/modules/core/src/opencl/runtime/generator/template/opencl_clamdfft.hpp.in new file mode 100644 index 0000000..8633580 --- /dev/null +++ b/modules/core/src/opencl/runtime/generator/template/opencl_clamdfft.hpp.in @@ -0,0 +1,11 @@ +#ifndef __OPENCV_CORE_OCL_RUNTIME_CLAMDFFT_HPP__ +#error "Invalid usage" +#endif + +@CLAMDFFT_REMAP_ORIGIN@ + +#include + +@CLAMDFFT_REMAP_DYNAMIC@ + +@CLAMDFFT_FN_DECLARATIONS@ diff --git a/modules/core/src/opencl/runtime/generator/template/opencl_clamdfft_impl.hpp.in b/modules/core/src/opencl/runtime/generator/template/opencl_clamdfft_impl.hpp.in new file mode 100644 index 0000000..11c834f --- /dev/null +++ b/modules/core/src/opencl/runtime/generator/template/opencl_clamdfft_impl.hpp.in @@ -0,0 +1,15 @@ +#ifndef ADDITIONAL_FN_DEFINITIONS +#define ADDITIONAL_FN_DEFINITIONS +#endif + +@CL_FN_ENUMS@ + +namespace { +@CL_FN_SWITCH@ +} + +@CL_FN_ENTRY_DEFINITIONS@ + +@CL_FN_ENTRY_LIST@ + +@CL_NUMBER_OF_ENABLED_FUNCTIONS@ diff --git a/modules/core/src/opencl/runtime/generator/template/opencl_core.hpp.in b/modules/core/src/opencl/runtime/generator/template/opencl_core.hpp.in new file mode 100644 index 0000000..4196622 --- /dev/null +++ b/modules/core/src/opencl/runtime/generator/template/opencl_core.hpp.in @@ -0,0 +1,15 @@ +#ifndef __OPENCV_CORE_OCL_RUNTIME_OPENCL_CORE_HPP__ +#error "Invalid usage" +#endif + +@CL_REMAP_ORIGIN@ + +#if defined __APPLE__ +#include +#else +#include +#endif + +@CL_REMAP_DYNAMIC@ + +@CL_FN_DECLARATIONS@ diff --git a/modules/core/src/opencl/runtime/generator/template/opencl_core_impl.hpp.in b/modules/core/src/opencl/runtime/generator/template/opencl_core_impl.hpp.in new file mode 100644 index 0000000..f3adb66 --- /dev/null +++ b/modules/core/src/opencl/runtime/generator/template/opencl_core_impl.hpp.in @@ -0,0 +1,16 @@ + +#ifndef ADDITIONAL_FN_DEFINITIONS +#define ADDITIONAL_FN_DEFINITIONS +#endif + +@CL_FN_ENUMS@ + +namespace { +@CL_FN_SWITCH@ +} // anonymous namespace + +@CL_FN_ENTRY_DEFINITIONS@ + +@CL_FN_ENTRY_LIST@ + +@CL_NUMBER_OF_ENABLED_FUNCTIONS@ diff --git a/modules/core/src/opencl/runtime/generator/template/opencl_core_wrappers.hpp.in b/modules/core/src/opencl/runtime/generator/template/opencl_core_wrappers.hpp.in new file mode 100644 index 0000000..847e67f --- /dev/null +++ b/modules/core/src/opencl/runtime/generator/template/opencl_core_wrappers.hpp.in @@ -0,0 +1,5 @@ +#ifndef __OPENCV_CORE_OCL_RUNTIME_OPENCL_WRAPPERS_HPP__ +#error "Invalid usage" +#endif + +@CL_FN_INLINE_WRAPPERS@ diff --git a/modules/core/src/opencl/runtime/opencl_clamdblas.cpp b/modules/core/src/opencl/runtime/opencl_clamdblas.cpp new file mode 100644 index 0000000..6296ef6 --- /dev/null +++ b/modules/core/src/opencl/runtime/opencl_clamdblas.cpp @@ -0,0 +1,125 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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 materials 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 OpenCV Foundation 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. +// +//M*/ + +#include "precomp.hpp" + +#ifdef HAVE_CLAMDBLAS + +#include "opencv2/core/opencl/runtime/opencl_core.hpp" +#include "opencv2/core/opencl/runtime/opencl_clamdblas.hpp" + +#if defined(_WIN32) +#include + + static void* WinGetProcAddress(const char* name) + { + static HMODULE opencl_module = NULL; + if (!opencl_module) + { + opencl_module = GetModuleHandleA("clAmdBlas.dll"); + if (!opencl_module) + { + opencl_module = LoadLibraryA("clAmdBlas.dll"); + if (!opencl_module) + return NULL; + } + } + return (void*)GetProcAddress(opencl_module, name); + } + #define CV_CL_GET_PROC_ADDRESS(name) WinGetProcAddress(name) +#endif // _WIN32 + +#if defined(linux) + #include + #include + + static void* GetProcAddress (const char* name) + { + static void* h = NULL; + if (!h) + { + h = dlopen("libclAmdBlas.so", RTLD_LAZY | RTLD_GLOBAL); + if (!h) + return NULL; + } + + return dlsym(h, name); + } + #define CV_CL_GET_PROC_ADDRESS(name) GetProcAddress(name) +#endif + +#ifndef CV_CL_GET_PROC_ADDRESS +#define CV_CL_GET_PROC_ADDRESS(name) NULL +#endif + +static void* openclamdblas_check_fn(int ID); + +#include "runtime_common.hpp" + +// +// BEGIN OF CUSTOM FUNCTIONS +// + +#define CUSTOM_FUNCTION_ID 1000 + +#undef ADDITIONAL_FN_DEFINITIONS + +// +// END OF CUSTOM FUNCTIONS HERE +// + +#include "autogenerated/opencl_clamdblas_impl.hpp" + +static void* openclamdblas_check_fn(int ID) +{ + ID = (ID <= CUSTOM_FUNCTION_ID) ? ID : ID - CUSTOM_FUNCTION_ID; + assert(ID >= 0 && ID < (int)(sizeof(openclamdblas_fn)/sizeof(openclamdblas_fn[0]))); + const struct DynamicFnEntry* e = openclamdblas_fn[ID]; + void* func = CV_CL_GET_PROC_ADDRESS(e->fnName); + if (!func) + { + CV_Error(cv::Error::OpenCLApiCallError, cv::format("OpenCL AMD BLAS function is not available: [%s]", e->fnName)); + } + *(e->ppFn) = func; + return func; +} + +#endif diff --git a/modules/core/src/opencl/runtime/opencl_clamdfft.cpp b/modules/core/src/opencl/runtime/opencl_clamdfft.cpp new file mode 100644 index 0000000..2514b0a --- /dev/null +++ b/modules/core/src/opencl/runtime/opencl_clamdfft.cpp @@ -0,0 +1,125 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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 materials 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 OpenCV Foundation 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. +// +//M*/ + +#include "precomp.hpp" + +#ifdef HAVE_CLAMDFFT + +#include "opencv2/core/opencl/runtime/opencl_core.hpp" +#include "opencv2/core/opencl/runtime/opencl_clamdfft.hpp" + +#if defined(_WIN32) +#include + + static void* WinGetProcAddress(const char* name) + { + static HMODULE opencl_module = NULL; + if (!opencl_module) + { + opencl_module = GetModuleHandleA("clAmdFft.Runtime.dll"); + if (!opencl_module) + { + opencl_module = LoadLibraryA("clAmdFft.Runtime.dll"); + if (!opencl_module) + return NULL; + } + } + return (void*)GetProcAddress(opencl_module, name); + } + #define CV_CL_GET_PROC_ADDRESS(name) WinGetProcAddress(name) +#endif // _WIN32 + +#if defined(linux) + #include + #include + + static void* GetProcAddress (const char* name) + { + static void* h = NULL; + if (!h) + { + h = dlopen("libclAmdFft.Runtime.so", RTLD_LAZY | RTLD_GLOBAL); + if (!h) + return NULL; + } + + return dlsym(h, name); + } + #define CV_CL_GET_PROC_ADDRESS(name) GetProcAddress(name) +#endif + +#ifndef CV_CL_GET_PROC_ADDRESS +#define CV_CL_GET_PROC_ADDRESS(name) NULL +#endif + +static void* openclamdfft_check_fn(int ID); + +#include "runtime_common.hpp" + +// +// BEGIN OF CUSTOM FUNCTIONS +// + +#define CUSTOM_FUNCTION_ID 1000 + +#undef ADDITIONAL_FN_DEFINITIONS + +// +// END OF CUSTOM FUNCTIONS HERE +// + +#include "autogenerated/opencl_clamdfft_impl.hpp" + +static void* openclamdfft_check_fn(int ID) +{ + ID = (ID <= CUSTOM_FUNCTION_ID) ? ID : ID - CUSTOM_FUNCTION_ID; + assert(ID >= 0 && ID < (int)(sizeof(openclamdfft_fn)/sizeof(openclamdfft_fn[0]))); + const struct DynamicFnEntry* e = openclamdfft_fn[ID]; + void* func = CV_CL_GET_PROC_ADDRESS(e->fnName); + if (!func) + { + CV_Error(cv::Error::OpenCLApiCallError, cv::format("OpenCL AMD FFT function is not available: [%s]", e->fnName)); + } + *(e->ppFn) = func; + return func; +} + +#endif diff --git a/modules/core/src/opencl/runtime/opencl_core.cpp b/modules/core/src/opencl/runtime/opencl_core.cpp new file mode 100644 index 0000000..d8f231b --- /dev/null +++ b/modules/core/src/opencl/runtime/opencl_core.cpp @@ -0,0 +1,200 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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 materials 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 OpenCV Foundation 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. +// +//M*/ + +#include "precomp.hpp" + +#if defined(HAVE_OPENCL) && !defined(HAVE_OPENCL_STATIC) + +#include "opencv2/core.hpp" // CV_Error + +#include "opencv2/core/opencl/runtime/opencl_core.hpp" + +static const char* funcToCheckOpenCL1_1 = "clEnqueueReadBufferRect"; +#define ERROR_MSG_CANT_LOAD "Failed to load OpenCL runtime\n" +#define ERROR_MSG_INVALID_VERSION "Failed to load OpenCL runtime (expected version 1.1+)\n" + +#if defined(__APPLE__) +#include + +static void* AppleCLGetProcAddress(const char* name) +{ + static bool initialized = false; + static void* handle = NULL; + if (!handle) + { + if(!initialized) + { + initialized = true; + const char* path = "/System/Library/Frameworks/OpenCL.framework/Versions/Current/OpenCL"; + const char* envPath = getenv("OPENCV_OPENCL_RUNTIME"); + if (envPath) + path = envPath; + handle = dlopen(oclpath, RTLD_LAZY | RTLD_GLOBAL); + if (handle == NULL) + { + fprintf(stderr, ERROR_MSG_CANT_LOAD); + } + else if (dlsym(handle, funcToCheckOpenCL1_1) == NULL) + { + fprintf(stderr, ERROR_MSG_INVALID_VERSION); + handle = NULL; + } + } + if (!handle) + return NULL; + } + return dlsym(handle, name); +} +#define CV_CL_GET_PROC_ADDRESS(name) AppleCLGetProcAddress(name) +#endif // __APPLE__ + +#if defined(_WIN32) +#include + +static void* WinGetProcAddress(const char* name) +{ + static bool initialized = false; + static HMODULE handle = NULL; + if (!handle) + { + if(!initialized) + { + initialized = true; + handle = GetModuleHandleA("OpenCL.dll"); + if (!handle) + { + const char* path = "OpenCL.dll"; + const char* envPath = getenv("OPENCV_OPENCL_RUNTIME"); + if (envPath) + path = envPath; + handle = LoadLibraryA(path); + if (!handle) + { + fprintf(stderr, ERROR_MSG_CANT_LOAD); + } + else if (GetProcAddress(handle, funcToCheckOpenCL1_1) == NULL) + { + fprintf(stderr, ERROR_MSG_INVALID_VERSION); + handle = NULL; + } + } + } + if (!handle) + return NULL; + } + return (void*)GetProcAddress(handle, name); +} +#define CV_CL_GET_PROC_ADDRESS(name) WinGetProcAddress(name) +#endif // _WIN32 + +#if defined(linux) +#include +#include + +static void* GetProcAddress(const char* name) +{ + static bool initialized = false; + static void* handle = NULL; + if (!handle) + { + if(!initialized) + { + initialized = true; + const char* path = "libOpenCL.so"; + const char* envPath = getenv("OPENCV_OPENCL_RUNTIME"); + if (envPath) + path = envPath; + handle = dlopen(path, RTLD_LAZY | RTLD_GLOBAL); + if (handle == NULL) + { + fprintf(stderr, ERROR_MSG_CANT_LOAD); + } + else if (dlsym(handle, funcToCheckOpenCL1_1) == NULL) + { + fprintf(stderr, ERROR_MSG_INVALID_VERSION); + handle = NULL; + } + } + if (!handle) + return NULL; + } + return dlsym(handle, name); +} +#define CV_CL_GET_PROC_ADDRESS(name) GetProcAddress(name) +#endif + +#ifndef CV_CL_GET_PROC_ADDRESS +#define CV_CL_GET_PROC_ADDRESS(name) NULL +#endif + +static void* opencl_check_fn(int ID); + +#include "runtime_common.hpp" + +// +// BEGIN OF CUSTOM FUNCTIONS +// + +#define CUSTOM_FUNCTION_ID 1000 + +#undef ADDITIONAL_FN_DEFINITIONS + +// +// END OF CUSTOM FUNCTIONS HERE +// + +#include "autogenerated/opencl_core_impl.hpp" + +static void* opencl_check_fn(int ID) +{ + ID = (ID <= CUSTOM_FUNCTION_ID) ? ID : ID - CUSTOM_FUNCTION_ID; + assert(ID >= 0 && ID < (int)(sizeof(opencl_fn_list)/sizeof(opencl_fn_list[0]))); + const struct DynamicFnEntry* e = opencl_fn_list[ID]; + void* func = CV_CL_GET_PROC_ADDRESS(e->fnName); + if (!func) + { + CV_Error(cv::Error::OpenCLApiCallError, cv::format("OpenCL function is not available: [%s]", e->fnName)); + } + *(e->ppFn) = func; + return func; +} + +#endif diff --git a/modules/core/src/opencl/runtime/runtime_common.hpp b/modules/core/src/opencl/runtime/runtime_common.hpp new file mode 100644 index 0000000..542ef25 --- /dev/null +++ b/modules/core/src/opencl/runtime/runtime_common.hpp @@ -0,0 +1,57 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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 materials 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 OpenCV Foundation 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. +// +//M*/ + +#ifndef __OPENCV_CORE_OCL_RUNTIME_COMMON_HPP__ +#define __OPENCV_CORE_OCL_RUNTIME_COMMON_HPP__ + +namespace cv { namespace ocl { namespace runtime { + +struct DynamicFnEntry +{ + const char* fnName; // "clCreateContext" + void** ppFn; // &clCreateContext_pfn +}; + +} } } // namespace cv::ocl::runtime + +using namespace cv::ocl::runtime; + +#endif // __OPENCV_CORE_OCL_RUNTIME_COMMON_HPP__ diff --git a/modules/core/src/precomp.hpp b/modules/core/src/precomp.hpp index 87104af..47d6fdc 100644 --- a/modules/core/src/precomp.hpp +++ b/modules/core/src/precomp.hpp @@ -244,6 +244,18 @@ struct TLSData static TLSData* get(); }; +#if defined(BUILD_SHARED_LIBS) +#if defined WIN32 || defined _WIN32 || defined WINCE +#define CL_RUNTIME_EXPORT __declspec(dllexport) +#elif defined __GNUC__ && __GNUC__ >= 4 +#define CL_RUNTIME_EXPORT __attribute__ ((visibility ("default"))) +#else +#define CL_RUNTIME_EXPORT +#endif +#else +#define CL_RUNTIME_EXPORT +#endif + namespace ocl { MatAllocator* getOpenCLAllocator(); diff --git a/modules/ocl/include/opencv2/ocl/private/opencl_utils.hpp b/modules/ocl/include/opencv2/ocl/private/opencl_utils.hpp index 08f980f..dfc6583 100644 --- a/modules/ocl/include/opencv2/ocl/private/opencl_utils.hpp +++ b/modules/ocl/include/opencv2/ocl/private/opencl_utils.hpp @@ -42,7 +42,7 @@ #ifndef __OPENCV_OCL_PRIVATE_OPENCL_UTILS_HPP__ #define __OPENCV_OCL_PRIVATE_OPENCL_UTILS_HPP__ -#include "opencv2/ocl/cl_runtime/cl_runtime.hpp" +#include "opencv2/core/opencl/runtime/opencl_core.hpp" #include #include diff --git a/modules/ocl/include/opencv2/ocl/private/util.hpp b/modules/ocl/include/opencv2/ocl/private/util.hpp index 98b734a..b1ceacd 100644 --- a/modules/ocl/include/opencv2/ocl/private/util.hpp +++ b/modules/ocl/include/opencv2/ocl/private/util.hpp @@ -46,7 +46,7 @@ #ifndef __OPENCV_OCL_PRIVATE_UTIL__ #define __OPENCV_OCL_PRIVATE_UTIL__ -#include "opencv2/ocl/cl_runtime/cl_runtime.hpp" +#include "opencv2/core/opencl/runtime/opencl_core.hpp" #include "opencv2/core/ocl_genbase.hpp" #include "opencv2/ocl.hpp" diff --git a/modules/ocl/src/fft.cpp b/modules/ocl/src/fft.cpp index e10d671..70711bb 100644 --- a/modules/ocl/src/fft.cpp +++ b/modules/ocl/src/fft.cpp @@ -57,7 +57,7 @@ namespace cv { namespace ocl { }} void cv::ocl::fft_teardown(){} #else -#include "opencv2/ocl/cl_runtime/clamdfft_runtime.hpp" +#include "opencv2/core/opencl/runtime/opencl_clamdfft.hpp" namespace cv { namespace ocl diff --git a/modules/ocl/src/gemm.cpp b/modules/ocl/src/gemm.cpp index 92119df..50a2fdc 100644 --- a/modules/ocl/src/gemm.cpp +++ b/modules/ocl/src/gemm.cpp @@ -72,7 +72,7 @@ void cv::ocl::clBlasTeardown() } #else -#include "opencv2/ocl/cl_runtime/clamdblas_runtime.hpp" +#include "opencv2/core/opencl/runtime/opencl_clamdblas.hpp" using namespace cv; static bool clBlasInitialized = false; diff --git a/modules/ocl/src/precomp.hpp b/modules/ocl/src/precomp.hpp index 3685fbe..25347ff 100644 --- a/modules/ocl/src/precomp.hpp +++ b/modules/ocl/src/precomp.hpp @@ -58,18 +58,6 @@ #include "cvconfig.h" -#if defined(BUILD_SHARED_LIBS) -#if defined WIN32 || defined _WIN32 || defined WINCE -#define CL_RUNTIME_EXPORT __declspec(dllexport) -#elif defined __GNUC__ && __GNUC__ >= 4 -#define CL_RUNTIME_EXPORT __attribute__ ((visibility ("default"))) -#else -#define CL_RUNTIME_EXPORT -#endif -#else -#define CL_RUNTIME_EXPORT -#endif - #include #include #include diff --git a/modules/ocl/src/safe_call.hpp b/modules/ocl/src/safe_call.hpp index e366541..14cbb6d 100644 --- a/modules/ocl/src/safe_call.hpp +++ b/modules/ocl/src/safe_call.hpp @@ -46,7 +46,7 @@ #ifndef __OPENCV_OPENCL_SAFE_CALL_HPP__ #define __OPENCV_OPENCL_SAFE_CALL_HPP__ -#include "opencv2/ocl/cl_runtime/cl_runtime.hpp" +#include "opencv2/core/opencl/runtime/opencl_core.hpp" #define openCLSafeCall(expr) ___openCLSafeCall(expr, __FILE__, __LINE__, CV_Func) #define openCLVerifyCall(res) ___openCLSafeCall(res, __FILE__, __LINE__, CV_Func) diff --git a/modules/ocl/test/test_api.cpp b/modules/ocl/test/test_api.cpp index a54ebb7..fe8bc49 100644 --- a/modules/ocl/test/test_api.cpp +++ b/modules/ocl/test/test_api.cpp @@ -40,7 +40,7 @@ //M*/ #include "test_precomp.hpp" -#include "opencv2/ocl/cl_runtime/cl_runtime.hpp" // for OpenCL types: cl_mem +#include "opencv2/core/opencl/runtime/opencl_core.hpp" // for OpenCL types: cl_mem #include "opencv2/core/ocl.hpp" TEST(TestAPI, openCLExecuteKernelInterop)