mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-01-18 06:03:15 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Merge branch '2.4'

Browse Source
This commit is contained in:
Andrey Kamaev 2013-01-25 16:30:36 +04:00
commit f4e27bcbbc
105 changed files with 4821 additions and 2883 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
CMakeLists.txt
View File

@ -415,10 +415,10 @@ if(WITH_OPENCL)
if(OPENCL_FOUND)
set(HAVE_OPENCL 1)
endif()
if(WITH_OPENCLAMDFFT)
if(WITH_OPENCLAMDFFT AND CLAMDFFT_INCLUDE_DIR)
set(HAVE_CLAMDFFT 1)
endif()
if(WITH_OPENCLAMDBLAS)
if(WITH_OPENCLAMDBLAS AND CLAMDBLAS_INCLUDE_DIR)
set(HAVE_CLAMDBLAS 1)
endif()
endif()

4
android/libinfo/info.c
View File

@ -6,8 +6,8 @@ const char* GetRevision(void);
const char* GetLibraryList(void);
JNIEXPORT jstring JNICALL Java_org_opencv_android_StaticHelper_getLibraryList(JNIEnv *, jclass);
#define PACKAGE_NAME "org.opencv.lib_v" CVAUX_STR(CV_MAJOR_VERSION) CVAUX_STR(CV_MINOR_VERSION) "_" ANDROID_PACKAGE_PLATFORM
#define PACKAGE_REVISION CVAUX_STR(CV_SUBMINOR_VERSION) "." CVAUX_STR(ANDROID_PACKAGE_RELEASE)
#define PACKAGE_NAME "org.opencv.lib_v" CVAUX_STR(CV_VERSION_EPOCH) CVAUX_STR(CV_VERSION_MAJOR) "_" ANDROID_PACKAGE_PLATFORM
#define PACKAGE_REVISION CVAUX_STR(CV_VERSION_MINOR) "." CVAUX_STR(ANDROID_PACKAGE_RELEASE)
const char* GetPackageName(void)
{

2
android/package/CMakeLists.txt
View File

@ -56,7 +56,7 @@ configure_file("${CMAKE_CURRENT_SOURCE_DIR}/${ANDROID_MANIFEST_FILE}" "${PACKAGE
configure_file("${CMAKE_CURRENT_SOURCE_DIR}/res/values/strings.xml" "${PACKAGE_DIR}/res/values/strings.xml" @ONLY)
configure_file("${CMAKE_CURRENT_SOURCE_DIR}/res/drawable/icon.png" "${PACKAGE_DIR}/res/drawable/icon.png" COPYONLY)
set(target_name "OpenCV_${OPENCV_VERSION_MAJOR}.${OPENCV_VERSION_MINOR}.${OPENCV_VERSION_PATCH}_binary_pack_${ANDROID_PACKAGE_PLATFORM}")
set(target_name "OpenCV_${OPENCV_VERSION}_binary_pack_${ANDROID_PACKAGE_PLATFORM}")
get_target_property(opencv_java_location opencv_java LOCATION)
set(android_proj_target_files ${ANDROID_PROJECT_FILES})

4
android/service/engine/AndroidManifest.xml
View File

@ -1,8 +1,8 @@
<?xml version="1.0" encoding="utf-8"?>
<manifest xmlns:android="http://schemas.android.com/apk/res/android"
package="org.opencv.engine"
android:versionCode="24@ANDROID_PLATFORM_VERSION_CODE@"
android:versionName="2.4" >
android:versionCode="25@ANDROID_PLATFORM_VERSION_CODE@"
android:versionName="2.5" >
<uses-sdk android:minSdkVersion="@ANDROID_NATIVE_API_LEVEL@" />
<uses-feature android:name="android.hardware.touchscreen" android:required="false"/>

80
android/service/engine/jni/BinderComponent/OpenCVEngine.cpp
View File

@ -15,60 +15,44 @@ using namespace android;
const int OpenCVEngine::Platform = DetectKnownPlatforms();
const int OpenCVEngine::CpuID = GetCpuID();
const int OpenCVEngine::KnownVersions[] = {2040000, 2040100, 2040200, 2040300, 2040301, 2040302};
std::set<std::string> OpenCVEngine::InitKnownOpenCVersions()
bool OpenCVEngine::ValidateVersion(int version)
{
std::set<std::string> result;
for (size_t i = 0; i < sizeof(KnownVersions)/sizeof(int); i++)
if (KnownVersions[i] == version)
return true;
result.insert("240");
result.insert("241");
result.insert("242");
result.insert("243");
return result;
return false;
}
const std::set<std::string> OpenCVEngine::KnownVersions = InitKnownOpenCVersions();
bool OpenCVEngine::ValidateVersionString(const std::string& version)
int OpenCVEngine::NormalizeVersionString(std::string version)
{
return (KnownVersions.find(version) != KnownVersions.end());
}
std::string OpenCVEngine::NormalizeVersionString(std::string version)
{
std::string result = "";
std::string suffix = "";
int result = 0;
if (version.empty())
{
return result;
}
if (('a' == version[version.size()-1]) || ('b' == version[version.size()-1]))
{
suffix = version[version.size()-1];
version.erase(version.size()-1);
}
std::vector<std::string> parts = SplitStringVector(version, '.');
if (parts.size() >= 2)
// Use only 4 digits of the version, i.e. 1.2.3.4.
// Other digits will be ignored.
if (parts.size() > 4)
parts.erase(parts.begin()+4, parts.end());
int multiplyer = 1000000;
for (std::vector<std::string>::const_iterator it = parts.begin(); it != parts.end(); ++it)
{
if (parts.size() >= 3)
{
result = parts[0] + parts[1] + parts[2] + suffix;
if (!ValidateVersionString(result))
result = "";
}
else
{
result = parts[0] + parts[1] + "0" + suffix;
if (!ValidateVersionString(result))
result = "";
}
int digit = atoi(it->c_str());
result += multiplyer*digit;
multiplyer /= 100;
}
if (!ValidateVersion(result))
result = 0;
return result;
}
@ -86,19 +70,19 @@ int32_t OpenCVEngine::GetVersion()
String16 OpenCVEngine::GetLibPathByVersion(android::String16 version)
{
std::string std_version(String8(version).string());
std::string norm_version;
int norm_version;
std::string path;
LOGD("OpenCVEngine::GetLibPathByVersion(%s) impl", String8(version).string());
norm_version = NormalizeVersionString(std_version);
if (!norm_version.empty())
if (0 != norm_version)
{
path = PackageManager->GetPackagePathByVersion(norm_version, Platform, CpuID);
if (path.empty())
{
LOGI("Package OpenCV of version %s is not installed. Try to install it :)", norm_version.c_str());
LOGI("Package OpenCV of version \"%s\" (%d) is not installed. Try to install it :)", String8(version).string(), norm_version);
}
else
{
@ -107,7 +91,7 @@ String16 OpenCVEngine::GetLibPathByVersion(android::String16 version)
}
else
{
LOGE("OpenCV version \"%s\" (%s) is not supported", String8(version).string(), norm_version.c_str());
LOGE("OpenCV version \"%s\" (%d) is not supported", String8(version).string(), norm_version);
}
return String16(path.c_str());
@ -116,11 +100,11 @@ String16 OpenCVEngine::GetLibPathByVersion(android::String16 version)
android::String16 OpenCVEngine::GetLibraryList(android::String16 version)
{
std::string std_version = String8(version).string();
std::string norm_version;
int norm_version;
String16 result;
norm_version = NormalizeVersionString(std_version);
if (!norm_version.empty())
if (0 != norm_version)
{
std::string tmp = PackageManager->GetPackagePathByVersion(norm_version, Platform, CpuID);
if (!tmp.empty())
@ -156,12 +140,12 @@ android::String16 OpenCVEngine::GetLibraryList(android::String16 version)
}
else
{
LOGI("Package OpenCV of version %s is not installed. Try to install it :)", norm_version.c_str());
LOGI("Package OpenCV of version \"%s\" (%d) is not installed. Try to install it :)", std_version.c_str(), norm_version);
}
}
else
{
LOGE("OpenCV version \"%s\" is not supported", norm_version.c_str());
LOGE("OpenCV version \"%s\" is not supported", std_version.c_str());
}
return result;
@ -170,21 +154,21 @@ android::String16 OpenCVEngine::GetLibraryList(android::String16 version)
bool OpenCVEngine::InstallVersion(android::String16 version)
{
std::string std_version = String8(version).string();
std::string norm_version;
int norm_version;
bool result = false;
LOGD("OpenCVEngine::InstallVersion() begin");
norm_version = NormalizeVersionString(std_version);
if (!norm_version.empty())
if (0 != norm_version)
{
LOGD("PackageManager->InstallVersion call");
result = PackageManager->InstallVersion(norm_version, Platform, CpuID);
}
else
{
LOGE("OpenCV version \"%s\" is not supported", norm_version.c_str());
LOGE("OpenCV version \"%s\" (%d) is not supported", std_version.c_str(), norm_version);
}
LOGD("OpenCVEngine::InstallVersion() end");

9
android/service/engine/jni/BinderComponent/OpenCVEngine.h
View File

@ -23,16 +23,15 @@ public:
protected:
IPackageManager* PackageManager;
static const std::set<std::string> KnownVersions;
static const int KnownVersions[];
OpenCVEngine();
static std::set<std::string> InitKnownOpenCVersions();
bool ValidateVersionString(const std::string& version);
std::string NormalizeVersionString(std::string version);
bool ValidateVersion(int version);
int NormalizeVersionString(std::string version);
bool FixPermissions(const std::string& path);
static const int Platform;
static const int CpuID;
};
#endif
#endif

39
android/service/engine/jni/NativeService/CommonPackageManager.cpp
View File

@ -11,22 +11,24 @@
using namespace std;
set<string> CommonPackageManager::GetInstalledVersions()
vector<int> CommonPackageManager::GetInstalledVersions()
{
set<string> result;
vector<int> result;
vector<PackageInfo> installed_packages = GetInstalledPackages();
for (vector<PackageInfo>::const_iterator it = installed_packages.begin(); it != installed_packages.end(); ++it)
result.resize(installed_packages.size());
for (size_t i = 0; i < installed_packages.size(); i++)
{
string version = it->GetVersion();
assert(!version.empty());
result.insert(version);
int version = installed_packages[i].GetVersion();
assert(version);
result[i] = version;
}
return result;
}
bool CommonPackageManager::CheckVersionInstalled(const std::string& version, int platform, int cpu_id)
bool CommonPackageManager::CheckVersionInstalled(int version, int platform, int cpu_id)
{
bool result = false;
LOGD("CommonPackageManager::CheckVersionInstalled() begin");
@ -48,14 +50,14 @@ bool CommonPackageManager::CheckVersionInstalled(const std::string& version, int
return result;
}
bool CommonPackageManager::InstallVersion(const std::string& version, int platform, int cpu_id)
bool CommonPackageManager::InstallVersion(int version, int platform, int cpu_id)
{
LOGD("CommonPackageManager::InstallVersion() begin");
PackageInfo package(version, platform, cpu_id);
return InstallPackage(package);
}
string CommonPackageManager::GetPackagePathByVersion(const std::string& version, int platform, int cpu_id)
string CommonPackageManager::GetPackagePathByVersion(int version, int platform, int cpu_id)
{
string result;
PackageInfo target_package(version, platform, cpu_id);
@ -64,7 +66,7 @@ string CommonPackageManager::GetPackagePathByVersion(const std::string& version,
for (vector<PackageInfo>::iterator it = all_packages.begin(); it != all_packages.end(); ++it)
{
LOGD("Check version \"%s\" compatibility with \"%s\"\n", version.c_str(), it->GetVersion().c_str());
LOGD("Check version \"%d\" compatibility with \"%d\"\n", version, it->GetVersion());
if (IsVersionCompatible(version, it->GetVersion()))
{
LOGD("Compatible");
@ -79,7 +81,7 @@ string CommonPackageManager::GetPackagePathByVersion(const std::string& version,
if (!packages.empty())
{
int OptRating = -1;
std::string OptVersion = "";
int OptVersion = 0;
std::vector<std::pair<int, int> >& group = CommonPackageManager::ArmRating;
if ((cpu_id & ARCH_X86) || (cpu_id & ARCH_X64))
@ -124,20 +126,13 @@ string CommonPackageManager::GetPackagePathByVersion(const std::string& version,
return result;
}
bool CommonPackageManager::IsVersionCompatible(const std::string& target_version, const std::string& package_version)
bool CommonPackageManager::IsVersionCompatible(int target_version, int package_version)
{
assert (target_version.size() == 3);
assert (package_version.size() == 3);
bool result = false;
assert(target_version);
assert(package_version);
// major version is the same and minor package version is above or the same as target.
if ((package_version[0] == target_version[0]) && (package_version[1] == target_version[1]) && (package_version[2] >= target_version[2]))
{
result = true;
}
return result;
return ( (package_version/10000 == target_version/10000) && (package_version%10000 >= target_version%10000) );
}
int CommonPackageManager::GetHardwareRating(int platform, int cpu_id, const std::vector<std::pair<int, int> >& group)

13
android/service/engine/jni/NativeService/CommonPackageManager.h
View File

@ -3,17 +3,16 @@
#include "IPackageManager.h"
#include "PackageInfo.h"
#include <set>
#include <vector>
#include <string>
class CommonPackageManager: public IPackageManager
{
public:
std::set<std::string> GetInstalledVersions();
bool CheckVersionInstalled(const std::string& version, int platform, int cpu_id);
bool InstallVersion(const std::string& version, int platform, int cpu_id);
std::string GetPackagePathByVersion(const std::string& version, int platform, int cpu_id);
std::vector<int> GetInstalledVersions();
bool CheckVersionInstalled(int version, int platform, int cpu_id);
bool InstallVersion(int version, int platform, int cpu_id);
std::string GetPackagePathByVersion(int version, int platform, int cpu_id);
virtual ~CommonPackageManager();
protected:
@ -23,7 +22,7 @@ protected:
static std::vector<std::pair<int, int> > InitArmRating();
static std::vector<std::pair<int, int> > InitIntelRating();
bool IsVersionCompatible(const std::string& target_version, const std::string& package_version);
bool IsVersionCompatible(int target_version, int package_version);
int GetHardwareRating(int platform, int cpu_id, const std::vector<std::pair<int, int> >& group);
virtual bool InstallPackage(const PackageInfo& package) = 0;
@ -31,4 +30,4 @@ protected:
};
#endif
#endif

46
android/service/engine/jni/NativeService/PackageInfo.cpp
View File

@ -124,14 +124,19 @@ inline int SplitIntelFeatures(const vector<string>& features)
return result;
}
inline string SplitVersion(const vector<string>& features, const string& package_version)
inline int SplitVersion(const vector<string>& features, const string& package_version)
{
string result;
int result = 0;
if ((features.size() > 1) && ('v' == features[1][0]))
{
result = features[1].substr(1);
result += SplitStringVector(package_version, '.')[0];
// Taking major and minor mart of library version from package name
string tmp1 = features[1].substr(1);
result += atoi(tmp1.substr(0,1).c_str())*1000000 + atoi(tmp1.substr(1,1).c_str())*10000;
// Taking release and build number from package revision
vector<string> tmp2 = SplitStringVector(package_version, '.');
result += atoi(tmp2[0].c_str())*100 + atoi(tmp2[1].c_str());
}
else
{
@ -186,9 +191,9 @@ inline int SplitPlatfrom(const vector<string>& features)
* Second part is version. Version starts from "v" symbol. After "v" symbol version nomber without dot symbol added.
* If platform is known third part is platform name
* If platform is unknown it is defined by hardware capabilities using pattern: <arch>_<floating point and vectorization features>_<other features>
* Example: armv7_neon, armv5_vfpv3
* Example: armv7_neon
*/
PackageInfo::PackageInfo(const string& version, int platform, int cpu_id, std::string install_path):
PackageInfo::PackageInfo(int version, int platform, int cpu_id, std::string install_path):
Version(version),
Platform(platform),
CpuID(cpu_id),
@ -198,7 +203,14 @@ InstallPath("")
Platform = PLATFORM_UNKNOWN;
#endif
FullName = BasePackageName + "_v" + Version.substr(0, Version.size()-1);
int major_version = version/1000000;
int minor_version = version/10000 - major_version*100;
char tmp[32];
sprintf(tmp, "%d%d", major_version, minor_version);
FullName = BasePackageName + std::string("_v") + std::string(tmp);
if (PLATFORM_UNKNOWN != Platform)
{
FullName += string("_") + JoinPlatform(platform);
@ -296,7 +308,7 @@ InstallPath("")
else
{
LOGD("PackageInfo::PackageInfo: package arch unknown");
Version.clear();
Version = 0;
CpuID = ARCH_UNKNOWN;
Platform = PLATFORM_UNKNOWN;
}
@ -304,7 +316,7 @@ InstallPath("")
else
{
LOGD("PackageInfo::PackageInfo: package arch unknown");
Version.clear();
Version = 0;
CpuID = ARCH_UNKNOWN;
Platform = PLATFORM_UNKNOWN;
}
@ -371,7 +383,7 @@ InstallPath(install_path)
{
LOGI("Info library not found in package");
LOGI("OpenCV Manager package does not contain any verison of OpenCV library");
Version.clear();
Version = 0;
CpuID = ARCH_UNKNOWN;
Platform = PLATFORM_UNKNOWN;
return;
@ -383,7 +395,7 @@ InstallPath(install_path)
if (!features.empty() && (BasePackageName == features[0]))
{
Version = SplitVersion(features, package_version);
if (Version.empty())
if (0 == Version)
{
CpuID = ARCH_UNKNOWN;
Platform = PLATFORM_UNKNOWN;
@ -410,7 +422,7 @@ InstallPath(install_path)
if (features.size() < 3)
{
LOGD("It is not OpenCV library package for this platform");
Version.clear();
Version = 0;
CpuID = ARCH_UNKNOWN;
Platform = PLATFORM_UNKNOWN;
return;
@ -444,7 +456,7 @@ InstallPath(install_path)
else
{
LOGD("It is not OpenCV library package for this platform");
Version.clear();
Version = 0;
CpuID = ARCH_UNKNOWN;
Platform = PLATFORM_UNKNOWN;
return;
@ -454,7 +466,7 @@ InstallPath(install_path)
else
{
LOGD("It is not OpenCV library package for this platform");
Version.clear();
Version = 0;
CpuID = ARCH_UNKNOWN;
Platform = PLATFORM_UNKNOWN;
return;
@ -463,7 +475,7 @@ InstallPath(install_path)
bool PackageInfo::IsValid() const
{
return !(Version.empty() && (PLATFORM_UNKNOWN == Platform) && (ARCH_UNKNOWN == CpuID));
return !((0 == Version) && (PLATFORM_UNKNOWN == Platform) && (ARCH_UNKNOWN == CpuID));
}
int PackageInfo::GetPlatform() const
@ -481,7 +493,7 @@ string PackageInfo::GetFullName() const
return FullName;
}
string PackageInfo::GetVersion() const
int PackageInfo::GetVersion() const
{
return Version;
}
@ -494,4 +506,4 @@ string PackageInfo::GetInstalationPath() const
bool PackageInfo::operator==(const PackageInfo& package) const
{
return (package.FullName == FullName);
}
}

8
android/service/engine/jni/NativeService/PackageInfo.h
View File

@ -30,10 +30,10 @@
class PackageInfo
{
public:
PackageInfo(const std::string& version, int platform, int cpu_id, std::string install_path = "/data/data/");
PackageInfo(int version, int platform, int cpu_id, std::string install_path = "/data/data/");
PackageInfo(const std::string& fullname, const std::string& install_path, std::string package_version = "0.0");
std::string GetFullName() const;
std::string GetVersion() const;
int GetVersion() const;
int GetPlatform() const;
int GetCpuID() const;
std::string GetInstalationPath() const;
@ -43,7 +43,7 @@ public:
protected:
static std::map<int, std::string> InitPlatformNameMap();
std::string Version;
int Version;
int Platform;
int CpuID;
std::string FullName;
@ -51,4 +51,4 @@ protected:
static const std::string BasePackageName;
};
#endif
#endif

12
android/service/engine/jni/include/IPackageManager.h
View File

@ -1,17 +1,17 @@
#ifndef __IPACKAGE_MANAGER__
#define __IPACKAGE_MANAGER__
#include <set>
#include <vector>
#include <string>
class IPackageManager
{
public:
virtual std::set<std::string> GetInstalledVersions() = 0;
virtual bool CheckVersionInstalled(const std::string& version, int platform, int cpu_id) = 0;
virtual bool InstallVersion(const std::string&, int platform, int cpu_id) = 0;
virtual std::string GetPackagePathByVersion(const std::string&, int platform, int cpu_id) = 0;
virtual std::vector<int> GetInstalledVersions() = 0;
virtual bool CheckVersionInstalled(int version, int platform, int cpu_id) = 0;
virtual bool InstallVersion(int version, int platform, int cpu_id) = 0;
virtual std::string GetPackagePathByVersion(int version, int platform, int cpu_id) = 0;
virtual ~IPackageManager(){};
};
#endif
#endif

2
android/service/engine/res/layout/info.xml
View File

@ -26,7 +26,7 @@
android:id="@+id/textView1"
android:layout_width="wrap_content"
android:layout_height="wrap_content"
android:text="Version: "
android:text="Library version: "
android:textAppearance="?android:attr/textAppearanceSmall" />
<TextView

2
android/service/engine/res/layout/main.xml
View File

@ -30,7 +30,7 @@
android:id="@+id/EngineVersionCaption"
android:layout_width="wrap_content"
android:layout_height="wrap_content"
android:text="Version: "
android:text="OpenCV Manager version: "
android:textAppearance="?android:attr/textAppearanceMedium" />
<TextView

2
cmake/OpenCVDetectAndroidSDK.cmake
View File

@ -302,7 +302,7 @@ macro(add_android_project target path)
COMMAND ${CMAKE_COMMAND} -E touch "${android_proj_bin_dir}/bin/${target}-debug.apk" # needed because ant does not update the timestamp of updated apk
WORKING_DIRECTORY "${android_proj_bin_dir}"
MAIN_DEPENDENCY "${android_proj_bin_dir}/${ANDROID_MANIFEST_FILE}"
DEPENDS "${OpenCV_BINARY_DIR}/bin/.classes.jar.dephelper" opencv_java # as we are part of OpenCV we can just force this dependency
DEPENDS "${OpenCV_BINARY_DIR}/bin/classes.jar.dephelper" opencv_java # as we are part of OpenCV we can just force this dependency
DEPENDS ${android_proj_file_deps} ${JNI_LIB_NAME})
endif()

212
cmake/OpenCVDetectOpenCL.cmake
View File

@ -1,78 +1,146 @@
if(APPLE)
set(OPENCL_FOUND YES)
set(OPENCL_LIBRARIES "-framework OpenCL")
set(OPENCL_FOUND YES)
set(OPENCL_LIBRARIES "-framework OpenCL")
else()
#find_package(OpenCL QUIET)
if(WITH_OPENCLAMDFFT)
find_path(CLAMDFFT_INCLUDE_DIR
NAMES clAmdFft.h)
find_library(CLAMDFFT_LIBRARIES
NAMES clAmdFft.Runtime)
find_package(OpenCL QUIET)
if(WITH_OPENCLAMDFFT)
set(CLAMDFFT_SEARCH_PATH $ENV{CLAMDFFT_PATH})
if(NOT CLAMDFFT_SEARCH_PATH)
if(WIN32)
set( CLAMDFFT_SEARCH_PATH "C:\\Program Files (x86)\\AMD\\clAmdFft" )
endif()
endif()
if(WITH_OPENCLAMDBLAS)
find_path(CLAMDBLAS_INCLUDE_DIR
NAMES clAmdBlas.h)
find_library(CLAMDBLAS_LIBRARIES
NAMES clAmdBlas)
endif()
# Try AMD/ATI Stream SDK
if (NOT OPENCL_FOUND)
set(ENV_AMDSTREAMSDKROOT $ENV{AMDAPPSDKROOT})
set(ENV_OPENCLROOT $ENV{OPENCLROOT})
set(ENV_CUDA_PATH $ENV{CUDA_PATH})
if(ENV_AMDSTREAMSDKROOT)
set(OPENCL_INCLUDE_SEARCH_PATH ${ENV_AMDSTREAMSDKROOT}/include)
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} ${ENV_AMDSTREAMSDKROOT}/lib/x86)
else()
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} ${ENV_AMDSTREAMSDKROOT}/lib/x86_64)
endif()
elseif(ENV_CUDA_PATH AND WIN32)
set(OPENCL_INCLUDE_SEARCH_PATH ${ENV_CUDA_PATH}/include)
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} ${ENV_CUDA_PATH}/lib/Win32)
else()
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} ${ENV_CUDA_PATH}/lib/x64)
endif()
elseif(ENV_OPENCLROOT AND UNIX)
set(OPENCL_INCLUDE_SEARCH_PATH ${ENV_OPENCLROOT}/inc)
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} /usr/lib)
else()
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} /usr/lib64)
endif()
endif()
if(OPENCL_INCLUDE_SEARCH_PATH)
find_path(OPENCL_INCLUDE_DIR
NAMES CL/cl.h OpenCL/cl.h
PATHS ${OPENCL_INCLUDE_SEARCH_PATH}
NO_DEFAULT_PATH)
else()
find_path(OPENCL_INCLUDE_DIR
NAMES CL/cl.h OpenCL/cl.h)
endif()
if(OPENCL_LIB_SEARCH_PATH)
find_library(OPENCL_LIBRARY NAMES OpenCL PATHS ${OPENCL_LIB_SEARCH_PATH} NO_DEFAULT_PATH)
else()
find_library(OPENCL_LIBRARY NAMES OpenCL)
endif()
include(FindPackageHandleStandardArgs)
find_package_handle_standard_args(
OPENCL
DEFAULT_MSG
OPENCL_LIBRARY OPENCL_INCLUDE_DIR
)
if(OPENCL_FOUND)
set(OPENCL_LIBRARIES ${OPENCL_LIBRARY})
set(HAVE_OPENCL 1)
else()
set(OPENCL_LIBRARIES)
endif()
set( CLAMDFFT_INCLUDE_SEARCH_PATH ${CLAMDFFT_SEARCH_PATH}/include )
if(UNIX)
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
set(CLAMDFFT_LIB_SEARCH_PATH /usr/lib)
else()
set(CLAMDFFT_LIB_SEARCH_PATH /usr/lib64)
endif()
else()
set(HAVE_OPENCL 1)
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
set(CLAMDFFT_LIB_SEARCH_PATH ${CLAMDFFT_SEARCH_PATH}\\lib32\\import)
else()
set(CLAMDFFT_LIB_SEARCH_PATH ${CLAMDFFT_SEARCH_PATH}\\lib64\\import)
endif()
endif()
find_path(CLAMDFFT_INCLUDE_DIR
NAMES clAmdFft.h
PATHS ${CLAMDFFT_INCLUDE_SEARCH_PATH}
PATH_SUFFIXES clAmdFft
NO_DEFAULT_PATH)
find_library(CLAMDFFT_LIBRARY
NAMES clAmdFft.Runtime
PATHS ${CLAMDFFT_LIB_SEARCH_PATH}
NO_DEFAULT_PATH)
if(CLAMDFFT_LIBRARY)
set(CLAMDFFT_LIBRARIES ${CLAMDFFT_LIBRARY})
else()
set(CLAMDFFT_LIBRARIES "")
endif()
endif()
if(WITH_OPENCLAMDBLAS)
set(CLAMDBLAS_SEARCH_PATH $ENV{CLAMDBLAS_PATH})
if(NOT CLAMDBLAS_SEARCH_PATH)
if(WIN32)
set( CLAMDBLAS_SEARCH_PATH "C:\\Program Files (x86)\\AMD\\clAmdBlas" )
endif()
endif()
set( CLAMDBLAS_INCLUDE_SEARCH_PATH ${CLAMDBLAS_SEARCH_PATH}/include )
if(UNIX)
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
set(CLAMDBLAS_LIB_SEARCH_PATH /usr/lib)
else()
set(CLAMDBLAS_LIB_SEARCH_PATH /usr/lib64)
endif()
else()
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
set(CLAMDBLAS_LIB_SEARCH_PATH ${CLAMDBLAS_SEARCH_PATH}\\lib32\\import)
else()
set(CLAMDBLAS_LIB_SEARCH_PATH ${CLAMDBLAS_SEARCH_PATH}\\lib64\\import)
endif()
endif()
find_path(CLAMDBLAS_INCLUDE_DIR
NAMES clAmdBlas.h
PATHS ${CLAMDBLAS_INCLUDE_SEARCH_PATH}
PATH_SUFFIXES clAmdBlas
NO_DEFAULT_PATH)
find_library(CLAMDBLAS_LIBRARY
NAMES clAmdBlas
PATHS ${CLAMDBLAS_LIB_SEARCH_PATH}
NO_DEFAULT_PATH)
if(CLAMDBLAS_LIBRARY)
set(CLAMDBLAS_LIBRARIES ${CLAMDBLAS_LIBRARY})
else()
set(CLAMDBLAS_LIBRARIES "")
endif()
endif()
# Try AMD/ATI Stream SDK
if (NOT OPENCL_FOUND)
set(ENV_AMDSTREAMSDKROOT $ENV{AMDAPPSDKROOT})
set(ENV_AMDAPPSDKROOT $ENV{AMDAPPSDKROOT})
set(ENV_OPENCLROOT $ENV{OPENCLROOT})
set(ENV_CUDA_PATH $ENV{CUDA_PATH})
if(ENV_AMDSTREAMSDKROOT)
set(OPENCL_INCLUDE_SEARCH_PATH ${ENV_AMDAPPSDKROOT}/include)
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} ${ENV_AMDAPPSDKROOT}/lib/x86)
else()
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} ${ENV_AMDAPPSDKROOT}/lib/x86_64)
endif()
elseif(ENV_AMDSTREAMSDKROOT)
set(OPENCL_INCLUDE_SEARCH_PATH ${ENV_AMDSTREAMSDKROOT}/include)
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} ${ENV_AMDSTREAMSDKROOT}/lib/x86)
else()
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} ${ENV_AMDSTREAMSDKROOT}/lib/x86_64)
endif()
elseif(ENV_CUDA_PATH AND WIN32)
set(OPENCL_INCLUDE_SEARCH_PATH ${ENV_CUDA_PATH}/include)
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} ${ENV_CUDA_PATH}/lib/Win32)
else()
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} ${ENV_CUDA_PATH}/lib/x64)
endif()
elseif(ENV_OPENCLROOT AND UNIX)
set(OPENCL_INCLUDE_SEARCH_PATH ${ENV_OPENCLROOT}/inc)
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} /usr/lib)
else()
set(OPENCL_LIB_SEARCH_PATH ${OPENCL_LIB_SEARCH_PATH} /usr/lib64)
endif()
endif()
if(OPENCL_INCLUDE_SEARCH_PATH)
find_path(OPENCL_INCLUDE_DIR
NAMES CL/cl.h OpenCL/cl.h
PATHS ${OPENCL_INCLUDE_SEARCH_PATH}
NO_DEFAULT_PATH)
else()
find_path(OPENCL_INCLUDE_DIR
NAMES CL/cl.h OpenCL/cl.h)
endif()
if(OPENCL_LIB_SEARCH_PATH)
find_library(OPENCL_LIBRARY NAMES OpenCL PATHS ${OPENCL_LIB_SEARCH_PATH} NO_DEFAULT_PATH)
else()
find_library(OPENCL_LIBRARY NAMES OpenCL)
endif()
include(FindPackageHandleStandardArgs)
find_package_handle_standard_args(
OPENCL
DEFAULT_MSG
OPENCL_LIBRARY OPENCL_INCLUDE_DIR
)
if(OPENCL_FOUND)
set(OPENCL_LIBRARIES ${OPENCL_LIBRARY})
set(HAVE_OPENCL 1)
else()
set(OPENCL_LIBRARIES)
endif()
else()
set(HAVE_OPENCL 1)
endif()
endif()

2
cmake/OpenCVExtraTargets.cmake
View File

@ -16,7 +16,7 @@ endif()
# Source package, for "make package_source"
# ----------------------------------------------------------------------------
if(BUILD_PACKAGE)
set(TARBALL_NAME "${CMAKE_PROJECT_NAME}-${OPENCV_VERSION_MAJOR}.${OPENCV_VERSION_MINOR}.${OPENCV_VERSION_PATCH}")
set(TARBALL_NAME "${CMAKE_PROJECT_NAME}-${OPENCV_VERSION}")
if (NOT WIN32)
if(APPLE)
set(TAR_CMD gnutar)

3
cmake/OpenCVModule.cmake
View File

@ -461,6 +461,7 @@ macro(ocv_create_module)
OUTPUT_NAME "${the_module}${OPENCV_DLLVERSION}"
DEBUG_POSTFIX "${OPENCV_DEBUG_POSTFIX}"
ARCHIVE_OUTPUT_DIRECTORY ${LIBRARY_OUTPUT_PATH}
LIBRARY_OUTPUT_DIRECTORY ${LIBRARY_OUTPUT_PATH}
RUNTIME_OUTPUT_DIRECTORY ${EXECUTABLE_OUTPUT_PATH}
INSTALL_NAME_DIR lib
)
@ -470,7 +471,7 @@ macro(ocv_create_module)
# Android SDK build scripts can include only .so files into final .apk
# As result we should not set version properties for Android
set_target_properties(${the_module} PROPERTIES
VERSION ${OPENCV_VERSION}
VERSION ${OPENCV_LIBVERSION}
SOVERSION ${OPENCV_SOVERSION}
)
endif()

14
cmake/OpenCVVersion.cmake
View File

@ -1,12 +1,18 @@
SET(OPENCV_VERSION_FILE "${CMAKE_CURRENT_SOURCE_DIR}/modules/core/include/opencv2/core/version.hpp")
FILE(STRINGS "${OPENCV_VERSION_FILE}" OPENCV_VERSION_PARTS REGEX "#define CV_.+OR_VERSION[ ]+[0-9]+" )
FILE(STRINGS "${OPENCV_VERSION_FILE}" OPENCV_VERSION_PARTS REGEX "#define CV_VERSION_[A-Z]+[ ]+[0-9]+" )
string(REGEX REPLACE ".+CV_MAJOR_VERSION[ ]+([0-9]+).*" "\\1" OPENCV_VERSION_MAJOR "${OPENCV_VERSION_PARTS}")
string(REGEX REPLACE ".+CV_MINOR_VERSION[ ]+([0-9]+).*" "\\1" OPENCV_VERSION_MINOR "${OPENCV_VERSION_PARTS}")
string(REGEX REPLACE ".+CV_SUBMINOR_VERSION[ ]+([0-9]+).*" "\\1" OPENCV_VERSION_PATCH "${OPENCV_VERSION_PARTS}")
string(REGEX REPLACE ".+CV_VERSION_EPOCH[ ]+([0-9]+).*" "\\1" OPENCV_VERSION_MAJOR "${OPENCV_VERSION_PARTS}")
string(REGEX REPLACE ".+CV_VERSION_MAJOR[ ]+([0-9]+).*" "\\1" OPENCV_VERSION_MINOR "${OPENCV_VERSION_PARTS}")
string(REGEX REPLACE ".+CV_VERSION_MINOR[ ]+([0-9]+).*" "\\1" OPENCV_VERSION_PATCH "${OPENCV_VERSION_PARTS}")
string(REGEX REPLACE ".+CV_VERSION_REVISION[ ]+([0-9]+).*" "\\1" OPENCV_VERSION_TWEAK "${OPENCV_VERSION_PARTS}")
set(OPENCV_VERSION "${OPENCV_VERSION_MAJOR}.${OPENCV_VERSION_MINOR}.${OPENCV_VERSION_PATCH}")
if(OPENCV_VERSION_TWEAK GREATER 0)
set(OPENCV_VERSION "${OPENCV_VERSION}.${OPENCV_VERSION_TWEAK}")
endif()
set(OPENCV_SOVERSION "${OPENCV_VERSION_MAJOR}.${OPENCV_VERSION_MINOR}")
set(OPENCV_LIBVERSION "${OPENCV_VERSION_MAJOR}.${OPENCV_VERSION_MINOR}.${OPENCV_VERSION_PATCH}")
# create a dependency on version file
# we never use output of the following command but cmake will rerun automatically if the version file changes

12
cmake/templates/OpenCVConfig.cmake.in
View File

@ -22,10 +22,11 @@
# - OpenCV_INCLUDE_DIRS : The OpenCV include directories.
# - OpenCV_COMPUTE_CAPABILITIES : The version of compute capability
# - OpenCV_ANDROID_NATIVE_API_LEVEL : Minimum required level of Android API
# - OpenCV_VERSION : The version of this OpenCV build. Example: "@OPENCV_VERSION@"
# - OpenCV_VERSION_MAJOR : Major version part of OpenCV_VERSION. Example: "@OPENCV_VERSION_MAJOR@"
# - OpenCV_VERSION_MINOR : Minor version part of OpenCV_VERSION. Example: "@OPENCV_VERSION_MINOR@"
# - OpenCV_VERSION_PATCH : Patch version part of OpenCV_VERSION. Example: "@OPENCV_VERSION_PATCH@"
# - OpenCV_VERSION : The version of this OpenCV build: "@OPENCV_VERSION@"
# - OpenCV_VERSION_MAJOR : Major version part of OpenCV_VERSION: "@OPENCV_VERSION_MAJOR@"
# - OpenCV_VERSION_MINOR : Minor version part of OpenCV_VERSION: "@OPENCV_VERSION_MINOR@"
# - OpenCV_VERSION_PATCH : Patch version part of OpenCV_VERSION: "@OPENCV_VERSION_PATCH@"
# - OpenCV_VERSION_TWEAK : Tweak version part of OpenCV_VERSION: "@OPENCV_VERSION_TWEAK@"
#
# Advanced variables:
# - OpenCV_SHARED
@ -99,6 +100,7 @@ SET(OpenCV_VERSION @OPENCV_VERSION@)
SET(OpenCV_VERSION_MAJOR @OPENCV_VERSION_MAJOR@)
SET(OpenCV_VERSION_MINOR @OPENCV_VERSION_MINOR@)
SET(OpenCV_VERSION_PATCH @OPENCV_VERSION_PATCH@)
SET(OpenCV_VERSION_TWEAK @OPENCV_VERSION_TWEAK@)
# ====================================================================
# Link libraries: e.g. libopencv_core.so, opencv_imgproc220d.lib, etc...
@ -183,7 +185,7 @@ set(OpenCV_FIND_COMPONENTS ${OpenCV_FIND_COMPONENTS_})
# Resolve dependencies
# ==============================================================
if(OpenCV_USE_MANGLED_PATHS)
set(OpenCV_LIB_SUFFIX ".${OpenCV_VERSION}")
set(OpenCV_LIB_SUFFIX ".${OpenCV_VERSION_MAJOR}.${OpenCV_VERSION_MINOR}.${OpenCV_VERSION_PATCH}")
else()
set(OpenCV_LIB_SUFFIX "")
endif()

25
doc/CMakeLists.txt
View File

@ -60,7 +60,7 @@ if(BUILD_DOCS AND HAVE_SPHINX)
configure_file("${OpenCV_SOURCE_DIR}/modules/refman.rst.in" "${OpenCV_SOURCE_DIR}/modules/refman.rst" IMMEDIATE @ONLY)
file(GLOB_RECURSE OPENCV_FILES_UG user_guide/*.rst)
file(GLOB_RECURSE OPENCV_FILES_UG user_guide/*.rst)
file(GLOB_RECURSE OPENCV_FILES_TUT tutorials/*.rst)
file(GLOB_RECURSE OPENCV_FILES_TUT_PICT tutorials/*.png tutorials/*.jpg)
@ -74,14 +74,21 @@ if(BUILD_DOCS AND HAVE_SPHINX)
COMMAND ${CMAKE_COMMAND} -E copy_if_different ${CMAKE_CURRENT_SOURCE_DIR}/mymath.sty ${CMAKE_CURRENT_BINARY_DIR}
COMMAND ${PYTHON_EXECUTABLE} "${CMAKE_CURRENT_SOURCE_DIR}/patch_refman_latex.py" opencv2refman.tex
COMMAND ${PYTHON_EXECUTABLE} "${CMAKE_CURRENT_SOURCE_DIR}/patch_refman_latex.py" opencv2manager.tex
COMMAND ${PDFLATEX_COMPILER} opencv2refman.tex
COMMAND ${PDFLATEX_COMPILER} opencv2refman.tex
COMMAND ${PDFLATEX_COMPILER} opencv2manager.tex
COMMAND ${PDFLATEX_COMPILER} opencv2manager.tex
COMMAND ${PDFLATEX_COMPILER} opencv_user.tex
COMMAND ${PDFLATEX_COMPILER} opencv_user.tex
COMMAND ${PDFLATEX_COMPILER} opencv_tutorials.tex
COMMAND ${PDFLATEX_COMPILER} opencv_tutorials.tex
COMMAND ${CMAKE_COMMAND} -E echo "Generating opencv2refman.pdf"
COMMAND ${PDFLATEX_COMPILER} -interaction=batchmode opencv2refman.tex
COMMAND ${PDFLATEX_COMPILER} -interaction=batchmode opencv2refman.tex
COMMAND ${CMAKE_COMMAND} -E echo "Generating opencv2manager.pdf"
COMMAND ${PDFLATEX_COMPILER} -interaction=batchmode opencv2manager.tex
COMMAND ${PDFLATEX_COMPILER} -interaction=batchmode opencv2manager.tex
COMMAND ${CMAKE_COMMAND} -E echo "Generating opencv_user.pdf"
COMMAND ${PDFLATEX_COMPILER} -interaction=batchmode opencv_user.tex
COMMAND ${PDFLATEX_COMPILER} -interaction=batchmode opencv_user.tex
COMMAND ${CMAKE_COMMAND} -E echo "Generating opencv_tutorials.pdf"
COMMAND ${PDFLATEX_COMPILER} -interaction=batchmode opencv_tutorials.tex
COMMAND ${PDFLATEX_COMPILER} -interaction=batchmode opencv_tutorials.tex
COMMAND ${CMAKE_COMMAND} -E echo "Generating opencv_cheatsheet.pdf"
COMMAND ${PDFLATEX_COMPILER} -interaction=batchmode "${CMAKE_CURRENT_SOURCE_DIR}/opencv_cheatsheet.tex"
COMMAND ${PDFLATEX_COMPILER} -interaction=batchmode "${CMAKE_CURRENT_SOURCE_DIR}/opencv_cheatsheet.tex"
DEPENDS ${OPENCV_DOC_DEPS}
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
COMMENT "Generating the PDF Manuals"

15
doc/conf.py
View File

@ -44,21 +44,24 @@ master_doc = 'index'
# General information about the project.
project = u'OpenCV'
copyright = u'2011-2012, opencv dev team'
copyright = u'2011-2013, opencv dev team'
# The version info for the project you're documenting, acts as replacement for
# |version| and |release|, also used in various other places throughout the
# built documents.
version_file = open("../modules/core/include/opencv2/core/version.hpp", "rt").read()
version_major = re.search("^W*#\W*define\W+CV_MAJOR_VERSION\W+(\d+)\W*$", version_file, re.MULTILINE).group(1)
version_minor = re.search("^W*#\W*define\W+CV_MINOR_VERSION\W+(\d+)\W*$", version_file, re.MULTILINE).group(1)
version_patch = re.search("^W*#\W*define\W+CV_SUBMINOR_VERSION\W+(\d+)\W*$", version_file, re.MULTILINE).group(1)
version_epoch = re.search("^W*#\W*define\W+CV_VERSION_EPOCH\W+(\d+)\W*$", version_file, re.MULTILINE).group(1)
version_major = re.search("^W*#\W*define\W+CV_VERSION_MAJOR\W+(\d+)\W*$", version_file, re.MULTILINE).group(1)
version_minor = re.search("^W*#\W*define\W+CV_VERSION_MINOR\W+(\d+)\W*$", version_file, re.MULTILINE).group(1)
version_patch = re.search("^W*#\W*define\W+CV_VERSION_REVISION\W+(\d+)\W*$", version_file, re.MULTILINE).group(1)
# The short X.Y version.
version = version_major + '.' + version_minor
version = version_epoch + '.' + version_major
# The full version, including alpha/beta/rc tags.
release = version_major + '.' + version_minor + '.' + version_patch
release = version_epoch + '.' + version_major + '.' + version_minor
if version_patch:
release = release + '.' + version_patch
# The language for content autogenerated by Sphinx. Refer to documentation
# for a list of supported languages.

1
doc/opencv_cheatsheet.tex
View File

@ -67,6 +67,7 @@
\usepackage[pdftex]{color,graphicx}
\usepackage[landscape]{geometry}
\usepackage{hyperref}
\usepackage[T1]{fontenc}
\hypersetup{colorlinks=true, filecolor=black, linkcolor=black, urlcolor=blue, citecolor=black}
\graphicspath{{./images/}}

2
doc/user_guide/ug_mat.rst
View File

@ -144,7 +144,7 @@ A call to ``waitKey()`` starts a message passing cycle that waits for a key stro
Mat img = imread("image.jpg");
Mat grey;
cvtColor(img, grey, CV_BGR2GREY);
cvtColor(img, grey, CV_BGR2GRAY);
Mat sobelx;
Sobel(grey, sobelx, CV_32F, 1, 0);

2
modules/androidcamera/src/camera_activity.cpp
View File

@ -269,7 +269,7 @@ void CameraWrapperConnector::fillListWrapperLibs(const string& folderPath, vecto
std::string CameraWrapperConnector::getDefaultPathLibFolder()
{
#define BIN_PACKAGE_NAME(x) "org.opencv.lib_v" CVAUX_STR(CV_MAJOR_VERSION) CVAUX_STR(CV_MINOR_VERSION) "_" x
#define BIN_PACKAGE_NAME(x) "org.opencv.lib_v" CVAUX_STR(CV_VERSION_EPOCH) CVAUX_STR(CV_VERSION_MAJOR) "_" x
const char* const packageList[] = {BIN_PACKAGE_NAME("armv7a"), OPENCV_ENGINE_PACKAGE};
for (size_t i = 0; i < sizeof(packageList)/sizeof(packageList[0]); i++)
{

19
modules/core/include/opencv2/core/version.hpp
View File

@ -47,12 +47,23 @@
#ifndef __OPENCV_VERSION_HPP__
#define __OPENCV_VERSION_HPP__
#define CV_MAJOR_VERSION 2
#define CV_MINOR_VERSION 4
#define CV_SUBMINOR_VERSION 9
#define CV_VERSION_EPOCH 2
#define CV_VERSION_MAJOR 4
#define CV_VERSION_MINOR 9
#define CV_VERSION_REVISION 0
#define CVAUX_STR_EXP(__A) #__A
#define CVAUX_STR(__A) CVAUX_STR_EXP(__A)
#define CV_VERSION CVAUX_STR(CV_MAJOR_VERSION) "." CVAUX_STR(CV_MINOR_VERSION) "." CVAUX_STR(CV_SUBMINOR_VERSION)
#if CV_VERSION_REVISION
# define CV_VERSION CVAUX_STR(CV_VERSION_EPOCH) "." CVAUX_STR(CV_VERSION_MAJOR) "." CVAUX_STR(CV_VERSION_MINOR) "." CVAUX_STR(CV_VERSION_REVISION)
#else
# define CV_VERSION CVAUX_STR(CV_VERSION_EPOCH) "." CVAUX_STR(CV_VERSION_MAJOR) "." CVAUX_STR(CV_VERSION_MINOR)
#endif
/* old style version constants*/
#define CV_MAJOR_VERSION CV_VERSION_EPOCH
#define CV_MINOR_VERSION CV_VERSION_MAJOR
#define CV_SUBMINOR_VERSION CV_VERSION_MINOR
#endif

2
modules/core/src/system.cpp
View File

@ -439,7 +439,7 @@ void error( const Exception& exc )
exc.func.c_str() : "unknown function", exc.file.c_str(), exc.line );
fprintf( stderr, "%s\n", buf );
fflush( stderr );
# ifdef ANDROID
# ifdef __ANDROID__
__android_log_print(ANDROID_LOG_ERROR, "cv::error()", "%s", buf);
# endif
}

7
modules/features2d/src/draw.cpp
View File

@ -200,10 +200,13 @@ void drawMatches( const Mat& img1, const vector<KeyPoint>& keypoints1,
// draw matches
for( size_t m = 0; m < matches1to2.size(); m++ )
{
int i1 = matches1to2[m].queryIdx;
int i2 = matches1to2[m].trainIdx;
if( matchesMask.empty() || matchesMask[m] )
{
int i1 = matches1to2[m].queryIdx;
int i2 = matches1to2[m].trainIdx;
CV_Assert(i1 >= 0 && i1 < static_cast<int>(keypoints1.size()));
CV_Assert(i2 >= 0 && i2 < static_cast<int>(keypoints2.size()));
const KeyPoint &kp1 = keypoints1[i1], &kp2 = keypoints2[i2];
_drawMatch( outImg, outImg1, outImg2, kp1, kp2, matchColor, flags );
}

2
modules/gpu/src/ffmpeg_video_source.cpp
View File

@ -64,7 +64,7 @@ namespace
{
#if defined WIN32 || defined _WIN32
const char* module_name = "opencv_ffmpeg"
CVAUX_STR(CV_MAJOR_VERSION) CVAUX_STR(CV_MINOR_VERSION) CVAUX_STR(CV_SUBMINOR_VERSION)
CVAUX_STR(CV_VERSION_EPOCH) CVAUX_STR(CV_VERSION_MAJOR) CVAUX_STR(CV_VERSION_MINOR)
#if (defined _MSC_VER && defined _M_X64) || (defined __GNUC__ && defined __x86_64__)
"_64"
#endif

2
modules/gpu/src/video_writer.cpp
View File

@ -767,7 +767,7 @@ namespace
{
#if defined WIN32 || defined _WIN32
const char* module_name = "opencv_ffmpeg"
CVAUX_STR(CV_MAJOR_VERSION) CVAUX_STR(CV_MINOR_VERSION) CVAUX_STR(CV_SUBMINOR_VERSION)
CVAUX_STR(CV_VERSION_EPOCH) CVAUX_STR(CV_VERSION_MAJOR) CVAUX_STR(CV_VERSION_MINOR)
#if (defined _MSC_VER && defined _M_X64) || (defined __GNUC__ && defined __x86_64__)
"_64"
#endif

2
modules/highgui/src/cap_ffmpeg.cpp
View File

@ -86,7 +86,7 @@ private:
{
#if defined WIN32 || defined _WIN32
const char* module_name = "opencv_ffmpeg"
CVAUX_STR(CV_MAJOR_VERSION) CVAUX_STR(CV_MINOR_VERSION) CVAUX_STR(CV_SUBMINOR_VERSION)
CVAUX_STR(CV_VERSION_EPOCH) CVAUX_STR(CV_VERSION_MAJOR) CVAUX_STR(CV_VERSION_MINOR)
#if (defined _MSC_VER && defined _M_X64) || (defined __GNUC__ && defined __x86_64__)
"_64"
#endif

20
modules/imgproc/perf/perf_resize.cpp
View File

@ -21,16 +21,14 @@ PERF_TEST_P(MatInfo_Size_Size, resizeUpLinear,
Size from = get<1>(GetParam());
Size to = get<2>(GetParam());
cv::Mat src(from, matType);
cv::Mat dst(to, matType);
declare.in(src, WARMUP_RNG).out(dst);
cv::Mat src(from, matType), dst(to, matType);
cvtest::fillGradient(src);
declare.in(src).out(dst);
TEST_CYCLE() resize(src, dst, to);
// Test case temporary disabled for Android Platform
#ifdef ANDROID
SANITY_CHECK(dst, 255); // TODO: Reimplement check in future versions
SANITY_CHECK(dst, 5);
#else
SANITY_CHECK(dst, 1 + 1e-6);
#endif
@ -50,16 +48,14 @@ PERF_TEST_P(MatInfo_Size_Size, resizeDownLinear,
Size from = get<1>(GetParam());
Size to = get<2>(GetParam());
cv::Mat src(from, matType);
cv::Mat dst(to, matType);
declare.in(src, WARMUP_RNG).out(dst);
cv::Mat src(from, matType), dst(to, matType);
cvtest::fillGradient(src);
declare.in(src).out(dst);
TEST_CYCLE() resize(src, dst, to);
// Test case temporary disabled for Android Platform
#ifdef ANDROID
SANITY_CHECK(dst, 255); // TODO: Reimplement check in future versions
SANITY_CHECK(dst, 5);
#else
SANITY_CHECK(dst, 1 + 1e-6);
#endif

59
modules/imgproc/perf/perf_warp.cpp
View File

@ -28,24 +28,23 @@ PERF_TEST_P( TestWarpAffine, WarpAffine,
)
)
{
Size sz;
Size sz, szSrc(512, 512);
int borderMode, interType;
sz = get<0>(GetParam());
interType = get<1>(GetParam());
borderMode = get<2>(GetParam());
Scalar borderColor = Scalar::all(150);
Mat src, img = imread(getDataPath("cv/shared/fruits.png"));
cvtColor(img, src, COLOR_BGR2RGBA, 4);
Mat src(szSrc,CV_8UC4), dst(sz, CV_8UC4);
cvtest::fillGradient(src);
if(borderMode == BORDER_CONSTANT) cvtest::smoothBorder(src, borderColor, 1);
Mat warpMat = getRotationMatrix2D(Point2f(src.cols/2.f, src.rows/2.f), 30., 2.2);
Mat dst(sz, CV_8UC4);
declare.in(src).out(dst);
TEST_CYCLE() warpAffine( src, dst, warpMat, sz, interType, borderMode, Scalar::all(150) );
TEST_CYCLE() warpAffine( src, dst, warpMat, sz, interType, borderMode, borderColor );
// Test case temporary disabled for Android Platform
#ifdef ANDROID
SANITY_CHECK(dst, 255); // TODO: Reimplement check in future versions
SANITY_CHECK(dst, interType==INTER_LINEAR? 5 : 10);
#else
SANITY_CHECK(dst, 1);
#endif
@ -59,15 +58,16 @@ PERF_TEST_P( TestWarpPerspective, WarpPerspective,
)
)
{
Size sz;
Size sz, szSrc(512, 512);
int borderMode, interType;
sz = get<0>(GetParam());
interType = get<1>(GetParam());
borderMode = get<2>(GetParam());
Scalar borderColor = Scalar::all(150);
Mat src, img = imread(getDataPath("cv/shared/fruits.png"));
cvtColor(img, src, COLOR_BGR2RGBA, 4);
Mat src(szSrc,CV_8UC4), dst(sz, CV_8UC4);
cvtest::fillGradient(src);
if(borderMode == BORDER_CONSTANT) cvtest::smoothBorder(src, borderColor, 1);
Mat rotMat = getRotationMatrix2D(Point2f(src.cols/2.f, src.rows/2.f), 30., 2.2);
Mat warpMat(3, 3, CV_64FC1);
for(int r=0; r<2; r++)
@ -76,13 +76,16 @@ PERF_TEST_P( TestWarpPerspective, WarpPerspective,
warpMat.at<double>(2, 0) = .3/sz.width;
warpMat.at<double>(2, 1) = .3/sz.height;
warpMat.at<double>(2, 2) = 1;
Mat dst(sz, CV_8UC4);
declare.in(src).out(dst);
TEST_CYCLE() warpPerspective( src, dst, warpMat, sz, interType, borderMode, Scalar::all(150) );
TEST_CYCLE() warpPerspective( src, dst, warpMat, sz, interType, borderMode, borderColor );
#ifdef ANDROID
SANITY_CHECK(dst, interType==INTER_LINEAR? 5 : 10);
#else
SANITY_CHECK(dst, 1);
#endif
}
PERF_TEST_P( TestWarpPerspectiveNear_t, WarpPerspectiveNear,
@ -105,24 +108,11 @@ PERF_TEST_P( TestWarpPerspectiveNear_t, WarpPerspectiveNear,
interType = get<1>(GetParam());
borderMode = get<2>(GetParam());
type = get<3>(GetParam());
Scalar borderColor = Scalar::all(150);
Mat src, img = imread(getDataPath("cv/shared/5MP.png"));
if( type == CV_8UC1 )
{
cvtColor(img, src, COLOR_BGR2GRAY, 1);
}
else if( type == CV_8UC4 )
{
cvtColor(img, src, COLOR_BGR2BGRA, 4);
}
else
{
FAIL();
}
resize(src, src, size);
Mat src(size, type), dst(size, type);
cvtest::fillGradient(src);
if(borderMode == BORDER_CONSTANT) cvtest::smoothBorder(src, borderColor, 1);
int shift = static_cast<int>(src.cols*0.04);
Mat srcVertices = (Mat_<Vec2f>(1, 4) << Vec2f(0, 0),
Vec2f(static_cast<float>(size.width-1), 0),
@ -134,19 +124,16 @@ PERF_TEST_P( TestWarpPerspectiveNear_t, WarpPerspectiveNear,
Vec2f(static_cast<float>(shift/2), static_cast<float>(size.height-1)));
Mat warpMat = getPerspectiveTransform(srcVertices, dstVertices);
Mat dst(size, type);
declare.in(src).out(dst);
declare.time(100);
TEST_CYCLE()
{
warpPerspective( src, dst, warpMat, size, interType, borderMode, Scalar::all(150) );
warpPerspective( src, dst, warpMat, size, interType, borderMode, borderColor );
}
// Test case temporary disabled for Android Platform
#ifdef ANDROID
SANITY_CHECK(dst, 255); // TODO: Reimplement check in future versions
SANITY_CHECK(dst, interType==INTER_LINEAR? 5 : 10);
#else
SANITY_CHECK(dst, 1);
#endif

451
modules/java/CMakeLists.txt
View File

@ -1,15 +1,10 @@
# ----------------------------------------------------------------------------
# CMake file for java support
# ----------------------------------------------------------------------------
if(IOS OR NOT PYTHON_EXECUTABLE OR NOT (JNI_FOUND OR (ANDROID AND ANDROID_NATIVE_API_LEVEL GREATER 7)))
if(IOS OR NOT PYTHON_EXECUTABLE OR NOT ANT_EXECUTABLE OR NOT (JNI_FOUND OR (ANDROID AND ANDROID_NATIVE_API_LEVEL GREATER 7)))
ocv_module_disable(java)
endif()
if(NOT ANDROID)
# disable java by default because java support on desktop is experimental
set(BUILD_opencv_java_INIT OFF)
endif()
set(the_description "The java bindings")
ocv_add_module(java BINDINGS opencv_core opencv_imgproc OPTIONAL opencv_objdetect opencv_features2d opencv_video opencv_highgui opencv_ml opencv_calib3d opencv_photo opencv_nonfree opencv_contrib)
ocv_module_include_directories("${CMAKE_CURRENT_SOURCE_DIR}/generator/src/cpp")
@ -18,6 +13,10 @@ if(NOT ANDROID)
include_directories(${JNI_INCLUDE_DIRS})
endif()
# output locations
set(JAVA_INSTALL_ROOT "sdk/java")
set(JNI_INSTALL_ROOT "sdk/native")
# get list of modules to wrap
string(REPLACE "opencv_" "" OPENCV_JAVA_MODULES "${OPENCV_MODULE_${the_module}_REQ_DEPS};${OPENCV_MODULE_${the_module}_OPT_DEPS}")
foreach(module ${OPENCV_JAVA_MODULES})
@ -26,20 +25,29 @@ foreach(module ${OPENCV_JAVA_MODULES})
endif()
endforeach()
set(GEN_JAVA "${CMAKE_CURRENT_SOURCE_DIR}/generator/gen_java.py")
set(HDR_PARSER "${CMAKE_CURRENT_SOURCE_DIR}/../python/src2/hdr_parser.py")
set(GEN_JAVADOC "${CMAKE_CURRENT_SOURCE_DIR}/generator/gen_javadoc.py")
set(RST_PARSER "${CMAKE_CURRENT_SOURCE_DIR}/generator/rst_parser.py")
######################################################################################################################################
# add dependencies to cmake (we should rerun cmake if any of these scripts is modified)
configure_file("${GEN_JAVA}" "${OpenCV_BINARY_DIR}/junk/gen_java.junk" COPYONLY)
configure_file("${HDR_PARSER}" "${OpenCV_BINARY_DIR}/junk/hdr_parser.junk" COPYONLY)
# scripts
set(scripts_gen_java "${CMAKE_CURRENT_SOURCE_DIR}/generator/gen_java.py")
set(scripts_hdr_parser "${CMAKE_CURRENT_SOURCE_DIR}/../python/src2/hdr_parser.py")
set(scripts_gen_javadoc "${CMAKE_CURRENT_SOURCE_DIR}/generator/gen_javadoc.py")
set(scripts_rst_parser "${CMAKE_CURRENT_SOURCE_DIR}/generator/rst_parser.py")
set(java_hdr_deps "")
set(generated_cpp_sources "")
set(generated_java_sources "")
# handwritten C/C++ and Java sources
file(GLOB handwrittren_h_sources "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/cpp/*.h" "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/cpp/*.hpp")
file(GLOB handwrittren_cpp_sources "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/cpp/*.cpp")
file(GLOB handwrittren_java_sources "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/java/*.java")
file(GLOB handwrittren_aidl_sources "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/java/*.aidl")
if(NOT ANDROID)
ocv_list_filterout(handwrittren_java_sources "/(engine|android)\\\\+")
ocv_list_filterout(handwrittren_aidl_sources "/(engine|android)\\\\+")
else()
file(GLOB_RECURSE handwrittren_lib_project_files_rel RELATIVE "${CMAKE_CURRENT_SOURCE_DIR}/android_lib/" "${CMAKE_CURRENT_SOURCE_DIR}/android_lib/*")
list(REMOVE_ITEM handwrittren_lib_project_files_rel "${ANDROID_MANIFEST_FILE}")
endif()
# setup raw java and cpp files generation (without javadoc and at temporary location)
# headers of OpenCV modules
set(opencv_public_headers "")
foreach(module ${OPENCV_JAVA_MODULES})
# get list of module headers
if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/generator/config/${module}.filelist")
@ -48,121 +56,213 @@ foreach(module ${OPENCV_JAVA_MODULES})
else()
set(module_headers "${OPENCV_MODULE_opencv_${module}_HEADERS}")
endif()
# C headers must go first
set(module_headers_cpp ${module_headers})
ocv_list_filterout(module_headers_cpp "\\\\.h$")
if(module_headers_cpp)
list(REMOVE_ITEM module_headers ${module_headers_cpp})
list(APPEND module_headers ${module_headers_cpp})
endif()
unset(module_headers_cpp)
# add dependencies to cmake (we should rerun cmake if any of these headers is modified)
foreach(header ${module_headers})
get_filename_component(header_name "${header}" NAME_WE)
configure_file("${header}" "${OpenCV_BINARY_DIR}/junk/${header_name}.junk" COPYONLY)
endforeach()
# first run (to get list of generated files)
if(module_headers)
file(REMOVE_RECURSE "${CMAKE_CURRENT_BINARY_DIR}/gen_java_out/")
file(MAKE_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/gen_java_out")
execute_process(COMMAND ${PYTHON_EXECUTABLE} "${GEN_JAVA}" "${HDR_PARSER}" ${module} ${module_headers}
WORKING_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/gen_java_out"
OUTPUT_QUIET ERROR_QUIET)
file(GLOB_RECURSE ${module}_generated_java_sources RELATIVE "${CMAKE_CURRENT_BINARY_DIR}/gen_java_out/" "${CMAKE_CURRENT_BINARY_DIR}/gen_java_out/*.java")
ocv_list_add_prefix(${module}_generated_java_sources "${CMAKE_CURRENT_BINARY_DIR}/")
# C headers must go first
set(module_headers_cpp ${module_headers})
ocv_list_filterout(module_headers_cpp "\\\\.h$")
if(module_headers_cpp)
list(REMOVE_ITEM module_headers ${module_headers_cpp})
list(APPEND module_headers ${module_headers_cpp})
endif()
unset(module_headers_cpp)
# second run (at build time)
add_custom_command(OUTPUT ${${module}_generated_java_sources} "${CMAKE_CURRENT_BINARY_DIR}/${module}.cpp"
COMMAND ${PYTHON_EXECUTABLE} "${GEN_JAVA}" "${HDR_PARSER}" ${module} ${module_headers}
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
DEPENDS "${GEN_JAVA}" "${HDR_PARSER}" ${module_headers})
list(APPEND java_hdr_deps ${module_headers})
list(APPEND generated_cpp_sources "${CMAKE_CURRENT_BINARY_DIR}/${module}.cpp")
list(APPEND generated_java_sources ${${module}_generated_java_sources})
set(opencv_public_headers_${module} ${module_headers})
list(APPEND opencv_public_headers ${module_headers})
else()
list(REMOVE_ITEM OPENCV_JAVA_MODULES ${module})
endif()
endforeach()
# get handwritten files used for wrappers generation
file(GLOB handwrittren_h_sources "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/cpp/*.h" "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/cpp/*.hpp")
file(GLOB handwrittren_cpp_sources "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/cpp/*.cpp")
file(GLOB handwrittren_java_sources "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/java/*.java")
file(GLOB handwrittren_aidl_sources "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/java/*.aidl")
# remove handwritten java files for disabled modules
foreach(jfile ${handwrittren_java_sources})
string(REGEX REPLACE "^.*/([^+]+)\\+.*\\.java$" "\\1" jmodname "${jfile}")
if(DEFINED HAVE_opencv_${jmodname} AND NOT HAVE_opencv_${jmodname})
list(REMOVE_ITEM handwrittren_java_sources "${jfile}")
endif()
# rst documentation used for javadoc generation
set(javadoc_rst_sources "")
foreach(module ${OPENCV_JAVA_MODULES})
file(GLOB_RECURSE refman_rst_headers "${OPENCV_MODULE_opencv_${module}_LOCATION}/*.rst")
list(APPEND javadoc_rst_sources ${refman_rst_headers})
endforeach()
# remove VideoCapture wrapper if highgui is disabled
if(NOT HAVE_opencv_highgui)
list(REMOVE_ITEM handwrittren_cpp_sources "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/cpp/VideoCapture.cpp")
endif()
# generated cpp files
set(generated_cpp_sources "")
foreach(module ${OPENCV_JAVA_MODULES})
list(APPEND generated_cpp_sources "${CMAKE_CURRENT_BINARY_DIR}/${module}.cpp")
endforeach()
# create list of javadoc documented files
unset(documented_java_files)
foreach(java_file ${handwrittren_java_sources} ${generated_java_sources})
# IMPORTANT: add dependencies to cmake (we should rerun cmake if any of these files is modified)
configure_file("${scripts_gen_java}" "${OpenCV_BINARY_DIR}/junk/gen_java.junk" COPYONLY)
configure_file("${scripts_hdr_parser}" "${OpenCV_BINARY_DIR}/junk/hdr_parser.junk" COPYONLY)
foreach(header ${opencv_public_headers})
get_filename_component(header_name "${header}" NAME)
configure_file("${header}" "${OpenCV_BINARY_DIR}/junk/${header_name}.junk" COPYONLY)
endforeach()
# generated java files
set(generated_java_sources "")
foreach(module ${OPENCV_JAVA_MODULES})
# first run of gen_java.py (to get list of generated files)
file(REMOVE_RECURSE "${CMAKE_CURRENT_BINARY_DIR}/gen_java_out/")
file(MAKE_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/gen_java_out")
execute_process(COMMAND ${PYTHON_EXECUTABLE} "${scripts_gen_java}" "${scripts_hdr_parser}" ${module} ${opencv_public_headers_${module}}
WORKING_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/gen_java_out"
OUTPUT_QUIET ERROR_QUIET)
unset(generated_java_sources_${module})
file(GLOB_RECURSE generated_java_sources_${module} RELATIVE "${CMAKE_CURRENT_BINARY_DIR}/gen_java_out/" "${CMAKE_CURRENT_BINARY_DIR}/gen_java_out/*.java")
ocv_list_add_prefix(generated_java_sources_${module} "${CMAKE_CURRENT_BINARY_DIR}/")
list(APPEND generated_java_sources ${generated_java_sources_${module}})
endforeach()
# generated java files with javadoc
set(documented_java_files "")
foreach(java_file ${generated_java_sources} ${handwrittren_java_sources})
get_filename_component(java_file_name "${java_file}" NAME_WE)
list(APPEND documented_java_files "${CMAKE_CURRENT_BINARY_DIR}/${java_file_name}-jdoc.java")
endforeach()
# generate javadoc files
file(GLOB_RECURSE refman_rst_headers "${CMAKE_CURRENT_SOURCE_DIR}/../*.rst")
set(java_documented_headers_deps ${handwrittren_java_sources} ${generated_java_sources} ${java_hdr_deps} ${refman_rst_headers}
"${GEN_JAVADOC}" "${RST_PARSER}" "${GEN_JAVA}" "${HDR_PARSER}")
######################################################################################################################################
# step 1: generate .cpp/.java from OpenCV headers
set(step1_depends "${scripts_gen_java}" "${scripts_hdr_parser}" ${opencv_public_headers})
foreach(module ${OPENCV_JAVA_MODULES})
# second run of gen_java.py (at build time)
add_custom_command(OUTPUT ${generated_java_sources_${module}} "${CMAKE_CURRENT_BINARY_DIR}/${module}.cpp"
COMMAND ${PYTHON_EXECUTABLE} "${scripts_gen_java}" "${scripts_hdr_parser}" ${module} ${opencv_public_headers_${module}}
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
DEPENDS "${scripts_gen_java}" "${scripts_hdr_parser}" ${opencv_public_headers_${module}}
)
endforeach()
# step 2: generate javadoc comments
set(step2_depends ${step1_depends} ${scripts_gen_javadoc} ${scripts_rst_parser} ${javadoc_rst_sources} ${generated_java_sources} ${handwrittren_java_sources})
string(REPLACE ";" "," OPENCV_JAVA_MODULES_STR "${OPENCV_JAVA_MODULES}")
add_custom_command(
OUTPUT ${documented_java_files}
COMMAND ${PYTHON_EXECUTABLE} "${GEN_JAVADOC}" --modules ${OPENCV_JAVA_MODULES_STR} "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/java" "${CMAKE_CURRENT_BINARY_DIR}" 2>"${CMAKE_CURRENT_BINARY_DIR}/get_javadoc_errors.log"
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
DEPENDS ${java_documented_headers_deps}
)
add_custom_command(OUTPUT ${documented_java_files}
COMMAND ${PYTHON_EXECUTABLE} "${scripts_gen_javadoc}" --modules ${OPENCV_JAVA_MODULES_STR} "${CMAKE_CURRENT_SOURCE_DIR}/generator/src/java" "${CMAKE_CURRENT_BINARY_DIR}" 2>"${CMAKE_CURRENT_BINARY_DIR}/get_javadoc_errors.log"
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
DEPENDS ${step2_depends}
)
# copy generated java files to the final location
set(JAVA_OUTPUT_DIR "src/org/opencv")
set(JAVA_INSTALL_ROOT "sdk/java")
set(JNI_INSTALL_ROOT "sdk/native")
# copy each documented header to the final destination
set(java_files "")
set(source_java_files ${documented_java_files} ${handwrittren_aidl_sources})
if(NOT ANDROID)
ocv_list_filterout(source_java_files "/(engine|android)\\\\+")
endif()
foreach(java_file ${source_java_files})
# step 3: copy files to destination
set(step3_input_files ${documented_java_files} ${handwrittren_aidl_sources})
set(copied_files "")
foreach(java_file ${step3_input_files})
get_filename_component(java_file_name "${java_file}" NAME)
string(REPLACE "-jdoc.java" ".java" java_file_name "${java_file_name}")
string(REPLACE "+" "/" java_file_name "${java_file_name}")
set(output_name "${OpenCV_BINARY_DIR}/src/org/opencv/${java_file_name}")
add_custom_command(
OUTPUT "${OpenCV_BINARY_DIR}/${JAVA_OUTPUT_DIR}/${java_file_name}"
COMMAND ${CMAKE_COMMAND} -E copy "${java_file}" "${OpenCV_BINARY_DIR}/${JAVA_OUTPUT_DIR}/${java_file_name}"
MAIN_DEPENDENCY "${java_file}"
DEPENDS ${java_documented_headers_deps}
COMMENT "Generating ${JAVA_OUTPUT_DIR}/${java_file_name}"
)
list(APPEND java_files "${OpenCV_BINARY_DIR}/${JAVA_OUTPUT_DIR}/${java_file_name}")
add_custom_command(OUTPUT "${output_name}"
COMMAND ${CMAKE_COMMAND} -E copy "${java_file}" "${output_name}"
MAIN_DEPENDENCY "${java_file}"
DEPENDS ${step2_depends}
COMMENT "Generating src/org/opencv/${java_file_name}"
)
list(APPEND copied_files "${output_name}")
if(ANDROID)
get_filename_component(install_subdir "${java_file_name}" PATH)
install(FILES "${OpenCV_BINARY_DIR}/${JAVA_OUTPUT_DIR}/${java_file_name}" DESTINATION ${JAVA_INSTALL_ROOT}/${JAVA_OUTPUT_DIR}/${install_subdir} COMPONENT main)
install(FILES "${output_name}" DESTINATION "${JAVA_INSTALL_ROOT}/src/org/opencv/${install_subdir}" COMPONENT main)
endif()
endforeach()
# custom target for java API
set(api_target ${the_module}_api)
add_custom_target(${api_target} DEPENDS ${java_files} ${documented_java_files} ${java_documented_headers_deps})
if(ANDROID)
set(android_copied_files "")
set(android_step3_input_files "")
foreach(file ${handwrittren_lib_project_files_rel})
add_custom_command(OUTPUT "${OpenCV_BINARY_DIR}/${file}"
COMMAND ${CMAKE_COMMAND} -E copy "${CMAKE_CURRENT_SOURCE_DIR}/android_lib/${file}" "${OpenCV_BINARY_DIR}/${file}"
MAIN_DEPENDENCY "${CMAKE_CURRENT_SOURCE_DIR}/android_lib/${file}"
COMMENT "Generating ${file}"
)
list(APPEND android_copied_files "${OpenCV_BINARY_DIR}/${file}")
list(APPEND android_step3_input_files "${CMAKE_CURRENT_SOURCE_DIR}/android_lib/${file}")
# add opencv_java library
add_library(${the_module} SHARED ${handwrittren_h_sources} ${handwrittren_cpp_sources} ${generated_cpp_sources})
if(NOT file MATCHES "jni/.+")
get_filename_component(install_subdir "${file}" PATH)
install(FILES "${OpenCV_BINARY_DIR}/${file}" DESTINATION "${JAVA_INSTALL_ROOT}/${install_subdir}" COMPONENT main)
endif()
endforeach()
# library project jni sources (nothing really depends on them so we will not add them to step3_input_files)
foreach(jni_file ${handwrittren_cpp_sources} ${handwrittren_h_sources} ${generated_cpp_sources})
get_filename_component(jni_file_name "${jni_file}" NAME)
add_custom_command(OUTPUT "${OpenCV_BINARY_DIR}/jni/${jni_file_name}"
COMMAND ${CMAKE_COMMAND} -E copy "${jni_file}" "${OpenCV_BINARY_DIR}/jni/${jni_file_name}"
DEPENDS "${jni_file}" ${java_hdr_deps}
COMMENT "Generating jni/${jni_file_name}"
)
list(APPEND android_copied_files "${OpenCV_BINARY_DIR}/jni/${jni_file_name}")
endforeach()
endif(ANDROID)
# step 3.5: generate Android library project
if(ANDROID AND ANDROID_EXECUTABLE)
set(lib_target_files ${ANDROID_LIB_PROJECT_FILES})
ocv_list_add_prefix(lib_target_files "${OpenCV_BINARY_DIR}/")
android_get_compatible_target(lib_target_sdk_target ${ANDROID_NATIVE_API_LEVEL} ${ANDROID_SDK_TARGET} 11)
configure_file("${CMAKE_CURRENT_SOURCE_DIR}/android_lib/${ANDROID_MANIFEST_FILE}" "${CMAKE_CURRENT_BINARY_DIR}/${ANDROID_MANIFEST_FILE}")
add_custom_command(OUTPUT ${lib_target_files} "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}"
COMMAND ${CMAKE_COMMAND} -E remove ${lib_target_files}
COMMAND ${CMAKE_COMMAND} -E copy "${CMAKE_CURRENT_BINARY_DIR}/${ANDROID_MANIFEST_FILE}" "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}"
COMMAND ${ANDROID_EXECUTABLE} --silent create lib-project --path \"${OpenCV_BINARY_DIR}\" --target \"${lib_target_sdk_target}\" --name OpenCV --package org.opencv 2>\"${CMAKE_CURRENT_BINARY_DIR}/create_lib_project.log\"
COMMAND ${CMAKE_COMMAND} -E copy "${CMAKE_CURRENT_BINARY_DIR}/${ANDROID_MANIFEST_FILE}" "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}"
MAIN_DEPENDENCY "${CMAKE_CURRENT_BINARY_DIR}/${ANDROID_MANIFEST_FILE}"
DEPENDS ${android_step3_input_files} ${android_copied_files}
COMMENT "Generating OpenCV Android library project. SDK target: ${lib_target_sdk_target}"
)
list(APPEND copied_files ${lib_target_files} "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}")
list(APPEND step3_input_files "${CMAKE_CURRENT_BINARY_DIR}/${ANDROID_MANIFEST_FILE}")
install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_PROJECT_PROPERTIES_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT main)
install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT main)
# creating empty 'gen' and 'res' folders
install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/${JAVA_INSTALL_ROOT}/gen\")" COMPONENT main)
install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/${JAVA_INSTALL_ROOT}/res\")" COMPONENT main)
endif(ANDROID AND ANDROID_EXECUTABLE)
set(step3_depends ${step2_depends} ${step3_input_files} ${copied_files})
# step 4: build jar
if(ANDROID)
set(JAR_FILE "${OpenCV_BINARY_DIR}/bin/classes.jar")
if(ANDROID_TOOLS_Pkg_Revision GREATER 13)
# build the library project
# normally we should do this after a native part, but for a library project we can build the java part first
add_custom_command(OUTPUT "${JAR_FILE}" "${JAR_FILE}.dephelper"
COMMAND ${ANT_EXECUTABLE} -q -noinput -k debug
COMMAND ${CMAKE_COMMAND} -E touch "${JAR_FILE}.dephelper" # can not rely on classes.jar because different versions of SDK update timestamp at different times
WORKING_DIRECTORY "${OpenCV_BINARY_DIR}"
DEPENDS ${step3_depends}
COMMENT "Building OpenCV Android library project"
)
else()
# ditto
add_custom_command(OUTPUT "${JAR_FILE}" "${JAR_FILE}.dephelper"
COMMAND ${CMAKE_COMMAND} -E touch "${JAR_FILE}"
COMMAND ${CMAKE_COMMAND} -E touch "${JAR_FILE}.dephelper"
WORKING_DIRECTORY "${OpenCV_BINARY_DIR}"
DEPENDS ${step3_depends}
COMMENT ""
)
endif()
else(ANDROID)
set(JAR_NAME opencv-${OPENCV_VERSION}.jar)
set(JAR_FILE "${OpenCV_BINARY_DIR}/bin/${JAR_NAME}")
configure_file("${CMAKE_CURRENT_SOURCE_DIR}/build.xml.in" "${OpenCV_BINARY_DIR}/build.xml" IMMEDIATE @ONLY)
list(APPEND step3_depends "${OpenCV_BINARY_DIR}/build.xml")
add_custom_command(OUTPUT "${JAR_FILE}" "${JAR_FILE}.dephelper"
COMMAND ${ANT_EXECUTABLE} -q -noinput -k jar
COMMAND ${CMAKE_COMMAND} -E touch "${JAR_FILE}.dephelper"
WORKING_DIRECTORY "${OpenCV_BINARY_DIR}"
DEPENDS ${step3_depends}
COMMENT "Generating ${JAR_NAME}"
)
endif(ANDROID)
# step 5: build native part
add_library(${the_module} SHARED ${handwrittren_h_sources} ${handwrittren_cpp_sources} ${generated_cpp_sources}
${copied_files}
"${JAR_FILE}" "${JAR_FILE}.dephelper")
if(BUILD_FAT_JAVA_LIB)
set(__deps ${OPENCV_MODULE_${the_module}_DEPS} ${OPENCV_MODULES_BUILD})
list(REMOVE_ITEM __deps ${the_module} opencv_ts)
@ -176,22 +276,7 @@ if(BUILD_FAT_JAVA_LIB)
else()
target_link_libraries(${the_module} ${OPENCV_MODULE_${the_module}_DEPS} ${OPENCV_LINKER_LIBS})
endif()
add_dependencies(${the_module} ${api_target})
# Additional target properties
set_target_properties(${the_module} PROPERTIES
OUTPUT_NAME "${the_module}"
ARCHIVE_OUTPUT_DIRECTORY ${LIBRARY_OUTPUT_PATH}
RUNTIME_OUTPUT_DIRECTORY ${EXECUTABLE_OUTPUT_PATH}
INSTALL_NAME_DIR ${OPENCV_LIB_INSTALL_PATH}
LINK_INTERFACE_LIBRARIES ""
)
install(TARGETS ${the_module}
LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main
ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main)
set(lib_target ${the_module}_library)
if(ANDROID)
target_link_libraries(${the_module} jnigraphics) # for Mat <=> Bitmap converters
@ -202,105 +287,37 @@ if(ANDROID)
if ( NOT (CMAKE_BUILD_TYPE MATCHES "Debug"))
add_custom_command(TARGET ${the_module} POST_BUILD COMMAND ${CMAKE_STRIP} --strip-unneeded "${__opencv_java_location}")
endif()
endif()
set(lib_proj_files "")
# manifest, jni, Eclipse project
file(GLOB_RECURSE android_lib_project_files RELATIVE "${CMAKE_CURRENT_SOURCE_DIR}/android_lib/" "${CMAKE_CURRENT_SOURCE_DIR}/android_lib/*")
list(REMOVE_ITEM android_lib_project_files "${ANDROID_MANIFEST_FILE}")
foreach(f ${android_lib_project_files})
if(NOT f MATCHES "\\.svn")
configure_file("${CMAKE_CURRENT_SOURCE_DIR}/android_lib/${f}" "${OpenCV_BINARY_DIR}/${f}")
list(APPEND lib_proj_files "${OpenCV_BINARY_DIR}/${f}")
if(NOT f MATCHES "jni/.+")
get_filename_component(install_subdir "${f}" PATH)
install(FILES "${OpenCV_BINARY_DIR}/${f}" DESTINATION "${JAVA_INSTALL_ROOT}/${install_subdir}" COMPONENT main)
endif()
endif()
endforeach()
# library project jni sources
foreach(jni_file ${handwrittren_cpp_sources} ${handwrittren_h_sources} ${generated_cpp_sources})
get_filename_component(jni_file_name "${jni_file}" NAME)
add_custom_command(
OUTPUT "${OpenCV_BINARY_DIR}/jni/${jni_file_name}"
COMMAND ${CMAKE_COMMAND} -E copy "${jni_file}" "${OpenCV_BINARY_DIR}/jni/${jni_file_name}"
DEPENDS "${jni_file}" ${java_hdr_deps}
COMMENT "Generating jni/${jni_file_name}"
)
list(APPEND lib_proj_files "${OpenCV_BINARY_DIR}/jni/${jni_file_name}")
endforeach()
# create Android library project in build folder
if(ANDROID_EXECUTABLE)
set(lib_target_files ${ANDROID_LIB_PROJECT_FILES})
ocv_list_add_prefix(lib_target_files "${OpenCV_BINARY_DIR}/")
android_get_compatible_target(lib_target_sdk_target ${ANDROID_NATIVE_API_LEVEL} ${ANDROID_SDK_TARGET} 11)
configure_file("${CMAKE_CURRENT_SOURCE_DIR}/android_lib/${ANDROID_MANIFEST_FILE}" "${CMAKE_CURRENT_BINARY_DIR}/${ANDROID_MANIFEST_FILE}")
add_custom_command(
OUTPUT ${lib_target_files} "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}"
COMMAND ${CMAKE_COMMAND} -E remove ${lib_target_files}
COMMAND ${CMAKE_COMMAND} -E copy "${CMAKE_CURRENT_BINARY_DIR}/${ANDROID_MANIFEST_FILE}" "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}"
COMMAND ${ANDROID_EXECUTABLE} --silent create lib-project --path \"${OpenCV_BINARY_DIR}\" --target \"${lib_target_sdk_target}\" --name OpenCV --package org.opencv 2>\"${CMAKE_CURRENT_BINARY_DIR}/create_lib_project.log\"
COMMAND ${CMAKE_COMMAND} -E copy "${CMAKE_CURRENT_BINARY_DIR}/${ANDROID_MANIFEST_FILE}" "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}"
MAIN_DEPENDENCY "${CMAKE_CURRENT_BINARY_DIR}/${ANDROID_MANIFEST_FILE}"
DEPENDS ${lib_proj_files}
COMMENT "Generating OpenCV Android library project. SDK target: ${lib_target_sdk_target}"
)
install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_PROJECT_PROPERTIES_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT main)
install(FILES "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}" DESTINATION ${JAVA_INSTALL_ROOT} COMPONENT main)
# creating empty 'gen' and 'res' folders
install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/sdk/java/gen\")" COMPONENT main)
install(CODE "MAKE_DIRECTORY(\"\$ENV{DESTDIR}\${CMAKE_INSTALL_PREFIX}/sdk/java/res\")" COMPONENT main)
if(ANT_EXECUTABLE AND ANDROID_TOOLS_Pkg_Revision GREATER 13)
# build the library project
# normally we should do this after a native part, but for a library project we can build the java part first
add_custom_command(
OUTPUT "${OpenCV_BINARY_DIR}/bin/classes.jar" "${OpenCV_BINARY_DIR}/bin/.classes.jar.dephelper"
COMMAND ${ANT_EXECUTABLE} -q -noinput -k debug
COMMAND ${CMAKE_COMMAND} -E touch "${OpenCV_BINARY_DIR}/bin/.classes.jar.dephelper" # can not rely on classes.jar because different versions of SDK update timestamp at different times
WORKING_DIRECTORY "${OpenCV_BINARY_DIR}"
DEPENDS ${lib_proj_files} ${lib_target_files} ${java_files}
COMMENT "Building OpenCV Android library project"
)
#install(FILES "${OpenCV_BINARY_DIR}/bin/classes.jar" "${OpenCV_BINARY_DIR}/bin/jarlist.cache" "${OpenCV_BINARY_DIR}/bin/build.prop" DESTINATION bin COMPONENT main)
#install(DIRECTORY "${OpenCV_BINARY_DIR}/bin/res" "${OpenCV_BINARY_DIR}/bin/classes" DESTINATION bin COMPONENT main)
list(APPEND lib_target_files "${OpenCV_BINARY_DIR}/bin/classes.jar")
endif()
add_custom_target(${lib_target} SOURCES ${lib_proj_files} ${lib_target_files} "${OpenCV_BINARY_DIR}/${ANDROID_MANIFEST_FILE}")
endif()
else(ANDROID)
configure_file("${CMAKE_CURRENT_SOURCE_DIR}/jar/build.xml" "${OpenCV_BINARY_DIR}/build.xml" IMMEDIATE @ONLY)
set(JAR_NAME opencv-${OPENCV_VERSION_MAJOR}.${OPENCV_VERSION_MINOR}.${OPENCV_VERSION_PATCH}.jar)
add_custom_command(
OUTPUT "${OpenCV_BINARY_DIR}/bin/${JAR_NAME}" "${OpenCV_BINARY_DIR}/bin/.${JAR_NAME}.dephelper"
COMMAND ${ANT_EXECUTABLE} -q -noinput -k jar
COMMAND ${CMAKE_COMMAND} -E touch "${OpenCV_BINARY_DIR}/bin/.${JAR_NAME}.dephelper"
WORKING_DIRECTORY "${OpenCV_BINARY_DIR}"
DEPENDS "${OpenCV_BINARY_DIR}/build.xml" ${java_files}
COMMENT "Generating ${JAR_NAME}"
# Additional target properties
set_target_properties(${the_module} PROPERTIES
OUTPUT_NAME "${the_module}"
ARCHIVE_OUTPUT_DIRECTORY ${LIBRARY_OUTPUT_PATH}
RUNTIME_OUTPUT_DIRECTORY ${EXECUTABLE_OUTPUT_PATH}
INSTALL_NAME_DIR ${OPENCV_LIB_INSTALL_PATH}
LINK_INTERFACE_LIBRARIES ""
)
add_custom_target(${lib_target} SOURCES "${OpenCV_BINARY_DIR}/bin/${JAR_NAME}")
endif(ANDROID)
add_dependencies(${lib_target} ${api_target})
add_dependencies(${the_module} ${lib_target})
# android test project
if(ANDROID AND BUILD_TESTS)
add_subdirectory(android_test)
if(ANDROID)
set_target_properties(${the_module} PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${LIBRARY_OUTPUT_PATH})
else()
set_target_properties(${the_module} PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${EXECUTABLE_OUTPUT_PATH})
endif()
# Desktop Java test project.
if((NOT ANDROID) AND BUILD_TESTS)
add_subdirectory(java_test)
if(ENABLE_SOLUTION_FOLDERS)
set_target_properties(${the_module} PROPERTIES FOLDER "bindings")
endif()
install(TARGETS ${the_module}
LIBRARY DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main
ARCHIVE DESTINATION ${OPENCV_LIB_INSTALL_PATH} COMPONENT main)
######################################################################################################################################
if(BUILD_TESTS)
if(ANDROID)
add_subdirectory(android_test)
else()
add_subdirectory(test)
endif()
endif()

2
modules/java/android_lib/.project
View File

@ -1,6 +1,6 @@
<?xml version="1.0" encoding="UTF-8"?>
<projectDescription>
<name>OpenCV Library - @OPENCV_VERSION_MAJOR@.@OPENCV_VERSION_MINOR@.@OPENCV_VERSION_PATCH@</name>
<name>OpenCV Library - @OPENCV_VERSION@</name>
<comment></comment>
<projects>
</projects>

4
modules/java/android_lib/AndroidManifest.xml
View File

@ -1,8 +1,8 @@
<?xml version="1.0" encoding="utf-8"?>
<manifest xmlns:android="http://schemas.android.com/apk/res/android"
package="org.opencv"
android:versionCode="@OPENCV_VERSION_MAJOR@@OPENCV_VERSION_MINOR@@OPENCV_VERSION_PATCH@"
android:versionName="@OPENCV_VERSION_MAJOR@.@OPENCV_VERSION_MINOR@.@OPENCV_VERSION_PATCH@">
android:versionCode="@OPENCV_VERSION_MAJOR@@OPENCV_VERSION_MINOR@@OPENCV_VERSION_PATCH@@OPENCV_VERSION_TWEAK@"
android:versionName="@OPENCV_VERSION@">
<uses-sdk android:minSdkVersion="8" />
</manifest>

2
modules/java/android_test/CMakeLists.txt
View File

@ -47,7 +47,7 @@ add_custom_command(
COMMAND ${CMAKE_COMMAND} -E touch "${opencv_test_java_bin_dir}/bin/OpenCVTest-debug.apk" # needed because ant does not update the timestamp of updated apk
WORKING_DIRECTORY "${opencv_test_java_bin_dir}"
MAIN_DEPENDENCY "${opencv_test_java_bin_dir}/${ANDROID_MANIFEST_FILE}"
DEPENDS "${OpenCV_BINARY_DIR}/bin/.classes.jar.dephelper" opencv_java
DEPENDS "${JAR_FILE}.dephelper" opencv_java
DEPENDS ${opencv_test_java_file_deps})
add_custom_target(${PROJECT_NAME} ALL SOURCES "${opencv_test_java_bin_dir}/bin/OpenCVTest-debug.apk" )

2
modules/java/jar/build.xml → modules/java/build.xml.in
View File

@ -10,6 +10,6 @@
<include name="**/*.java"/>
</javac>
<jar basedir="src" destfile="bin/opencv-@OPENCV_VERSION_MAJOR@.@OPENCV_VERSION_MINOR@.@OPENCV_VERSION_PATCH@.jar"/>
<jar basedir="src" destfile="bin/@JAR_NAME@"/>
</target>
</project>

17
modules/java/generator/gen_java.py
View File

@ -871,22 +871,9 @@ public class %(jc)s {
// This file is auto-generated, please don't edit!
//
#include <jni.h>
#include "converters.h"
#if defined DEBUG && defined ANDROID
# include <android/log.h>
# define MODULE_LOG_TAG "OpenCV.%(m)s"
# define LOGD(...) ((void)__android_log_print(ANDROID_LOG_DEBUG, MODULE_LOG_TAG, __VA_ARGS__))
#else //DEBUG
# define LOGD(...)
#endif //DEBUG
#ifdef _MSC_VER
# pragma warning(disable:4800 4244)
#endif
#define LOG_TAG "org.opencv.%(m)s"
#include "common.h"
#include "opencv2/%(m)s/%(m)s.hpp"
using namespace cv;

23
modules/java/generator/src/cpp/Mat.cpp
View File

@ -1,27 +1,6 @@
#include <jni.h>
#include "converters.h"
#ifdef ANDROID
#include <android/log.h>
#define LOG_TAG "org.opencv.core.Mat"
#define LOGE(...) ((void)__android_log_print(ANDROID_LOG_ERROR, LOG_TAG, __VA_ARGS__))
#ifdef DEBUG
#define LOGD(...) ((void)__android_log_print(ANDROID_LOG_DEBUG, LOG_TAG, __VA_ARGS__))
#else //!DEBUG
#define LOGD(...)
#endif //DEBUG
#else
#define LOGE(...)
#define LOGD(...)
#endif
#ifdef _MSC_VER
# pragma warning(disable:4800)
#endif
#include "common.h"
#include "opencv2/core/core.hpp"
using namespace cv;

11
modules/java/generator/src/cpp/VideoCapture.cpp
View File

@ -1,12 +1,5 @@
#include <jni.h>
#if defined DEBUG && defined ANDROID
#include <android/log.h>
#define MODULE_LOG_TAG "OpenCV.highgui"
#define LOGD(...) ((void)__android_log_print(ANDROID_LOG_DEBUG, MODULE_LOG_TAG, __VA_ARGS__))
#else
#define LOGD(...)
#endif
#define LOG_TAG "org.opencv.highgui.VideoCapture"
#include "common.h"
#include "opencv2/opencv_modules.hpp"
#ifdef HAVE_OPENCV_HIGHGUI

33
modules/java/generator/src/cpp/common.h Normal file
View File

@ -0,0 +1,33 @@
#ifndef __JAVA_COMMON_H__
#define __JAVA_COMMON_H__
#if !defined(__ppc__)
// to suppress warning from jni.h on OS X
# define TARGET_RT_MAC_CFM 0
#endif
#include <jni.h>
#ifdef __ANDROID__
# include <android/log.h>
# define LOGE(...) ((void)__android_log_print(ANDROID_LOG_ERROR, LOG_TAG, __VA_ARGS__))
# ifdef DEBUG
# define LOGD(...) ((void)__android_log_print(ANDROID_LOG_DEBUG, LOG_TAG, __VA_ARGS__))
# else
# define LOGD(...)
# endif
#else
# define LOGE(...)
# define LOGD(...)
#endif
#include "converters.h"
#include "core_manual.hpp"
#include "features2d_manual.hpp"
#ifdef _MSC_VER
# pragma warning(disable:4800 4244)
#endif
#endif //__JAVA_COMMON_H__

11
modules/java/generator/src/cpp/converters.cpp
View File

@ -1,12 +1,5 @@
#include "converters.h"
#if defined DEBUG && defined ANDROID
#include <android/log.h>
#define MODULE_LOG_TAG "OpenCV.converters"
#define LOGD(...) ((void)__android_log_print(ANDROID_LOG_DEBUG, MODULE_LOG_TAG, __VA_ARGS__))
#else //DEBUG
#define LOGD(...)
#endif //DEBUG
#define LOG_TAG "org.opencv.utils.Converters"
#include "common.h"
using namespace cv;

2
modules/java/generator/src/cpp/converters.h
View File

@ -1,5 +1,3 @@
#include <jni.h>
#include "opencv2/opencv_modules.hpp"
#include "opencv2/core/core.hpp"
#include "features2d_manual.hpp"

15
modules/java/generator/src/cpp/core_manual.cpp Normal file
View File

@ -0,0 +1,15 @@
#define LOG_TAG "org.opencv.core.Core"
#include "common.h"
static int quietCallback( int, const char*, const char*, const char*, int, void* )
{
return 0;
}
void cv::setErrorVerbosity(bool verbose)
{
if(verbose)
cv::redirectError(0);
else
cv::redirectError((cv::ErrorCallback)quietCallback);
}

7
modules/java/generator/src/cpp/core_manual.hpp
View File

@ -2,6 +2,13 @@
#include "opencv2/core/core.hpp"
namespace cv
{
CV_EXPORTS_W void setErrorVerbosity(bool verbose);
}
#if 0
namespace cv

2
modules/java/generator/src/cpp/jni_part.cpp
View File

@ -1,4 +1,4 @@
#include <jni.h>
#include "common.h"
#include "opencv2/opencv_modules.hpp"

18
modules/java/generator/src/cpp/utils.cpp
View File

@ -1,24 +1,14 @@
#include <jni.h>
#define LOG_TAG "org.opencv.android.Utils"
#include "common.h"
#include "opencv2/core/core.hpp"
#include "opencv2/imgproc/imgproc.hpp"
#ifdef ANDROID
#ifdef __ANDROID__
#include <android/bitmap.h>
#include <android/log.h>
#define LOG_TAG "org.opencv.android.Utils"
#define LOGE(...) ((void)__android_log_print(ANDROID_LOG_ERROR, LOG_TAG, __VA_ARGS__))
#ifdef DEBUG
#define LOGD(...) ((void)__android_log_print(ANDROID_LOG_DEBUG, LOG_TAG, __VA_ARGS__))
#else //!DEBUG
#define LOGD(...)
#endif //DEBUG
using namespace cv;
extern "C" {
/*
@ -168,4 +158,4 @@ JNIEXPORT void JNICALL Java_org_opencv_android_Utils_nMatToBitmap
} // extern "C"
#endif //ANDROID
#endif //__ANDROID__

80
modules/java/java_test/CMakeLists.txt
View File

@ -1,80 +0,0 @@
ocv_check_dependencies(opencv_java ${OPENCV_MODULE_opencv_java_OPT_DEPS} ${OPENCV_MODULE_opencv_java_REQ_DEPS})
if(NOT OCV_DEPENDENCIES_FOUND OR NOT ANT_EXECUTABLE)
return()
endif()
# TODO: This has the same name as the Android test project. That project should
# probably be renamed.
project(opencv_test_java)
set(opencv_test_java_bin_dir "${CMAKE_CURRENT_BINARY_DIR}/.build")
set(android_source_dir "${CMAKE_CURRENT_SOURCE_DIR}/../android_test")
set(java_source_dir ${CMAKE_CURRENT_SOURCE_DIR})
# get project sources
file(GLOB_RECURSE opencv_test_java_files RELATIVE "${android_source_dir}" "${android_source_dir}/res/*" "${android_source_dir}/src/*")
ocv_list_filterout(opencv_test_java_files ".svn")
ocv_list_filterout(opencv_test_java_files ".*#.*")
# These are the files that need to be updated for pure Java.
ocv_list_filterout(opencv_test_java_files ".*OpenCVTestCase.*")
ocv_list_filterout(opencv_test_java_files ".*OpenCVTestRunner.*")
# These files aren't for desktop Java.
ocv_list_filterout(opencv_test_java_files ".*android.*")
# These are files updated for pure Java.
file(GLOB_RECURSE modified_files RELATIVE "${java_source_dir}" "${java_source_dir}/src/*")
ocv_list_filterout(modified_files ".svn")
ocv_list_filterout(modified_files ".*#.*")
# These are extra jars needed to run the tests.
file(GLOB_RECURSE lib_files RELATIVE "${java_source_dir}" "${java_source_dir}/lib/*")
ocv_list_filterout(lib_files ".svn")
ocv_list_filterout(lib_files ".*#.*")
# copy sources out from the build tree
set(opencv_test_java_file_deps "")
foreach(f ${opencv_test_java_files})
add_custom_command(
OUTPUT "${opencv_test_java_bin_dir}/${f}"
COMMAND ${CMAKE_COMMAND} -E copy "${android_source_dir}/${f}" "${opencv_test_java_bin_dir}/${f}"
MAIN_DEPENDENCY "${android_source_dir}/${f}"
COMMENT "Copying ${f}")
list(APPEND opencv_test_java_file_deps "${android_source_dir}/${f}" "${opencv_test_java_bin_dir}/${f}")
endforeach()
# Overwrite select Android sources with Java-specific sources.
# Also, copy over the libs we'll need for testing.
foreach(f ${modified_files} ${lib_files})
add_custom_command(
OUTPUT "${opencv_test_java_bin_dir}/${f}"
COMMAND ${CMAKE_COMMAND} -E copy "${java_source_dir}/${f}" "${opencv_test_java_bin_dir}/${f}"
MAIN_DEPENDENCY "${java_source_dir}/${f}"
COMMENT "Copying ${f}")
list(APPEND opencv_test_java_file_deps "${java_source_dir}/${f}")
endforeach()
# Copy the OpenCV jar after it has been generated.
add_custom_command(
OUTPUT "${opencv_test_java_bin_dir}/bin/${JAR_NAME}"
COMMAND ${CMAKE_COMMAND} -E copy "${CMAKE_BINARY_DIR}/bin/${JAR_NAME}" "${opencv_test_java_bin_dir}/bin/${JAR_NAME}"
COMMENT "Copying the OpenCV jar")
add_custom_target(copy_opencv_jar ALL SOURCES "${opencv_test_java_bin_dir}/bin/${JAR_NAME}")
# ${the_module} is the target for the Java jar.
add_dependencies(copy_opencv_jar ${the_module})
# Copy the ant build file.
file(COPY "${CMAKE_CURRENT_SOURCE_DIR}/build.xml" DESTINATION "${opencv_test_java_bin_dir}")
# Create a script for running the Java tests and place it in build/bin.
if(WIN32)
file(WRITE "${CMAKE_BINARY_DIR}/bin/opencv_test_java.cmd" "cd ${opencv_test_java_bin_dir}\nset PATH=${EXECUTABLE_OUTPUT_PATH}/Release;%PATH%\nant -DjavaLibraryPath=${EXECUTABLE_OUTPUT_PATH}/Release buildAndTest")
file(WRITE "${CMAKE_BINARY_DIR}/bin/opencv_test_java_D.cmd" "cd ${opencv_test_java_bin_dir}\nset PATH=${EXECUTABLE_OUTPUT_PATH}/Debug;%PATH%\nant -DjavaLibraryPath=${EXECUTABLE_OUTPUT_PATH}/Debug buildAndTest")
else()
file(WRITE "${CMAKE_BINARY_DIR}/bin/opencv_test_java.sh" "cd ${opencv_test_java_bin_dir};\nant -DjavaLibraryPath=${LIBRARY_OUTPUT_PATH} buildAndTest;\ncd -")
endif()
add_custom_target(${PROJECT_NAME} ALL SOURCES ${opencv_test_java_file_deps})
add_dependencies(opencv_tests ${PROJECT_NAME})

77
modules/java/test/CMakeLists.txt Normal file
View File

@ -0,0 +1,77 @@
ocv_check_dependencies(opencv_java ${OPENCV_MODULE_opencv_java_OPT_DEPS} ${OPENCV_MODULE_opencv_java_REQ_DEPS})
if(NOT OCV_DEPENDENCIES_FOUND)
return()
endif()
project(opencv_test_java)
set(opencv_test_java_bin_dir "${CMAKE_CURRENT_BINARY_DIR}/.build")
set(android_source_dir "${CMAKE_CURRENT_SOURCE_DIR}/../android_test")
set(java_source_dir ${CMAKE_CURRENT_SOURCE_DIR})
# get project sources
file(GLOB_RECURSE opencv_test_java_files RELATIVE "${android_source_dir}" "${android_source_dir}/res/*" "${android_source_dir}/src/*.java")
# These are the files that need to be updated for pure Java.
ocv_list_filterout(opencv_test_java_files "OpenCVTest(Case|Runner).java")
# These files aren't for desktop Java.
ocv_list_filterout(opencv_test_java_files "/android/")
# These are files updated for pure Java.
file(GLOB_RECURSE modified_files RELATIVE "${java_source_dir}" "${java_source_dir}/src/*")
# These are extra jars needed to run the tests.
file(GLOB_RECURSE lib_files RELATIVE "${java_source_dir}" "${java_source_dir}/lib/*.jar")
# copy sources out from the build tree
set(opencv_test_java_file_deps "")
foreach(f ${opencv_test_java_files})
add_custom_command(OUTPUT "${opencv_test_java_bin_dir}/${f}"
COMMAND ${CMAKE_COMMAND} -E copy "${android_source_dir}/${f}" "${opencv_test_java_bin_dir}/${f}"
DEPENDS "${android_source_dir}/${f}"
COMMENT "Copying ${f}"
)
list(APPEND opencv_test_java_file_deps "${android_source_dir}/${f}" "${opencv_test_java_bin_dir}/${f}")
endforeach()
# Overwrite select Android sources with Java-specific sources.
# Also, copy over the libs we'll need for testing.
foreach(f ${modified_files} ${lib_files})
add_custom_command(OUTPUT "${opencv_test_java_bin_dir}/${f}"
COMMAND ${CMAKE_COMMAND} -E copy "${java_source_dir}/${f}" "${opencv_test_java_bin_dir}/${f}"
DEPENDS "${java_source_dir}/${f}"
COMMENT "Copying ${f}"
)
list(APPEND opencv_test_java_file_deps "${java_source_dir}/${f}" "${opencv_test_java_bin_dir}/${f}")
endforeach()
# Copy the OpenCV jar after it has been generated.
add_custom_command(OUTPUT "${opencv_test_java_bin_dir}/bin/${JAR_NAME}"
COMMAND ${CMAKE_COMMAND} -E copy "${JAR_FILE}" "${opencv_test_java_bin_dir}/bin/${JAR_NAME}"
DEPENDS "${JAR_FILE}"
COMMENT "Copying the OpenCV jar"
)
add_custom_command(OUTPUT "${opencv_test_java_bin_dir}/build.xml"
COMMAND ${CMAKE_COMMAND} -E copy "${CMAKE_CURRENT_SOURCE_DIR}/build.xml" "${opencv_test_java_bin_dir}/build.xml"
DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/build.xml"
COMMENT "Copying build.xml"
)
# Create a script for running the Java tests and place it in build/bin.
#if(WIN32)
#file(WRITE "${CMAKE_BINARY_DIR}/bin/opencv_test_java.cmd" "cd ${opencv_test_java_bin_dir}\nset PATH=${EXECUTABLE_OUTPUT_PATH}/Release;%PATH%\nant -DjavaLibraryPath=${EXECUTABLE_OUTPUT_PATH}/Release buildAndTest")
#file(WRITE "${CMAKE_BINARY_DIR}/bin/opencv_test_java_D.cmd" "cd ${opencv_test_java_bin_dir}\nset PATH=${EXECUTABLE_OUTPUT_PATH}/Debug;%PATH%\nant -DjavaLibraryPath=${EXECUTABLE_OUTPUT_PATH}/Debug buildAndTest")
#else()
#file(WRITE "${CMAKE_BINARY_DIR}/bin/opencv_test_java.sh" "cd ${opencv_test_java_bin_dir};\nant -DjavaLibraryPath=${LIBRARY_OUTPUT_PATH} buildAndTest;\ncd -")
#endif()
add_custom_command(OUTPUT "${opencv_test_java_bin_dir}/build/jar/opencv-test.jar"
COMMAND "${ANT_EXECUTABLE}" build
WORKING_DIRECTORY "${opencv_test_java_bin_dir}"
DEPENDS ${opencv_test_java_file_deps} "${opencv_test_java_bin_dir}/build.xml" "${CMAKE_CURRENT_SOURCE_DIR}/build.xml" "${JAR_FILE}" "${opencv_test_java_bin_dir}/bin/${JAR_NAME}"
COMMENT "Build Java tests"
)
add_custom_target(${PROJECT_NAME} ALL SOURCES "${opencv_test_java_bin_dir}/build/jar/opencv-test.jar")
add_dependencies(${PROJECT_NAME} ${the_module})

18
modules/java/java_test/build.xml → modules/java/test/build.xml
View File

@ -16,7 +16,7 @@
<target name="compile">
<mkdir dir="build/classes"/>
<javac sourcepath="" srcdir="src" destdir="build/classes" >
<javac sourcepath="" srcdir="src" destdir="build/classes" includeantruntime="false" >
<include name="**/*.java"/>
<classpath refid="master-classpath"/>
</javac>
@ -26,15 +26,16 @@
<mkdir dir="build/jar"/>
<jar destfile="build/jar/opencv-test.jar" basedir="build/classes">
<manifest>
<attribute name="Main-Class" value="org.opencv.test.OpenCVTestRunner"/>
<attribute name="Main-Class" value="org.opencv.test.OpenCVTestRunner"/>
</manifest>
</jar>
</target>
<target name="test">
<mkdir dir="testResults"/>
<junit printsummary="false" haltonfailure="false" haltonerror="false" showoutput="false" logfailedtests="true" maxmemory="256m">
<junit printsummary="true" haltonfailure="false" haltonerror="false" showoutput="false" logfailedtests="true" maxmemory="256m">
<sysproperty key="java.library.path" path="${javaLibraryPath}"/>
<env key="PATH" path="${javaLibraryPath}"/>
<classpath refid="master-classpath"/>
<classpath>
<pathelement location="build/classes"/>
@ -43,13 +44,18 @@
<formatter type="xml"/>
<batchtest fork="yes" todir="testResults">
<zipfileset src="build/jar/opencv-test.jar" includes="**/*.class" excludes="**/OpenCVTest*">
<exclude name="**/*$*.class"/>
</zipfileset>
<zipfileset src="build/jar/opencv-test.jar" includes="**/*.class" excludes="**/OpenCVTest*">
<exclude name="**/*$*.class"/>
</zipfileset>
</batchtest>
</junit>
</target>
<target name="build">
<antcall target="compile"/>
<antcall target="jar"/>
</target>
<target name="buildAndTest">
<antcall target="compile"/>
<antcall target="jar"/>

0
modules/java/java_test/lib/junit-4.11.jar → modules/java/test/lib/junit-4.11.jar
View File

46
modules/java/java_test/src/org/opencv/test/OpenCVTestCase.java → modules/java/test/src/org/opencv/test/OpenCVTestCase.java
View File

@ -96,31 +96,33 @@ public class OpenCVTestCase extends TestCase {
protected void setUp() throws Exception {
super.setUp();
try {
System.loadLibrary("opencv_java");
} catch (SecurityException e) {
System.out.println(e.toString());
System.exit(-1);
} catch (UnsatisfiedLinkError e) {
System.out.println(e.toString());
System.exit(-1);
}
try {
System.loadLibrary("opencv_java");
} catch (SecurityException e) {
System.out.println(e.toString());
System.exit(-1);
} catch (UnsatisfiedLinkError e) {
System.out.println(e.toString());
System.exit(-1);
}
String pwd;
try {
pwd = new File(".").getCanonicalPath() + File.separator;
} catch (IOException e) {
System.out.println(e);
return;
}
Core.setErrorVerbosity(false);
OpenCVTestRunner.LENA_PATH = pwd + "res/drawable/lena.jpg";
OpenCVTestRunner.CHESS_PATH = pwd + "res/drawable/chessboard.jpg";
OpenCVTestRunner.LBPCASCADE_FRONTALFACE_PATH = pwd + "res/raw/lbpcascade_frontalface.xml";
String pwd;
try {
pwd = new File(".").getCanonicalPath() + File.separator;
} catch (IOException e) {
System.out.println(e);
return;
}
assert(new File(OpenCVTestRunner.LENA_PATH).exists());
assert(new File(OpenCVTestRunner.CHESS_PATH).exists());
assert(new File(OpenCVTestRunner.LBPCASCADE_FRONTALFACE_PATH).exists());
OpenCVTestRunner.LENA_PATH = pwd + "res/drawable/lena.jpg";
OpenCVTestRunner.CHESS_PATH = pwd + "res/drawable/chessboard.jpg";
OpenCVTestRunner.LBPCASCADE_FRONTALFACE_PATH = pwd + "res/raw/lbpcascade_frontalface.xml";
assert(new File(OpenCVTestRunner.LENA_PATH).exists());
assert(new File(OpenCVTestRunner.CHESS_PATH).exists());
assert(new File(OpenCVTestRunner.LBPCASCADE_FRONTALFACE_PATH).exists());
dst = new Mat();
assertTrue(dst.empty());

0
modules/java/java_test/src/org/opencv/test/OpenCVTestRunner.java → modules/java/test/src/org/opencv/test/OpenCVTestRunner.java
View File

2
modules/nonfree/src/surf.cpp
View File

@ -482,6 +482,8 @@ static void fastHessianDetector( const Mat& sum, const Mat& mask_sum, vector<Key
vector<int> sampleSteps(nTotalLayers);
vector<int> middleIndices(nMiddleLayers);
keypoints.clear();
// Allocate space and calculate properties of each layer
int index = 0, middleIndex = 0, step = SAMPLE_STEP0;

19
modules/nonfree/test/test_rotation_and_scale_invariance.cpp
View File

@ -668,3 +668,22 @@ TEST(Features2d_ScaleInvariance_Descriptor_SIFT, regression)
0.87f);
test.safe_run();
}
TEST(Features2d_RotationInvariance2_Detector_SURF, regression)
{
Mat cross(100, 100, CV_8UC1, Scalar(255));
line(cross, Point(30, 50), Point(69, 50), Scalar(100), 3);
line(cross, Point(50, 30), Point(50, 69), Scalar(100), 3);
SURF surf(8000., 3, 4, true, false);
vector<KeyPoint> keypoints;
surf(cross, noArray(), keypoints);
ASSERT_EQ(keypoints.size(), (vector<KeyPoint>::size_type) 5);
ASSERT_LT( fabs(keypoints[1].response - keypoints[2].response), 1e-6);
ASSERT_LT( fabs(keypoints[1].response - keypoints[3].response), 1e-6);
ASSERT_LT( fabs(keypoints[1].response - keypoints[4].response), 1e-6);
}

3
modules/objdetect/src/cascadedetect.cpp
View File

@ -237,7 +237,7 @@ class MeanshiftGrouping
{
public:
MeanshiftGrouping(const Point3d& densKer, const vector<Point3d>& posV,
const vector<double>& wV, double, int maxIter = 20)
const vector<double>& wV, double eps, int maxIter = 20)
{
densityKernel = densKer;
weightsV = wV;
@ -246,6 +246,7 @@ public:
meanshiftV.resize(positionsCount);
distanceV.resize(positionsCount);
iterMax = maxIter;
modeEps = eps;
for (unsigned i = 0; i<positionsV.size(); i++)
{

2
modules/ocl/CMakeLists.txt
View File

@ -39,6 +39,8 @@ if (HAVE_OPENCL)
endif()
endif()
ocv_warnings_disable(CMAKE_CXX_FLAGS -Wshadow)
ocv_set_module_sources(HEADERS ${lib_hdrs} SOURCES ${lib_int_hdrs} ${lib_srcs} ${kernels_cpp})
ocv_create_module(${ocl_link_libs})

11
modules/ocl/include/opencv2/ocl/ocl.hpp
View File

@ -1754,12 +1754,13 @@ namespace cv
}
}
#if _MSC_VER >= 1200
#pragma warning( push)
#pragma warning( disable: 4267)
#if defined _MSC_VER && _MSC_VER >= 1200
# pragma warning( push)
# pragma warning( disable: 4267)
#endif
#include "opencv2/ocl/matrix_operations.hpp"
#if _MSC_VER >= 1200
#pragma warning( pop)
#if defined _MSC_VER && _MSC_VER >= 1200
# pragma warning( pop)
#endif
#endif /* __OPENCV_OCL_HPP__ */

10
modules/ocl/perf/main.cpp
View File

@ -97,12 +97,12 @@ int main(int argc, char **argv)
unsigned int pid = cmd.get<unsigned int>("p");
int device = cmd.get<int>("d");
print_info();
int flag = CVCL_DEVICE_TYPE_GPU;
// int flag = CVCL_DEVICE_TYPE_GPU;
if(type == "cpu")
{
flag = CVCL_DEVICE_TYPE_CPU;
}
// if(type == "cpu")
// {
// flag = CVCL_DEVICE_TYPE_CPU;
// }
std::vector<cv::ocl::Info> oclinfo;
int devnums = getDevice(oclinfo);
if(devnums <= device || device < 0)

4
modules/ocl/perf/perf_arithm.cpp
View File

@ -2597,13 +2597,13 @@ TEST_P(Sum, MAT)
Has_roi(k);
t0 = (double)cvGetTickCount();//cpu start
Scalar cpures = cv::sum(mat1_roi);
cv::sum(mat1_roi);
t0 = (double)cvGetTickCount() - t0;//cpu end
t1 = (double)cvGetTickCount();//gpu start1
gmat1 = mat1_roi;
t2 = (double)cvGetTickCount(); //kernel
Scalar gpures = cv::ocl::sum(gmat1);
cv::ocl::sum(gmat1);
t2 = (double)cvGetTickCount() - t2;//kernel
t1 = (double)cvGetTickCount() - t1;//gpu end1
if(j == 0)

18
modules/ocl/perf/perf_haar.cpp
View File

@ -109,15 +109,15 @@ TEST_F(Haar, FaceDetect)
double t = 0;
vector<Rect> faces, oclfaces;
const static Scalar colors[] = { CV_RGB(0, 0, 255),
CV_RGB(0, 128, 255),
CV_RGB(0, 255, 255),
CV_RGB(0, 255, 0),
CV_RGB(255, 128, 0),
CV_RGB(255, 255, 0),
CV_RGB(255, 0, 0),
CV_RGB(255, 0, 255)
} ;
// const static Scalar colors[] = { CV_RGB(0, 0, 255),
// CV_RGB(0, 128, 255),
// CV_RGB(0, 255, 255),
// CV_RGB(0, 255, 0),
// CV_RGB(255, 128, 0),
// CV_RGB(255, 255, 0),
// CV_RGB(255, 0, 0),
// CV_RGB(255, 0, 255)
// } ;
Mat gray, smallImg(cvRound (img.rows / scale), cvRound(img.cols / scale), CV_8UC1 );
MemStorage storage(cvCreateMemStorage(0));

7
modules/ocl/perf/perf_imgproc.cpp
View File

@ -379,7 +379,7 @@ TEST_P(bilateralFilter, Mat)
}
else
{
for(int i = 0; i < sizeof(bordertype) / sizeof(int); i++)
for(size_t i = 0; i < sizeof(bordertype) / sizeof(int); i++)
{
cout << borderstr[i] << endl;
#ifndef PRINT_KERNEL_RUN_TIME
@ -397,7 +397,7 @@ TEST_P(bilateralFilter, Mat)
for(int j = 0; j < LOOP_TIMES + 1; j ++)
{
Has_roi(k);
if(((bordertype[i] != cv::BORDER_CONSTANT) && (bordertype[i] != cv::BORDER_REPLICATE)) && (mat1_roi.cols <= radius) || (mat1_roi.cols <= radius) || (mat1_roi.rows <= radius) || (mat1_roi.rows <= radius))
if(((bordertype[i] != cv::BORDER_CONSTANT) && (bordertype[i] != cv::BORDER_REPLICATE) && (mat1_roi.cols <= radius)) || (mat1_roi.cols <= radius) || (mat1_roi.rows <= radius) || (mat1_roi.rows <= radius))
{
continue;
}
@ -482,7 +482,7 @@ TEST_P(CopyMakeBorder, Mat)
}
else
{
for(int i = 0; i < sizeof(bordertype) / sizeof(int); i++)
for(size_t i = 0; i < sizeof(bordertype) / sizeof(int); i++)
{
#ifndef PRINT_KERNEL_RUN_TIME
double totalcputick = 0;
@ -1133,7 +1133,6 @@ PARAM_TEST_CASE(Remap, MatType, MatType, MatType, int, int)
cv::RNG &rng = TS::ptr()->get_rng();
cv::Size srcSize = cv::Size(MWIDTH, MHEIGHT);
cv::Size dstSize = cv::Size(MWIDTH, MHEIGHT);
cv::Size map1Size = cv::Size(MWIDTH, MHEIGHT);
double min = 5, max = 16;

9
modules/ocl/perf/precomp.hpp
View File

@ -38,6 +38,15 @@
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#ifdef __GNUC__
# pragma GCC diagnostic ignored "-Wmissing-declarations"
# if defined __clang__ || defined __APPLE__
# pragma GCC diagnostic ignored "-Wmissing-prototypes"
# pragma GCC diagnostic ignored "-Wextra"
# endif
#endif
#ifndef __OPENCV_TEST_PRECOMP_HPP__
#define __OPENCV_TEST_PRECOMP_HPP__

12
modules/ocl/perf/utility.hpp
View File

@ -84,12 +84,12 @@ double checkSimilarity(const cv::Mat &m1, const cv::Mat &m2);
EXPECT_LE(checkNorm(cv::Mat(mat)), eps) \
}
//#define EXPECT_MAT_NEAR(mat1, mat2, eps) \
//{ \
// ASSERT_EQ(mat1.type(), mat2.type()); \
// ASSERT_EQ(mat1.size(), mat2.size()); \
// EXPECT_LE(checkNorm(cv::Mat(mat1), cv::Mat(mat2)), eps); \
//}
/*#define EXPECT_MAT_NEAR(mat1, mat2, eps) \
{ \
ASSERT_EQ(mat1.type(), mat2.type()); \
ASSERT_EQ(mat1.size(), mat2.size()); \
EXPECT_LE(checkNorm(cv::Mat(mat1), cv::Mat(mat2)), eps); \
}*/
#define EXPECT_MAT_NEAR(mat1, mat2, eps,s) \
{ \

60
modules/ocl/src/arithm.cpp
View File

@ -362,11 +362,11 @@ void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string
openCLExecuteKernel(clCxt, kernelString, kernelName, globalThreads, localThreads, args, -1, depth);
}
void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName, const char **kernelString)
static void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName, const char **kernelString)
{
arithmetic_run<char>(src1, src2, dst, kernelName, kernelString, (void *)NULL);
}
void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString)
static void arithmetic_run(const oclMat &src1, const oclMat &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString)
{
if(src1.clCxt -> impl -> double_support == 0 && src1.type() == CV_64F)
{
@ -476,7 +476,9 @@ void arithmetic_scalar_run(const oclMat &src1, const Scalar &src2, oclMat &dst,
//CV_Assert(src1.depth() != CV_8S);
if(mask.data)
{
CV_Assert(mask.type() == CV_8U && src1.rows == mask.rows && src1.cols == mask.cols);
}
Context *clCxt = src1.clCxt;
int channels = dst.oclchannels();
@ -530,7 +532,7 @@ void arithmetic_scalar_run(const oclMat &src1, const Scalar &src2, oclMat &dst,
openCLExecuteKernel(clCxt, kernelString, kernelName, globalThreads, localThreads, args, channels, depth);
}
void arithmetic_scalar_run(const oclMat &src, oclMat &dst, string kernelName, const char **kernelString, double scalar)
static void arithmetic_scalar_run(const oclMat &src, oclMat &dst, string kernelName, const char **kernelString, double scalar)
{
if(src.clCxt -> impl -> double_support == 0 && src.type() == CV_64F)
{
@ -590,7 +592,7 @@ void arithmetic_scalar_run(const oclMat &src, oclMat &dst, string kernelName, co
typedef void (*ArithmeticFuncS)(const oclMat &src1, const Scalar &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString, int isMatSubScalar);
void arithmetic_scalar(const oclMat &src1, const Scalar &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString, int isMatSubScalar)
static void arithmetic_scalar(const oclMat &src1, const Scalar &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString, int isMatSubScalar)
{
static ArithmeticFuncS tab[8] =
{
@ -608,7 +610,7 @@ void arithmetic_scalar(const oclMat &src1, const Scalar &src2, oclMat &dst, cons
cv::ocl::error("Unsupported arithmetic operation", __FILE__, __LINE__);
func(src1, src2, dst, mask, kernelName, kernelString, isMatSubScalar);
}
void arithmetic_scalar(const oclMat &src1, const Scalar &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString)
static void arithmetic_scalar(const oclMat &src1, const Scalar &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString)
{
arithmetic_scalar(src1, src2, dst, mask, kernelName, kernelString, 0);
}
@ -660,7 +662,7 @@ void cv::ocl::absdiff(const oclMat &src1, const Scalar &src2, oclMat &dst)
//////////////////////////////////////////////////////////////////////////////
///////////////////////////////// compare ///////////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void compare_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName, const char **kernelString)
static void compare_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName, const char **kernelString)
{
dst.create(src1.size(), CV_8UC1);
CV_Assert(src1.oclchannels() == 1);
@ -739,7 +741,7 @@ void cv::ocl::compare(const oclMat &src1, const oclMat &src2, oclMat &dst , int
//////////////////////////////////////////////////////////////////////////////
//type = 0 sum,type = 1 absSum,type = 2 sqrSum
void arithmetic_sum_buffer_run(const oclMat &src, cl_mem &dst, int vlen , int groupnum, int type = 0)
static void arithmetic_sum_buffer_run(const oclMat &src, cl_mem &dst, int vlen , int groupnum, int type = 0)
{
vector<pair<size_t , const void *> > args;
int all_cols = src.step / (vlen * src.elemSize1());
@ -872,7 +874,7 @@ void cv::ocl::meanStdDev(const oclMat &src, Scalar &mean, Scalar &stddev)
//////////////////////////////////////////////////////////////////////////////
//////////////////////////////////// minMax /////////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void arithmetic_minMax_run(const oclMat &src, const oclMat &mask, cl_mem &dst, int vlen , int groupnum, string kernelName)
static void arithmetic_minMax_run(const oclMat &src, const oclMat &mask, cl_mem &dst, int vlen , int groupnum, string kernelName)
{
vector<pair<size_t , const void *> > args;
int all_cols = src.step / (vlen * src.elemSize1());
@ -909,7 +911,7 @@ void arithmetic_minMax_run(const oclMat &src, const oclMat &mask, cl_mem &dst, i
}
void arithmetic_minMax_mask_run(const oclMat &src, const oclMat &mask, cl_mem &dst, int vlen, int groupnum, string kernelName)
static void arithmetic_minMax_mask_run(const oclMat &src, const oclMat &mask, cl_mem &dst, int vlen, int groupnum, string kernelName)
{
vector<pair<size_t , const void *> > args;
size_t gt[3] = {groupnum * 256, 1, 1}, lt[3] = {256, 1, 1};
@ -1063,7 +1065,7 @@ double cv::ocl::norm(const oclMat &src1, const oclMat &src2, int normType)
//////////////////////////////////////////////////////////////////////////////
////////////////////////////////// flip //////////////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void arithmetic_flip_rows_run(const oclMat &src, oclMat &dst, string kernelName)
static void arithmetic_flip_rows_run(const oclMat &src, oclMat &dst, string kernelName)
{
if(src.clCxt -> impl -> double_support == 0 && src.type() == CV_64F)
{
@ -1112,7 +1114,7 @@ void arithmetic_flip_rows_run(const oclMat &src, oclMat &dst, string kernelName)
openCLExecuteKernel(clCxt, &arithm_flip, kernelName, globalThreads, localThreads, args, -1, depth);
}
void arithmetic_flip_cols_run(const oclMat &src, oclMat &dst, string kernelName, bool isVertical)
static void arithmetic_flip_cols_run(const oclMat &src, oclMat &dst, string kernelName, bool isVertical)
{
if(src.clCxt -> impl -> double_support == 0 && src.type() == CV_64F)
{
@ -1183,7 +1185,7 @@ void cv::ocl::flip(const oclMat &src, oclMat &dst, int flipCode)
//////////////////////////////////////////////////////////////////////////////
////////////////////////////////// LUT //////////////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void arithmetic_lut_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName)
static void arithmetic_lut_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName)
{
Context *clCxt = src1.clCxt;
int channels = src1.oclchannels();
@ -1284,7 +1286,7 @@ void cv::ocl::LUT(const oclMat &src, const oclMat &lut, oclMat &dst)
//////////////////////////////////////////////////////////////////////////////
//////////////////////////////// exp log /////////////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void arithmetic_exp_log_run(const oclMat &src, oclMat &dst, string kernelName, const char **kernelString)
static void arithmetic_exp_log_run(const oclMat &src, oclMat &dst, string kernelName, const char **kernelString)
{
dst.create(src.size(), src.type());
CV_Assert(src.cols == dst.cols &&
@ -1333,7 +1335,7 @@ void cv::ocl::log(const oclMat &src, oclMat &dst)
//////////////////////////////////////////////////////////////////////////////
////////////////////////////// magnitude phase ///////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void arithmetic_magnitude_phase_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName)
static void arithmetic_magnitude_phase_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName)
{
if(src1.clCxt -> impl -> double_support == 0 && src1.type() == CV_64F)
{
@ -1381,7 +1383,7 @@ void cv::ocl::magnitude(const oclMat &src1, const oclMat &src2, oclMat &dst)
arithmetic_magnitude_phase_run(src1, src2, dst, "arithm_magnitude");
}
void arithmetic_phase_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName, const char **kernelString)
static void arithmetic_phase_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName, const char **kernelString)
{
if(src1.clCxt -> impl -> double_support == 0 && src1.type() == CV_64F)
{
@ -1444,7 +1446,7 @@ void cv::ocl::phase(const oclMat &x, const oclMat &y, oclMat &Angle , bool angle
//////////////////////////////////////////////////////////////////////////////
////////////////////////////////// cartToPolar ///////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void arithmetic_cartToPolar_run(const oclMat &src1, const oclMat &src2, oclMat &dst_mag, oclMat &dst_cart,
static void arithmetic_cartToPolar_run(const oclMat &src1, const oclMat &src2, oclMat &dst_mag, oclMat &dst_cart,
string kernelName, bool angleInDegrees)
{
if(src1.clCxt -> impl -> double_support == 0 && src1.type() == CV_64F)
@ -1499,7 +1501,7 @@ void cv::ocl::cartToPolar(const oclMat &x, const oclMat &y, oclMat &mag, oclMat
//////////////////////////////////////////////////////////////////////////////
////////////////////////////////// polarToCart ///////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void arithmetic_ptc_run(const oclMat &src1, const oclMat &src2, oclMat &dst1, oclMat &dst2, bool angleInDegrees,
static void arithmetic_ptc_run(const oclMat &src1, const oclMat &src2, oclMat &dst1, oclMat &dst2, bool angleInDegrees,
string kernelName)
{
if(src1.clCxt -> impl -> double_support == 0 && src1.type() == CV_64F)
@ -1564,7 +1566,7 @@ void cv::ocl::polarToCart(const oclMat &magnitude, const oclMat &angle, oclMat &
//////////////////////////////////////////////////////////////////////////////
/////////////////////////////////// minMaxLoc ////////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void arithmetic_minMaxLoc_run(const oclMat &src, cl_mem &dst, int vlen , int groupnum)
static void arithmetic_minMaxLoc_run(const oclMat &src, cl_mem &dst, int vlen , int groupnum)
{
vector<pair<size_t , const void *> > args;
int all_cols = src.step / (vlen * src.elemSize1());
@ -1588,7 +1590,7 @@ void arithmetic_minMaxLoc_run(const oclMat &src, cl_mem &dst, int vlen , int gro
openCLExecuteKernel(src.clCxt, &arithm_minMaxLoc, "arithm_op_minMaxLoc", gt, lt, args, -1, -1, build_options);
}
void arithmetic_minMaxLoc_mask_run(const oclMat &src, const oclMat &mask, cl_mem &dst, int vlen, int groupnum)
static void arithmetic_minMaxLoc_mask_run(const oclMat &src, const oclMat &mask, cl_mem &dst, int vlen, int groupnum)
{
vector<pair<size_t , const void *> > args;
size_t gt[3] = {groupnum * 256, 1, 1}, lt[3] = {256, 1, 1};
@ -1702,7 +1704,7 @@ void cv::ocl::minMaxLoc(const oclMat &src, double *minVal, double *maxVal,
//////////////////////////////////////////////////////////////////////////////
///////////////////////////// countNonZero ///////////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void arithmetic_countNonZero_run(const oclMat &src, cl_mem &dst, int vlen , int groupnum, string kernelName)
static void arithmetic_countNonZero_run(const oclMat &src, cl_mem &dst, int vlen , int groupnum, string kernelName)
{
vector<pair<size_t , const void *> > args;
int all_cols = src.step / (vlen * src.elemSize1());
@ -1759,7 +1761,7 @@ int cv::ocl::countNonZero(const oclMat &src)
//////////////////////////////////////////////////////////////////////////////
////////////////////////////////bitwise_op////////////////////////////////////
//////////////////////////////////////////////////////////////////////////////
void bitwise_run(const oclMat &src1, oclMat &dst, string kernelName, const char **kernelString)
static void bitwise_run(const oclMat &src1, oclMat &dst, string kernelName, const char **kernelString)
{
dst.create(src1.size(), src1.type());
@ -1853,11 +1855,11 @@ void bitwise_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string ker
openCLExecuteKernel(clCxt, kernelString, kernelName, globalThreads, localThreads, args, -1, depth);
}
void bitwise_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName, const char **kernelString)
static void bitwise_run(const oclMat &src1, const oclMat &src2, oclMat &dst, string kernelName, const char **kernelString)
{
bitwise_run<char>(src1, src2, dst, kernelName, kernelString, (void *)NULL);
}
void bitwise_run(const oclMat &src1, const oclMat &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString)
static void bitwise_run(const oclMat &src1, const oclMat &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString)
{
dst.create(src1.size(), src1.type());
CV_Assert(src1.cols == src2.cols && src2.cols == dst.cols &&
@ -1919,7 +1921,9 @@ void bitwise_scalar_run(const oclMat &src1, const Scalar &src2, oclMat &dst, con
if(mask.data)
{
CV_Assert(mask.type() == CV_8U && src1.rows == mask.rows && src1.cols == mask.cols);
}
Context *clCxt = src1.clCxt;
int channels = dst.oclchannels();
@ -1977,7 +1981,7 @@ void bitwise_scalar_run(const oclMat &src1, const Scalar &src2, oclMat &dst, con
typedef void (*BitwiseFuncS)(const oclMat &src1, const Scalar &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString, int isMatSubScalar);
void bitwise_scalar(const oclMat &src1, const Scalar &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString, int isMatSubScalar)
static void bitwise_scalar(const oclMat &src1, const Scalar &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString, int isMatSubScalar)
{
static BitwiseFuncS tab[8] =
{
@ -2007,7 +2011,7 @@ void bitwise_scalar(const oclMat &src1, const Scalar &src2, oclMat &dst, const o
cv::ocl::error("Unsupported arithmetic operation", __FILE__, __LINE__);
func(src1, src2, dst, mask, kernelName, kernelString, isMatSubScalar);
}
void bitwise_scalar(const oclMat &src1, const Scalar &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString)
static void bitwise_scalar(const oclMat &src1, const Scalar &src2, oclMat &dst, const oclMat &mask, string kernelName, const char **kernelString)
{
bitwise_scalar(src1, src2, dst, mask, kernelName, kernelString, 0);
}
@ -2153,7 +2157,7 @@ cv::ocl::oclMat cv::ocl::operator ^ (const oclMat &src1, const oclMat &src2)
//////////////////////////////////////////////////////////////////////////////
#define TILE_DIM (32)
#define BLOCK_ROWS (256/TILE_DIM)
void transpose_run(const oclMat &src, oclMat &dst, string kernelName)
static void transpose_run(const oclMat &src, oclMat &dst, string kernelName)
{
if(src.clCxt -> impl -> double_support == 0 && src.type() == CV_64F)
{
@ -2368,7 +2372,7 @@ void cv::ocl::magnitudeSqr(const oclMat &src1, oclMat &dst)
openCLExecuteKernel(clCxt, &arithm_magnitudeSqr, "magnitudeSqr", globalThreads, localThreads, args, 2, depth);
}
void arithmetic_pow_run(const oclMat &src1, double p, oclMat &dst, string kernelName, const char **kernelString)
static void arithmetic_pow_run(const oclMat &src1, double p, oclMat &dst, string kernelName, const char **kernelString)
{
CV_Assert(src1.cols == dst.cols && src1.rows == dst.rows);
CV_Assert(src1.type() == dst.type());
@ -2417,7 +2421,7 @@ void cv::ocl::pow(const oclMat &x, double p, oclMat &y)
return;
}
CV_Assert(x.type() == y.type() && x.size() == y.size() && x.depth() == CV_32F || x.depth() == CV_64F);
CV_Assert((x.type() == y.type() && x.size() == y.size() && x.depth() == CV_32F) || x.depth() == CV_64F);
y.create(x.size(), x.type());
string kernelName = "arithm_pow";

4
modules/ocl/src/brute_force_matcher.cpp
View File

@ -44,7 +44,7 @@
//M*/
#include "precomp.hpp"
#include <functional>
#include <iterator>
#include <vector>
using namespace cv;
@ -821,7 +821,7 @@ void findKnnMatch(int k, const oclMat &trainIdx, const oclMat &distance, const o
}
}
void findKnnMatchDispatcher(int k, const oclMat &trainIdx, const oclMat &distance, const oclMat &allDist, int distType)
static void findKnnMatchDispatcher(int k, const oclMat &trainIdx, const oclMat &distance, const oclMat &allDist, int distType)
{
findKnnMatch<256>(k, trainIdx, distance, allDist, distType);
}

268
modules/ocl/src/color.cpp
View File

@ -16,6 +16,7 @@
//
// @Authors
// Wang Weiyan, wangweiyanster@gmail.com
// Peng Xiao, pengxiao@multicorewareinc.com
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
@ -70,79 +71,218 @@ void cv::ocl::cvtColor(const oclMat &, oclMat &, int, int, const Stream &)
namespace cv
{
namespace ocl
{
extern const char *cvt_color;
}
namespace ocl
{
extern const char *cvt_color;
}
}
namespace
{
void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx)
{
vector<pair<size_t , const void *> > args;
int channels = src.oclchannels();
char build_options[50];
//printf("depth:%d,channels:%d,bidx:%d\n",src.depth(),src.oclchannels(),bidx);
sprintf(build_options, "-D DEPTH_%d", src.depth());
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&channels));
args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data));
size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1};
openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2Gray", gt, lt, args, -1, -1, build_options);
}
void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int /*dcn*/)
{
Size sz = src.size();
int scn = src.oclchannels(), depth = src.depth(), bidx;
void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx)
{
vector<pair<size_t , const void *> > args;
int channels = src.oclchannels();
char build_options[50];
sprintf(build_options, "-D DEPTH_%d", src.depth());
//printf("depth:%d,channels:%d,bidx:%d\n",src.depth(),src.oclchannels(),bidx);
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&channels));
args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data));
size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1};
openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2Gray", gt, lt, args, -1, -1, build_options);
}
void Gray2RGB_caller(const oclMat &src, oclMat &dst)
{
vector<pair<size_t , const void *> > args;
char build_options[50];
sprintf(build_options, "-D DEPTH_%d", src.depth());
//printf("depth:%d,channels:%d,bidx:%d\n",src.depth(),src.oclchannels(),bidx);
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.step));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data));
size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1};
openCLExecuteKernel(src.clCxt, &cvt_color, "Gray2RGB", gt, lt, args, -1, -1, build_options);
}
void RGB2YUV_caller(const oclMat &src, oclMat &dst, int bidx)
{
vector<pair<size_t , const void *> > args;
int channels = src.oclchannels();
char build_options[50];
sprintf(build_options, "-D DEPTH_%d", src.depth());
//printf("depth:%d,channels:%d,bidx:%d\n",src.depth(),src.oclchannels(),bidx);
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&channels));
args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data));
size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1};
openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2YUV", gt, lt, args, -1, -1, build_options);
}
void YUV2RGB_caller(const oclMat &src, oclMat &dst, int bidx)
{
vector<pair<size_t , const void *> > args;
int channels = src.oclchannels();
char build_options[50];
sprintf(build_options, "-D DEPTH_%d", src.depth());
//printf("depth:%d,channels:%d,bidx:%d\n",src.depth(),src.oclchannels(),bidx);
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&channels));
args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data));
size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1};
openCLExecuteKernel(src.clCxt, &cvt_color, "YUV2RGB", gt, lt, args, -1, -1, build_options);
}
void YUV2RGB_NV12_caller(const oclMat &src, oclMat &dst, int bidx)
{
vector<pair<size_t , const void *> > args;
char build_options[50];
sprintf(build_options, "-D DEPTH_%d", src.depth());
//printf("depth:%d,channels:%d,bidx:%d\n",src.depth(),src.oclchannels(),bidx);
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data));
size_t gt[3] = {dst.cols / 2, dst.rows / 2, 1}, lt[3] = {16, 16, 1};
openCLExecuteKernel(src.clCxt, &cvt_color, "YUV2RGBA_NV12", gt, lt, args, -1, -1, build_options);
}
void RGB2YCrCb_caller(const oclMat &src, oclMat &dst, int bidx)
{
vector<pair<size_t , const void *> > args;
int channels = src.oclchannels();
char build_options[50];
sprintf(build_options, "-D DEPTH_%d", src.depth());
//printf("depth:%d,channels:%d,bidx:%d\n",src.depth(),src.oclchannels(),bidx);
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&channels));
args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data));
size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1};
openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2YCrCb", gt, lt, args, -1, -1, build_options);
}
void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
{
Size sz = src.size();
int scn = src.oclchannels(), depth = src.depth(), bidx;
CV_Assert(depth == CV_8U || depth == CV_16U);
CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32F);
switch (code)
{
/*
case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR:
case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA:
case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:
case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555:
case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:
case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA:
*/
case CV_BGR2GRAY:
case CV_BGRA2GRAY:
case CV_RGB2GRAY:
case CV_RGBA2GRAY:
{
CV_Assert(scn == 3 || scn == 4);
bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, 1));
RGB2Gray_caller(src, dst, bidx);
break;
}
switch (code)
{
/*
case CV_BGR5652GRAY: case CV_BGR5552GRAY:
case CV_GRAY2BGR: case CV_GRAY2BGRA:
case CV_GRAY2BGR565: case CV_GRAY2BGR555:
case CV_BGR2YCrCb: case CV_RGB2YCrCb:
case CV_BGR2YUV: case CV_RGB2YUV:
case CV_YCrCb2BGR: case CV_YCrCb2RGB:
case CV_YUV2BGR: case CV_YUV2RGB:
case CV_BGR2XYZ: case CV_RGB2XYZ:
case CV_XYZ2BGR: case CV_XYZ2RGB:
case CV_BGR2HSV: case CV_RGB2HSV: case CV_BGR2HSV_FULL: case CV_RGB2HSV_FULL:
case CV_BGR2HLS: case CV_RGB2HLS: case CV_BGR2HLS_FULL: case CV_RGB2HLS_FULL:
case CV_HSV2BGR: case CV_HSV2RGB: case CV_HSV2BGR_FULL: case CV_HSV2RGB_FULL:
case CV_HLS2BGR: case CV_HLS2RGB: case CV_HLS2BGR_FULL: case CV_HLS2RGB_FULL:
case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR:
case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA:
case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:
case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555:
case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:
case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA:
*/
default:
CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );
}
case CV_BGR2GRAY:
case CV_BGRA2GRAY:
case CV_RGB2GRAY:
case CV_RGBA2GRAY:
{
CV_Assert(scn == 3 || scn == 4);
bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, 1));
RGB2Gray_caller(src, dst, bidx);
break;
}
case CV_GRAY2BGR:
case CV_GRAY2BGRA:
{
CV_Assert(scn == 1);
dcn = code == CV_GRAY2BGRA ? 4 : 3;
dst.create(sz, CV_MAKETYPE(depth, dcn));
Gray2RGB_caller(src, dst);
break;
}
case CV_BGR2YUV:
case CV_RGB2YUV:
{
CV_Assert(scn == 3 || scn == 4);
bidx = code == CV_BGR2YUV ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, 3));
RGB2YUV_caller(src, dst, bidx);
break;
}
case CV_YUV2BGR:
case CV_YUV2RGB:
{
CV_Assert(scn == 3 || scn == 4);
bidx = code == CV_YUV2BGR ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, 3));
YUV2RGB_caller(src, dst, bidx);
break;
}
case CV_YUV2RGB_NV12:
case CV_YUV2BGR_NV12:
case CV_YUV2RGBA_NV12:
case CV_YUV2BGRA_NV12:
{
CV_Assert(scn == 1);
CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U );
dcn = code == CV_YUV2BGRA_NV12 || code == CV_YUV2RGBA_NV12 ? 4 : 3;
bidx = code == CV_YUV2BGRA_NV12 || code == CV_YUV2BGR_NV12 ? 0 : 2;
Size dstSz(sz.width, sz.height * 2 / 3);
dst.create(dstSz, CV_MAKETYPE(depth, dcn));
YUV2RGB_NV12_caller(src, dst, bidx);
break;
}
case CV_BGR2YCrCb:
case CV_RGB2YCrCb:
{
CV_Assert(scn == 3 || scn == 4);
bidx = code == CV_BGR2YCrCb ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, 3));
RGB2YCrCb_caller(src, dst, bidx);
break;
}
case CV_YCrCb2BGR:
case CV_YCrCb2RGB:
{
break;
}
/*
case CV_BGR5652GRAY: case CV_BGR5552GRAY:
case CV_GRAY2BGR565: case CV_GRAY2BGR555:
case CV_BGR2YCrCb: case CV_RGB2YCrCb:
case CV_BGR2XYZ: case CV_RGB2XYZ:
case CV_XYZ2BGR: case CV_XYZ2RGB:
case CV_BGR2HSV: case CV_RGB2HSV: case CV_BGR2HSV_FULL: case CV_RGB2HSV_FULL:
case CV_BGR2HLS: case CV_RGB2HLS: case CV_BGR2HLS_FULL: case CV_RGB2HLS_FULL:
case CV_HSV2BGR: case CV_HSV2RGB: case CV_HSV2BGR_FULL: case CV_HSV2RGB_FULL:
case CV_HLS2BGR: case CV_HLS2RGB: case CV_HLS2BGR_FULL: case CV_HLS2RGB_FULL:
*/
default:
CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );
}
}
}
void cv::ocl::cvtColor(const oclMat &src, oclMat &dst, int code, int dcn)

844
modules/ocl/src/filtering.cpp
View File

File diff suppressed because it is too large Load Diff

1579
modules/ocl/src/haar.cpp
View File

File diff suppressed because it is too large Load Diff

50
modules/ocl/src/imgproc.cpp
View File

@ -171,7 +171,7 @@ namespace cv
typedef void (*gpuThresh_t)(const oclMat &src, oclMat &dst, double thresh, double maxVal, int type);
void threshold_8u(const oclMat &src, oclMat &dst, double thresh, double maxVal, int type)
static void threshold_8u(const oclMat &src, oclMat &dst, double thresh, double maxVal, int type)
{
CV_Assert( (src.cols == dst.cols) && (src.rows == dst.rows) );
Context *clCxt = src.clCxt;
@ -202,7 +202,7 @@ namespace cv
openCLExecuteKernel(clCxt, &imgproc_threshold, kernelName, globalThreads, localThreads, args, src.oclchannels(), src.depth());
}
void threshold_32f(const oclMat &src, oclMat &dst, double thresh, double maxVal, int type)
static void threshold_32f(const oclMat &src, oclMat &dst, double thresh, double maxVal, int type)
{
CV_Assert( (src.cols == dst.cols) && (src.rows == dst.rows) );
Context *clCxt = src.clCxt;
@ -388,7 +388,7 @@ namespace cv
////////////////////////////////////////////////////////////////////////////////////////////
// resize
void resize_gpu( const oclMat &src, oclMat &dst, double fx, double fy, int interpolation)
static void resize_gpu( const oclMat &src, oclMat &dst, double fx, double fy, int interpolation)
{
CV_Assert( (src.channels() == dst.channels()) );
Context *clCxt = src.clCxt;
@ -593,7 +593,7 @@ namespace cv
int dstOffset = dst.offset / dst.elemSize();
int __bordertype[] = {cv::BORDER_CONSTANT, cv::BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101};
const char *borderstr[] = {"BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101"};
int bordertype_index;
size_t bordertype_index;
for(bordertype_index = 0; bordertype_index < sizeof(__bordertype) / sizeof(int); bordertype_index++)
{
if(__bordertype[bordertype_index] == bordertype)
@ -826,9 +826,9 @@ namespace cv
{
#define Sd(y,x) (Sd[y*3+x])
#define Dd(y,x) (Dd[y*3+x])
#define det3(m) (m(0,0)*(m(1,1)*m(2,2) - m(1,2)*m(2,1)) - \
m(0,1)*(m(1,0)*m(2,2) - m(1,2)*m(2,0)) + \
m(0,2)*(m(1,0)*m(2,1) - m(1,1)*m(2,0)))
#define det3(m) (m(0,0)*(m(1,1)*m(2,2) - m(1,2)*m(2,1)) - \
m(0,1)*(m(1,0)*m(2,2) - m(1,2)*m(2,0)) + \
m(0,2)*(m(1,0)*m(2,1) - m(1,1)*m(2,0)))
double *Sd = M;
double *Dd = M;
double d = det3(Sd);
@ -1018,12 +1018,19 @@ namespace cv
int warpInd = (flags & WARP_INVERSE_MAP) >> 4;
F coeffs[2][3];
Mat coeffsMat(2, 3, CV_64F, (void *)coeffs);
double coeffsM[2*3];
Mat coeffsMat(2, 3, CV_64F, (void *)coeffsM);
M.convertTo(coeffsMat, coeffsMat.type());
if(!warpInd)
{
convert_coeffs((F *)(&coeffs[0][0]));
convert_coeffs(coeffsM);
}
for(int i = 0; i < 2; ++i)
for(int j = 0; j < 3; ++j)
coeffs[i][j] = coeffsM[i*3+j];
warpAffine_gpu(src, dst, coeffs, interpolation);
}
@ -1041,13 +1048,19 @@ namespace cv
int warpInd = (flags & WARP_INVERSE_MAP) >> 4;
double coeffs[3][3];
Mat coeffsMat(3, 3, CV_64F, (void *)coeffs);
double coeffsM[3*3];
Mat coeffsMat(3, 3, CV_64F, (void *)coeffsM);
M.convertTo(coeffsMat, coeffsMat.type());
if(!warpInd)
{
invert((double *)(&coeffs[0][0]));
invert(coeffsM);
}
for(int i = 0; i < 3; ++i)
for(int j = 0; j < 3; ++j)
coeffs[i][j] = coeffsM[i*3+j];
warpPerspective_gpu(src, dst, coeffs, interpolation);
}
@ -1144,7 +1157,7 @@ namespace cv
}
/////////////////////// corner //////////////////////////////
void extractCovData(const oclMat &src, oclMat &Dx, oclMat &Dy,
static void extractCovData(const oclMat &src, oclMat &Dx, oclMat &Dy,
int blockSize, int ksize, int borderType)
{
CV_Assert(src.type() == CV_8UC1 || src.type() == CV_32FC1);
@ -1174,7 +1187,7 @@ namespace cv
CV_Assert(Dx.offset == 0 && Dy.offset == 0);
}
void corner_ocl(const char *src_str, string kernelName, int block_size, float k, oclMat &Dx, oclMat &Dy,
static void corner_ocl(const char *src_str, string kernelName, int block_size, float k, oclMat &Dx, oclMat &Dy,
oclMat &dst, int border_type)
{
char borderType[30];
@ -1258,7 +1271,7 @@ namespace cv
corner_ocl(imgproc_calcMinEigenVal, "calcMinEigenVal", blockSize, 0, Dx, Dy, dst, borderType);
}
/////////////////////////////////// MeanShiftfiltering ///////////////////////////////////////////////
void meanShiftFiltering_gpu(const oclMat &src, oclMat dst, int sp, int sr, int maxIter, float eps)
static void meanShiftFiltering_gpu(const oclMat &src, oclMat dst, int sp, int sr, int maxIter, float eps)
{
CV_Assert( (src.cols == dst.cols) && (src.rows == dst.rows) );
CV_Assert( !(dst.step & 0x3) );
@ -1321,7 +1334,7 @@ namespace cv
}
void meanShiftProc_gpu(const oclMat &src, oclMat dstr, oclMat dstsp, int sp, int sr, int maxIter, float eps)
static void meanShiftProc_gpu(const oclMat &src, oclMat dstr, oclMat dstsp, int sp, int sr, int maxIter, float eps)
{
//sanity checks
CV_Assert( (src.cols == dstr.cols) && (src.rows == dstr.rows) &&
@ -1398,7 +1411,7 @@ namespace cv
const int HISTOGRAM256_BIN_COUNT = 256;
}
///////////////////////////////calcHist/////////////////////////////////////////////////////////////////
void calc_sub_hist(const oclMat &mat_src, const oclMat &mat_sub_hist)
static void calc_sub_hist(const oclMat &mat_src, const oclMat &mat_sub_hist)
{
using namespace histograms;
@ -1477,7 +1490,7 @@ namespace cv
openCLExecuteKernel(clCxt, &imgproc_histogram, kernelName, globalThreads, localThreads, args, -1, depth);
}
}
void merge_sub_hist(const oclMat &sub_hist, oclMat &mat_hist)
static void merge_sub_hist(const oclMat &sub_hist, oclMat &mat_hist)
{
using namespace histograms;
@ -1535,7 +1548,6 @@ namespace cv
{
int cn = src.channels();
int i, j, maxk, radius;
Size size = src.size();
CV_Assert( (src.channels() == 1 || src.channels() == 3) &&
src.type() == dst.type() && src.size() == dst.size() &&
@ -1632,7 +1644,7 @@ inline int divUp(int total, int grain)
{
return (total + grain - 1) / grain;
}
void convolve_run(const oclMat &src, const oclMat &temp1, oclMat &dst, string kernelName, const char **kernelString)
static void convolve_run(const oclMat &src, const oclMat &temp1, oclMat &dst, string kernelName, const char **kernelString)
{
CV_Assert(src.depth() == CV_32FC1);
CV_Assert(temp1.depth() == CV_32F);

14
modules/ocl/src/initialization.cpp
View File

@ -237,23 +237,13 @@ namespace cv
int getDevice(std::vector<Info> &oclinfo, int devicetype)
{
cl_device_type _devicetype;
switch(devicetype)
{
case CVCL_DEVICE_TYPE_DEFAULT:
_devicetype = CL_DEVICE_TYPE_DEFAULT;
break;
case CVCL_DEVICE_TYPE_CPU:
_devicetype = CL_DEVICE_TYPE_CPU;
break;
case CVCL_DEVICE_TYPE_GPU:
_devicetype = CL_DEVICE_TYPE_GPU;
break;
case CVCL_DEVICE_TYPE_ACCELERATOR:
_devicetype = CL_DEVICE_TYPE_ACCELERATOR;
break;
case CVCL_DEVICE_TYPE_ALL:
_devicetype = CL_DEVICE_TYPE_ALL;
break;
default:
CV_Error(CV_GpuApiCallError, "Unkown device type");
@ -336,7 +326,7 @@ namespace cv
size_t extends_size;
openCLSafeCall(clGetDeviceInfo(oclinfo.impl->devices[devnum], CL_DEVICE_EXTENSIONS,
EXT_LEN, (void *)extends_set, &extends_size));
CV_Assert(extends_size < EXT_LEN);
CV_Assert(extends_size < (size_t)EXT_LEN);
extends_set[EXT_LEN - 1] = 0;
memset(oclinfo.impl->extra_options, 0, 512);
oclinfo.impl->double_support = 0;
@ -592,7 +582,7 @@ namespace cv
size_t binarySize = ftell(fp);
fseek(fp, 0, SEEK_SET);
char *binary = new char[binarySize];
fread(binary, binarySize, 1, fp);
CV_Assert(1 == fread(binary, binarySize, 1, fp));
fclose(fp);
cl_int status = 0;
program = clCreateProgramWithBinary(clCxt->impl->clContext,

2
modules/ocl/src/interpolate_frames.cpp
View File

@ -282,7 +282,7 @@ void interpolate::bindImgTex(const oclMat &img, cl_mem &texture)
openCLFree(texture);
}
#if CL_VERSION_1_2
#ifdef CL_VERSION_1_2
cl_image_desc desc;
desc.image_type = CL_MEM_OBJECT_IMAGE2D;
desc.image_width = img.step / img.elemSize();

20
modules/ocl/src/kernels/arithm_mul.cl
View File

@ -16,6 +16,7 @@
//
// @Authors
// Jia Haipeng, jiahaipeng95@gmail.com
// Dachuan Zhao, dachuan@multicorewareinc.com
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
@ -260,3 +261,22 @@ __kernel void arithm_mul_D6 (__global double *src1, int src1_step, int src1_offs
}
}
#endif
__kernel void arithm_muls_D5 (__global float *src1, int src1_step, int src1_offset,
__global float *dst, int dst_step, int dst_offset,
int rows, int cols, int dst_step1, float scalar)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < cols && y < rows)
{
int src1_index = mad24(y, src1_step, (x << 2) + src1_offset);
int dst_index = mad24(y, dst_step, (x << 2) + dst_offset);
float data1 = *((__global float *)((__global char *)src1 + src1_index));
float tmp = data1 * scalar;
*((__global float *)((__global char *)dst + dst_index)) = tmp;
}
}

219
modules/ocl/src/kernels/cvt_color.cl
View File

@ -16,6 +16,7 @@
//
// @Authors
// Jia Haipeng, jiahaipeng95@gmail.com
// Peng Xiao, pengxiao@multicorewareinc.com
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
@ -48,13 +49,33 @@
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#endif
#define DATA_TYPE UNDEFINED
#if defined (DEPTH_0)
#undef DATA_TYPE
#define DATA_TYPE uchar
#define MAX_NUM 255
#define HALF_MAX 128
#define SAT_CAST(num) convert_uchar_sat(num)
#endif
#if defined (DEPTH_2)
#undef DATA_TYPE
#define DATA_TYPE ushort
#define MAX_NUM 65535
#define HALF_MAX 32768
#define SAT_CAST(num) convert_ushort_sat(num)
#endif
#if defined (DEPTH_5)
#undef DATA_TYPE
#define DATA_TYPE float
#define MAX_NUM 1.0f
#define HALF_MAX 0.5f
#define SAT_CAST(num) (num)
#endif
#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n))
enum
{
@ -65,6 +86,7 @@ enum
B2Y = 1868,
BLOCK_SIZE = 256
};
///////////////////////////////////// RGB <-> GRAY //////////////////////////////////////
__kernel void RGB2Gray(int cols,int rows,int src_step,int dst_step,int channels,
int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst)
@ -72,10 +94,203 @@ __kernel void RGB2Gray(int cols,int rows,int src_step,int dst_step,int channels,
const int x = get_global_id(0);
const int y = get_global_id(1);
src_step /= sizeof(DATA_TYPE);
dst_step /= sizeof(DATA_TYPE);
if (y < rows && x < cols)
{
int src_idx = y * src_step + x * channels * sizeof(DATA_TYPE);
int dst_idx = y * dst_step + x * sizeof(DATA_TYPE);
int src_idx = y * src_step + x * channels;
int dst_idx = y * dst_step + x;
#if defined (DEPTH_5)
dst[dst_idx] = src[src_idx + bidx] * 0.114f + src[src_idx + 1] * 0.587f + src[src_idx + (bidx^2)] * 0.299f;
#else
dst[dst_idx] = (DATA_TYPE)CV_DESCALE((src[src_idx + bidx] * B2Y + src[src_idx + 1] * G2Y + src[src_idx + (bidx^2)] * R2Y), yuv_shift);
#endif
}
}
__kernel void Gray2RGB(int cols,int rows,int src_step,int dst_step,
__global const DATA_TYPE* src, __global DATA_TYPE* dst)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
src_step /= sizeof(DATA_TYPE);
dst_step /= sizeof(DATA_TYPE);
if (y < rows && x < cols)
{
int src_idx = y * src_step + x;
int dst_idx = y * dst_step + x * 4;
DATA_TYPE val = src[src_idx];
dst[dst_idx++] = val;
dst[dst_idx++] = val;
dst[dst_idx++] = val;
dst[dst_idx] = MAX_NUM;
}
}
///////////////////////////////////// RGB <-> YUV //////////////////////////////////////
__constant float c_RGB2YUVCoeffs_f[5] = { 0.114f, 0.587f, 0.299f, 0.492f, 0.877f };
__constant int c_RGB2YUVCoeffs_i[5] = { B2Y, G2Y, R2Y, 8061, 14369 };
__kernel void RGB2YUV(int cols,int rows,int src_step,int dst_step,int channels,
int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
src_step /= sizeof(DATA_TYPE);
dst_step /= sizeof(DATA_TYPE);
if (y < rows && x < cols)
{
int src_idx = y * src_step + x * channels;
int dst_idx = y * dst_step + x * channels;
dst += dst_idx;
const DATA_TYPE rgb[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]};
#if defined (DEPTH_5)
__constant float * coeffs = c_RGB2YUVCoeffs_f;
const DATA_TYPE Y = rgb[0] * coeffs[bidx] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx^2];
const DATA_TYPE Cr = (rgb[bidx] - Y) * coeffs[3] + HALF_MAX;
const DATA_TYPE Cb = (rgb[bidx^2] - Y) * coeffs[4] + HALF_MAX;
#else
__constant int * coeffs = c_RGB2YUVCoeffs_i;
const int delta = HALF_MAX * (1 << yuv_shift);
const int Y = CV_DESCALE(rgb[0] * coeffs[bidx] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx^2], yuv_shift);
const int Cr = CV_DESCALE((rgb[bidx] - Y) * coeffs[3] + delta, yuv_shift);
const int Cb = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[4] + delta, yuv_shift);
#endif
dst[0] = SAT_CAST( Y );
dst[1] = SAT_CAST( Cr );
dst[2] = SAT_CAST( Cb );
}
}
__constant float c_YUV2RGBCoeffs_f[5] = { 2.032f, -0.395f, -0.581f, 1.140f };
__constant int c_YUV2RGBCoeffs_i[5] = { 33292, -6472, -9519, 18678 };
__kernel void YUV2RGB(int cols,int rows,int src_step,int dst_step,int channels,
int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
src_step /= sizeof(DATA_TYPE);
dst_step /= sizeof(DATA_TYPE);
if (y < rows && x < cols)
{
int src_idx = y * src_step + x * channels;
int dst_idx = y * dst_step + x * channels;
dst += dst_idx;
const DATA_TYPE yuv[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]};
#if defined (DEPTH_5)
__constant float * coeffs = c_YUV2RGBCoeffs_f;
const float b = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[3];
const float g = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1];
const float r = yuv[0] + (yuv[1] - HALF_MAX) * coeffs[0];
#else
__constant int * coeffs = c_YUV2RGBCoeffs_i;
const int b = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[3], yuv_shift);
const int g = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1], yuv_shift);
const int r = yuv[0] + CV_DESCALE((yuv[1] - HALF_MAX) * coeffs[0], yuv_shift);
#endif
dst[bidx^2] = SAT_CAST( b );
dst[1] = SAT_CAST( g );
dst[bidx] = SAT_CAST( r );
}
}
__constant int ITUR_BT_601_CY = 1220542;
__constant int ITUR_BT_601_CUB = 2116026;
__constant int ITUR_BT_601_CUG = -409993;
__constant int ITUR_BT_601_CVG = -852492;
__constant int ITUR_BT_601_CVR = 1673527;
__constant int ITUR_BT_601_SHIFT = 20;
__kernel void YUV2RGBA_NV12(int cols,int rows,int src_step,int dst_step,
int bidx, int width, int height, __global const uchar* src, __global uchar* dst)
{
const int x = get_global_id(0); // max_x = width / 2
const int y = get_global_id(1); // max_y = height/ 2
if (y < height / 2 && x < width / 2 )
{
__global const uchar* ysrc = src + (y << 1) * src_step + (x << 1);
__global const uchar* usrc = src + (height + y) * src_step + (x << 1);
__global uchar* dst1 = dst + (y << 1) * dst_step + (x << 3);
__global uchar* dst2 = dst + ((y << 1) + 1) * dst_step + (x << 3);
int Y1 = ysrc[0];
int Y2 = ysrc[1];
int Y3 = ysrc[src_step];
int Y4 = ysrc[src_step + 1];
int U = usrc[0] - 128;
int V = usrc[1] - 128;
int ruv = (1 << (ITUR_BT_601_SHIFT - 1)) + ITUR_BT_601_CVR * V;
int guv = (1 << (ITUR_BT_601_SHIFT - 1)) + ITUR_BT_601_CVG * V + ITUR_BT_601_CUG * U;
int buv = (1 << (ITUR_BT_601_SHIFT - 1)) + ITUR_BT_601_CUB * U;
Y1 = max(0, Y1 - 16) * ITUR_BT_601_CY;
dst1[2 - bidx] = convert_uchar_sat((Y1 + ruv) >> ITUR_BT_601_SHIFT);
dst1[1] = convert_uchar_sat((Y1 + guv) >> ITUR_BT_601_SHIFT);
dst1[bidx] = convert_uchar_sat((Y1 + buv) >> ITUR_BT_601_SHIFT);
dst1[3] = 255;
Y2 = max(0, Y2 - 16) * ITUR_BT_601_CY;
dst1[6 - bidx] = convert_uchar_sat((Y2 + ruv) >> ITUR_BT_601_SHIFT);
dst1[5] = convert_uchar_sat((Y2 + guv) >> ITUR_BT_601_SHIFT);
dst1[4 + bidx] = convert_uchar_sat((Y2 + buv) >> ITUR_BT_601_SHIFT);
dst1[7] = 255;
Y3 = max(0, Y3 - 16) * ITUR_BT_601_CY;
dst2[2 - bidx] = convert_uchar_sat((Y3 + ruv) >> ITUR_BT_601_SHIFT);
dst2[1] = convert_uchar_sat((Y3 + guv) >> ITUR_BT_601_SHIFT);
dst2[bidx] = convert_uchar_sat((Y3 + buv) >> ITUR_BT_601_SHIFT);
dst2[3] = 255;
Y4 = max(0, Y4 - 16) * ITUR_BT_601_CY;
dst2[6 - bidx] = convert_uchar_sat((Y4 + ruv) >> ITUR_BT_601_SHIFT);
dst2[5] = convert_uchar_sat((Y4 + guv) >> ITUR_BT_601_SHIFT);
dst2[4 + bidx] = convert_uchar_sat((Y4 + buv) >> ITUR_BT_601_SHIFT);
dst2[7] = 255;
}
}
///////////////////////////////////// RGB <-> YUV //////////////////////////////////////
__constant float c_RGB2YCrCbCoeffs_f[5] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f};
__constant int c_RGB2YCrCbCoeffs_i[5] = {R2Y, G2Y, B2Y, 11682, 9241};
__kernel void RGB2YCrCb(int cols,int rows,int src_step,int dst_step,int channels,
int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
src_step /= sizeof(DATA_TYPE);
dst_step /= sizeof(DATA_TYPE);
if (y < rows && x < cols)
{
int src_idx = y * src_step + x * channels;
int dst_idx = y * dst_step + x * channels;
dst += dst_idx;
const DATA_TYPE rgb[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]};
#if defined (DEPTH_5)
__constant float * coeffs = c_RGB2YCrCbCoeffs_f;
const DATA_TYPE Y = rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx];
const DATA_TYPE Cr = (rgb[bidx^2] - Y) * coeffs[3] + HALF_MAX;
const DATA_TYPE Cb = (rgb[bidx] - Y) * coeffs[4] + HALF_MAX;
#else
__constant int * coeffs = c_RGB2YCrCbCoeffs_i;
const int delta = HALF_MAX * (1 << yuv_shift);
const int Y = CV_DESCALE(rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx], yuv_shift);
const int Cr = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[3] + delta, yuv_shift);
const int Cb = CV_DESCALE((rgb[bidx] - Y) * coeffs[4] + delta, yuv_shift);
#endif
dst[0] = SAT_CAST( Y );
dst[1] = SAT_CAST( Cr );
dst[2] = SAT_CAST( Cb );
}
}

460
modules/ocl/src/kernels/haarobjectdetect_scaled2.cl
View File

@ -44,75 +44,75 @@
//M*/
// Enter your kernel in this window
#pragma OPENCL EXTENSION cl_amd_printf:enable
//#pragma OPENCL EXTENSION cl_amd_printf:enable
#define CV_HAAR_FEATURE_MAX 3
typedef int sumtype;
typedef float sqsumtype;
typedef struct __attribute__((aligned (128))) GpuHidHaarFeature
typedef struct __attribute__((aligned(128))) GpuHidHaarFeature
{
struct __attribute__((aligned (32)))
{
int p0 __attribute__((aligned (4)));
int p1 __attribute__((aligned (4)));
int p2 __attribute__((aligned (4)));
int p3 __attribute__((aligned (4)));
float weight __attribute__((aligned (4)));
}
rect[CV_HAAR_FEATURE_MAX] __attribute__((aligned (32)));
struct __attribute__((aligned(32)))
{
int p0 __attribute__((aligned(4)));
int p1 __attribute__((aligned(4)));
int p2 __attribute__((aligned(4)));
int p3 __attribute__((aligned(4)));
float weight __attribute__((aligned(4)));
}
rect[CV_HAAR_FEATURE_MAX] __attribute__((aligned(32)));
}
GpuHidHaarFeature;
typedef struct __attribute__((aligned (128) )) GpuHidHaarTreeNode
typedef struct __attribute__((aligned(128))) GpuHidHaarTreeNode
{
int p[CV_HAAR_FEATURE_MAX][4] __attribute__((aligned (64)));
int p[CV_HAAR_FEATURE_MAX][4] __attribute__((aligned(64)));
float weight[CV_HAAR_FEATURE_MAX] /*__attribute__((aligned (16)))*/;
float threshold /*__attribute__((aligned (4)))*/;
float alpha[2] __attribute__((aligned (8)));
int left __attribute__((aligned (4)));
int right __attribute__((aligned (4)));
float alpha[2] __attribute__((aligned(8)));
int left __attribute__((aligned(4)));
int right __attribute__((aligned(4)));
}
GpuHidHaarTreeNode;
typedef struct __attribute__((aligned (32))) GpuHidHaarClassifier
typedef struct __attribute__((aligned(32))) GpuHidHaarClassifier
{
int count __attribute__((aligned (4)));
GpuHidHaarTreeNode* node __attribute__((aligned (8)));
float* alpha __attribute__((aligned (8)));
int count __attribute__((aligned(4)));
GpuHidHaarTreeNode *node __attribute__((aligned(8)));
float *alpha __attribute__((aligned(8)));
}
GpuHidHaarClassifier;
typedef struct __attribute__((aligned (64))) GpuHidHaarStageClassifier
typedef struct __attribute__((aligned(64))) GpuHidHaarStageClassifier
{
int count __attribute__((aligned (4)));
float threshold __attribute__((aligned (4)));
int two_rects __attribute__((aligned (4)));
int reserved0 __attribute__((aligned (8)));
int reserved1 __attribute__((aligned (8)));
int reserved2 __attribute__((aligned (8)));
int reserved3 __attribute__((aligned (8)));
int count __attribute__((aligned(4)));
float threshold __attribute__((aligned(4)));
int two_rects __attribute__((aligned(4)));
int reserved0 __attribute__((aligned(8)));
int reserved1 __attribute__((aligned(8)));
int reserved2 __attribute__((aligned(8)));
int reserved3 __attribute__((aligned(8)));
}
GpuHidHaarStageClassifier;
typedef struct __attribute__((aligned (64))) GpuHidHaarClassifierCascade
typedef struct __attribute__((aligned(64))) GpuHidHaarClassifierCascade
{
int count __attribute__((aligned (4)));
int is_stump_based __attribute__((aligned (4)));
int has_tilted_features __attribute__((aligned (4)));
int is_tree __attribute__((aligned (4)));
int pq0 __attribute__((aligned (4)));
int pq1 __attribute__((aligned (4)));
int pq2 __attribute__((aligned (4)));
int pq3 __attribute__((aligned (4)));
int p0 __attribute__((aligned (4)));
int p1 __attribute__((aligned (4)));
int p2 __attribute__((aligned (4)));
int p3 __attribute__((aligned (4)));
float inv_window_area __attribute__((aligned (4)));
}GpuHidHaarClassifierCascade;
int count __attribute__((aligned(4)));
int is_stump_based __attribute__((aligned(4)));
int has_tilted_features __attribute__((aligned(4)));
int is_tree __attribute__((aligned(4)));
int pq0 __attribute__((aligned(4)));
int pq1 __attribute__((aligned(4)));
int pq2 __attribute__((aligned(4)));
int pq3 __attribute__((aligned(4)));
int p0 __attribute__((aligned(4)));
int p1 __attribute__((aligned(4)));
int p2 __attribute__((aligned(4)));
int p3 __attribute__((aligned(4)));
float inv_window_area __attribute__((aligned(4)));
} GpuHidHaarClassifierCascade;
__kernel void gpuRunHaarClassifierCascade_scaled2(
global GpuHidHaarStageClassifier * stagecascadeptr,
global int4 * info,
global GpuHidHaarTreeNode * nodeptr,
global const int * restrict sum,
global const float * restrict sqsum,
global int4 * candidate,
global GpuHidHaarStageClassifier *stagecascadeptr,
global int4 *info,
global GpuHidHaarTreeNode *nodeptr,
global const int *restrict sum,
global const float *restrict sqsum,
global int4 *candidate,
const int step,
const int loopcount,
const int start_stage,
@ -120,215 +120,167 @@ __kernel void gpuRunHaarClassifierCascade_scaled2(
const int end_stage,
const int startnode,
const int splitnode,
global int4 * p,
//const int4 * pq,
global float * correction,
const int nodecount)
global int4 *p,
//const int4 * pq,
global float *correction,
const int nodecount)
{
int grpszx = get_local_size(0);
int grpszy = get_local_size(1);
int grpnumx = get_num_groups(0);
int grpidx=get_group_id(0);
int lclidx = get_local_id(0);
int lclidy = get_local_id(1);
int lcl_sz = mul24(grpszx,grpszy);
int lcl_id = mad24(lclidy,grpszx,lclidx);
__local int lclshare[1024];
__local int* glboutindex=lclshare+0;
__local int* lclcount=glboutindex+1;
__local int* lcloutindex=lclcount+1;
__local float* partialsum=(__local float*)(lcloutindex+(lcl_sz<<1));
glboutindex[0]=0;
int outputoff = mul24(grpidx,256);
candidate[outputoff+(lcl_id<<2)] = (int4)0;
candidate[outputoff+(lcl_id<<2)+1] = (int4)0;
candidate[outputoff+(lcl_id<<2)+2] = (int4)0;
candidate[outputoff+(lcl_id<<2)+3] = (int4)0;
for(int scalei = 0; scalei <loopcount; scalei++)
{
int4 scaleinfo1;
scaleinfo1 = info[scalei];
int width = (scaleinfo1.x & 0xffff0000) >> 16;
int height = scaleinfo1.x & 0xffff;
int grpnumperline =(scaleinfo1.y & 0xffff0000) >> 16;
int totalgrp = scaleinfo1.y & 0xffff;
float factor = as_float(scaleinfo1.w);
float correction_t=correction[scalei];
int ystep=(int)(max(2.0f,factor)+0.5f);
for(int grploop=get_group_id(0);grploop<totalgrp;grploop+=grpnumx){
int4 cascadeinfo=p[scalei];
int grpidy = grploop / grpnumperline;
int grpidx = grploop - mul24(grpidy, grpnumperline);
int ix = mad24(grpidx,grpszx,lclidx);
int iy = mad24(grpidy,grpszy,lclidy);
int x=ix*ystep;
int y=iy*ystep;
lcloutindex[lcl_id]=0;
lclcount[0]=0;
int result=1,nodecounter;
float mean,variance_norm_factor;
//if((ix < width) && (iy < height))
{
const int p_offset = mad24(y, step, x);
cascadeinfo.x +=p_offset;
cascadeinfo.z +=p_offset;
mean = (sum[mad24(cascadeinfo.y,step,cascadeinfo.x)] - sum[mad24(cascadeinfo.y,step,cascadeinfo.z)] -
sum[mad24(cascadeinfo.w,step,cascadeinfo.x)] + sum[mad24(cascadeinfo.w,step,cascadeinfo.z)])
*correction_t;
variance_norm_factor =sqsum[mad24(cascadeinfo.y,step, cascadeinfo.x)] - sqsum[mad24(cascadeinfo.y, step, cascadeinfo.z)] -
sqsum[mad24(cascadeinfo.w, step, cascadeinfo.x)] + sqsum[mad24(cascadeinfo.w, step, cascadeinfo.z)];
variance_norm_factor = variance_norm_factor * correction_t - mean * mean;
variance_norm_factor = variance_norm_factor >=0.f ? sqrt(variance_norm_factor) : 1.f;
result = 1;
nodecounter = startnode+nodecount*scalei;
for(int stageloop = start_stage; stageloop < split_stage&&result; stageloop++ )
{
float stage_sum = 0.f;
int4 stageinfo = *(global int4*)(stagecascadeptr+stageloop);
float stagethreshold = as_float(stageinfo.y);
for(int nodeloop = 0; nodeloop < stageinfo.x; nodeloop++ )
{
__global GpuHidHaarTreeNode* currentnodeptr = (nodeptr + nodecounter);
int4 info1 = *(__global int4*)(&(currentnodeptr->p[0][0]));
int4 info2 = *(__global int4*)(&(currentnodeptr->p[1][0]));
int4 info3 = *(__global int4*)(&(currentnodeptr->p[2][0]));
float4 w = *(__global float4*)(&(currentnodeptr->weight[0]));
float2 alpha2 = *(__global float2*)(&(currentnodeptr->alpha[0]));
float nodethreshold = w.w * variance_norm_factor;
info1.x +=p_offset;
info1.z +=p_offset;
info2.x +=p_offset;
info2.z +=p_offset;
float classsum = (sum[mad24(info1.y,step,info1.x)] - sum[mad24(info1.y,step,info1.z)] -
sum[mad24(info1.w,step,info1.x)] + sum[mad24(info1.w,step,info1.z)]) * w.x;
classsum += (sum[mad24(info2.y,step,info2.x)] - sum[mad24(info2.y,step,info2.z)] -
sum[mad24(info2.w,step,info2.x)] + sum[mad24(info2.w,step,info2.z)]) * w.y;
info3.x +=p_offset;
info3.z +=p_offset;
classsum += (sum[mad24(info3.y,step,info3.x)] - sum[mad24(info3.y,step,info3.z)] -
sum[mad24(info3.w,step,info3.x)] + sum[mad24(info3.w,step,info3.z)]) * w.z;
stage_sum += classsum >= nodethreshold ? alpha2.y : alpha2.x;
nodecounter++;
}
result=(stage_sum>=stagethreshold);
}
if(result&&(ix<width)&&(iy<height))
{
int queueindex=atomic_inc(lclcount);
lcloutindex[queueindex<<1]=(y<<16)|x;
lcloutindex[(queueindex<<1)+1]=as_int(variance_norm_factor);
}
barrier(CLK_LOCAL_MEM_FENCE);
int queuecount=lclcount[0];
nodecounter=splitnode+nodecount*scalei;
for(int stageloop=split_stage;stageloop<end_stage&&queuecount>0;stageloop++)
{
lclcount[0]=0;
barrier(CLK_LOCAL_MEM_FENCE);
int2 stageinfo=*(global int2*)(stagecascadeptr+stageloop);
float stagethreshold=as_float(stageinfo.y);
int perfscale=queuecount>4?3:2;
int queuecount_loop=(queuecount+(1<<perfscale)-1)>>perfscale;
int lcl_compute_win=lcl_sz>>perfscale;
int lcl_compute_win_id=(lcl_id>>(6-perfscale));
int lcl_loops=(stageinfo.x+lcl_compute_win-1)>>(6-perfscale);
int lcl_compute_id=lcl_id-(lcl_compute_win_id<<(6-perfscale));
for(int queueloop=0;queueloop<queuecount_loop&&lcl_compute_win_id<queuecount;queueloop++)
{
float stage_sum=0.f;
int temp_coord=lcloutindex[lcl_compute_win_id<<1];
float variance_norm_factor=as_float(lcloutindex[(lcl_compute_win_id<<1)+1]);
int queue_offset=mad24(((temp_coord&(int)0xffff0000)>>16),step,temp_coord&0xffff);
int tempnodecounter=lcl_compute_id;
float part_sum=0.f;
for(int lcl_loop=0;lcl_loop<lcl_loops&&tempnodecounter<stageinfo.x;lcl_loop++)
{
__global GpuHidHaarTreeNode* currentnodeptr = (nodeptr + nodecounter + tempnodecounter);
int4 info1 = *(__global int4*)(&(currentnodeptr->p[0][0]));
int4 info2 = *(__global int4*)(&(currentnodeptr->p[1][0]));
int4 info3 = *(__global int4*)(&(currentnodeptr->p[2][0]));
float4 w = *(__global float4*)(&(currentnodeptr->weight[0]));
float2 alpha2 = *(__global float2*)(&(currentnodeptr->alpha[0]));
float nodethreshold = w.w * variance_norm_factor;
info1.x +=queue_offset;
info1.z +=queue_offset;
info2.x +=queue_offset;
info2.z +=queue_offset;
float classsum = (sum[mad24(info1.y,step,info1.x)] - sum[mad24(info1.y,step,info1.z)] -
sum[mad24(info1.w,step,info1.x)] + sum[mad24(info1.w,step,info1.z)]) * w.x;
classsum += (sum[mad24(info2.y,step,info2.x)] - sum[mad24(info2.y,step,info2.z)] -
sum[mad24(info2.w,step,info2.x)] + sum[mad24(info2.w,step,info2.z)]) * w.y;
int grpszx = get_local_size(0);
int grpszy = get_local_size(1);
int grpnumx = get_num_groups(0);
int grpidx = get_group_id(0);
int lclidx = get_local_id(0);
int lclidy = get_local_id(1);
int lcl_sz = mul24(grpszx, grpszy);
int lcl_id = mad24(lclidy, grpszx, lclidx);
__local int lclshare[1024];
__local int *glboutindex = lclshare + 0;
__local int *lclcount = glboutindex + 1;
__local int *lcloutindex = lclcount + 1;
__local float *partialsum = (__local float *)(lcloutindex + (lcl_sz << 1));
glboutindex[0] = 0;
int outputoff = mul24(grpidx, 256);
candidate[outputoff + (lcl_id << 2)] = (int4)0;
candidate[outputoff + (lcl_id << 2) + 1] = (int4)0;
candidate[outputoff + (lcl_id << 2) + 2] = (int4)0;
candidate[outputoff + (lcl_id << 2) + 3] = (int4)0;
info3.x +=queue_offset;
info3.z +=queue_offset;
classsum += (sum[mad24(info3.y,step,info3.x)] - sum[mad24(info3.y,step,info3.z)] -
sum[mad24(info3.w,step,info3.x)] + sum[mad24(info3.w,step,info3.z)]) * w.z;
part_sum += classsum >= nodethreshold ? alpha2.y : alpha2.x;
tempnodecounter+=lcl_compute_win;
}
partialsum[lcl_id]=part_sum;
barrier(CLK_LOCAL_MEM_FENCE);
for(int i=0;i<lcl_compute_win&&(lcl_compute_id==0);i++)
{
stage_sum+=partialsum[lcl_id+i];
}
if(stage_sum>=stagethreshold&&(lcl_compute_id==0))
{
int queueindex=atomic_inc(lclcount);
lcloutindex[queueindex<<1]=temp_coord;
lcloutindex[(queueindex<<1)+1]=as_int(variance_norm_factor);
}
lcl_compute_win_id+=(1<<perfscale);
barrier(CLK_LOCAL_MEM_FENCE);
}
queuecount=lclcount[0];
nodecounter+=stageinfo.x;
}
if(lcl_id<queuecount)
{
int temp=lcloutindex[lcl_id<<1];
int x=temp&0xffff;
int y=(temp&(int)0xffff0000)>>16;
temp=glboutindex[0];
int4 candidate_result;
candidate_result.zw=(int2)convert_int_rtn(factor*20.f);
candidate_result.x=x;
candidate_result.y=y;
atomic_inc(glboutindex);
candidate[outputoff+temp+lcl_id]=candidate_result;
}
barrier(CLK_LOCAL_MEM_FENCE);
}
for (int scalei = 0; scalei < loopcount; scalei++)
{
int4 scaleinfo1;
scaleinfo1 = info[scalei];
int width = (scaleinfo1.x & 0xffff0000) >> 16;
int height = scaleinfo1.x & 0xffff;
int grpnumperline = (scaleinfo1.y & 0xffff0000) >> 16;
int totalgrp = scaleinfo1.y & 0xffff;
float factor = as_float(scaleinfo1.w);
float correction_t = correction[scalei];
int ystep = (int)(max(2.0f, factor) + 0.5f);
for (int grploop = get_group_id(0); grploop < totalgrp; grploop += grpnumx)
{
int4 cascadeinfo = p[scalei];
int grpidy = grploop / grpnumperline;
int grpidx = grploop - mul24(grpidy, grpnumperline);
int ix = mad24(grpidx, grpszx, lclidx);
int iy = mad24(grpidy, grpszy, lclidy);
int x = ix * ystep;
int y = iy * ystep;
lcloutindex[lcl_id] = 0;
lclcount[0] = 0;
int result = 1, nodecounter;
float mean, variance_norm_factor;
//if((ix < width) && (iy < height))
{
const int p_offset = mad24(y, step, x);
cascadeinfo.x += p_offset;
cascadeinfo.z += p_offset;
mean = (sum[mad24(cascadeinfo.y, step, cascadeinfo.x)] - sum[mad24(cascadeinfo.y, step, cascadeinfo.z)] -
sum[mad24(cascadeinfo.w, step, cascadeinfo.x)] + sum[mad24(cascadeinfo.w, step, cascadeinfo.z)])
* correction_t;
variance_norm_factor = sqsum[mad24(cascadeinfo.y, step, cascadeinfo.x)] - sqsum[mad24(cascadeinfo.y, step, cascadeinfo.z)] -
sqsum[mad24(cascadeinfo.w, step, cascadeinfo.x)] + sqsum[mad24(cascadeinfo.w, step, cascadeinfo.z)];
variance_norm_factor = variance_norm_factor * correction_t - mean * mean;
variance_norm_factor = variance_norm_factor >= 0.f ? sqrt(variance_norm_factor) : 1.f;
result = 1;
nodecounter = startnode + nodecount * scalei;
for (int stageloop = start_stage; stageloop < end_stage && result; stageloop++)
{
float stage_sum = 0.f;
int4 stageinfo = *(global int4 *)(stagecascadeptr + stageloop);
float stagethreshold = as_float(stageinfo.y);
for (int nodeloop = 0; nodeloop < stageinfo.x; nodeloop++)
{
__global GpuHidHaarTreeNode *currentnodeptr = (nodeptr + nodecounter);
int4 info1 = *(__global int4 *)(&(currentnodeptr->p[0][0]));
int4 info2 = *(__global int4 *)(&(currentnodeptr->p[1][0]));
int4 info3 = *(__global int4 *)(&(currentnodeptr->p[2][0]));
float4 w = *(__global float4 *)(&(currentnodeptr->weight[0]));
float2 alpha2 = *(__global float2 *)(&(currentnodeptr->alpha[0]));
float nodethreshold = w.w * variance_norm_factor;
info1.x += p_offset;
info1.z += p_offset;
info2.x += p_offset;
info2.z += p_offset;
float classsum = (sum[mad24(info1.y, step, info1.x)] - sum[mad24(info1.y, step, info1.z)] -
sum[mad24(info1.w, step, info1.x)] + sum[mad24(info1.w, step, info1.z)]) * w.x;
classsum += (sum[mad24(info2.y, step, info2.x)] - sum[mad24(info2.y, step, info2.z)] -
sum[mad24(info2.w, step, info2.x)] + sum[mad24(info2.w, step, info2.z)]) * w.y;
info3.x += p_offset;
info3.z += p_offset;
classsum += (sum[mad24(info3.y, step, info3.x)] - sum[mad24(info3.y, step, info3.z)] -
sum[mad24(info3.w, step, info3.x)] + sum[mad24(info3.w, step, info3.z)]) * w.z;
stage_sum += classsum >= nodethreshold ? alpha2.y : alpha2.x;
nodecounter++;
}
result = (stage_sum >= stagethreshold);
}
if (result && (ix < width) && (iy < height))
{
int queueindex = atomic_inc(lclcount);
lcloutindex[queueindex << 1] = (y << 16) | x;
lcloutindex[(queueindex << 1) + 1] = as_int(variance_norm_factor);
}
barrier(CLK_LOCAL_MEM_FENCE);
int queuecount = lclcount[0];
nodecounter = splitnode + nodecount * scalei;
if (lcl_id < queuecount)
{
int temp = lcloutindex[lcl_id << 1];
int x = temp & 0xffff;
int y = (temp & (int)0xffff0000) >> 16;
temp = glboutindex[0];
int4 candidate_result;
candidate_result.zw = (int2)convert_int_rtn(factor * 20.f);
candidate_result.x = x;
candidate_result.y = y;
atomic_inc(glboutindex);
candidate[outputoff + temp + lcl_id] = candidate_result;
}
barrier(CLK_LOCAL_MEM_FENCE);
}
}
}
}
}
__kernel void gpuscaleclassifier(global GpuHidHaarTreeNode * orinode, global GpuHidHaarTreeNode * newnode,float scale,float weight_scale,int nodenum)
__kernel void gpuscaleclassifier(global GpuHidHaarTreeNode *orinode, global GpuHidHaarTreeNode *newnode, float scale, float weight_scale, int nodenum)
{
int counter=get_global_id(0);
int tr_x[3],tr_y[3],tr_h[3],tr_w[3],i=0;
GpuHidHaarTreeNode t1 = *(orinode + counter);
#pragma unroll
for(i=0;i<3;i++){
tr_x[i]=(int)(t1.p[i][0]*scale+0.5f);
tr_y[i]=(int)(t1.p[i][1]*scale+0.5f);
tr_w[i]=(int)(t1.p[i][2]*scale+0.5f);
tr_h[i]=(int)(t1.p[i][3]*scale+0.5f);
}
t1.weight[0]=t1.p[2][0]?-(t1.weight[1]*tr_h[1]*tr_w[1]+t1.weight[2]*tr_h[2]*tr_w[2])/(tr_h[0]*tr_w[0]):-t1.weight[1]*tr_h[1]*tr_w[1]/(tr_h[0]*tr_w[0]);
counter+=nodenum;
#pragma unroll
for(i=0;i<3;i++)
{
newnode[counter].p[i][0]=tr_x[i];
newnode[counter].p[i][1]=tr_y[i];
newnode[counter].p[i][2]=tr_x[i]+tr_w[i];
newnode[counter].p[i][3]=tr_y[i]+tr_h[i];
newnode[counter].weight[i]=t1.weight[i]*weight_scale;
}
newnode[counter].left=t1.left;
newnode[counter].right=t1.right;
newnode[counter].threshold=t1.threshold;
newnode[counter].alpha[0]=t1.alpha[0];
newnode[counter].alpha[1]=t1.alpha[1];
int counter = get_global_id(0);
int tr_x[3], tr_y[3], tr_h[3], tr_w[3], i = 0;
GpuHidHaarTreeNode t1 = *(orinode + counter);
#pragma unroll
for (i = 0; i < 3; i++)
{
tr_x[i] = (int)(t1.p[i][0] * scale + 0.5f);
tr_y[i] = (int)(t1.p[i][1] * scale + 0.5f);
tr_w[i] = (int)(t1.p[i][2] * scale + 0.5f);
tr_h[i] = (int)(t1.p[i][3] * scale + 0.5f);
}
t1.weight[0] = t1.p[2][0] ? -(t1.weight[1] * tr_h[1] * tr_w[1] + t1.weight[2] * tr_h[2] * tr_w[2]) / (tr_h[0] * tr_w[0]) : -t1.weight[1] * tr_h[1] * tr_w[1] / (tr_h[0] * tr_w[0]);
counter += nodenum;
#pragma unroll
for (i = 0; i < 3; i++)
{
newnode[counter].p[i][0] = tr_x[i];
newnode[counter].p[i][1] = tr_y[i];
newnode[counter].p[i][2] = tr_x[i] + tr_w[i];
newnode[counter].p[i][3] = tr_y[i] + tr_h[i];
newnode[counter].weight[i] = t1.weight[i] * weight_scale;
}
newnode[counter].left = t1.left;
newnode[counter].right = t1.right;
newnode[counter].threshold = t1.threshold;
newnode[counter].alpha[0] = t1.alpha[0];
newnode[counter].alpha[1] = t1.alpha[1];
}

8
modules/ocl/src/kernels/imgproc_canny.cl
View File

@ -691,7 +691,7 @@ __kernel
}
barrier(CLK_LOCAL_MEM_FENCE);
int ind = grp_idy * get_num_groups(0) + grp_idx;
int ind = mad24(grp_idy, (int)get_local_size(0), grp_idx);
if(ind < count)
{
@ -714,10 +714,10 @@ __kernel
}
barrier(CLK_LOCAL_MEM_FENCE);
while (s_counter > 0 && s_counter <= stack_size - get_num_groups(0))
while (s_counter > 0 && s_counter <= stack_size - get_local_size(0))
{
const int subTaskIdx = lidx >> 3;
const int portion = min(s_counter, get_num_groups(0) >> 3);
const int portion = min(s_counter, get_local_size(0)>> 3);
pos.x = pos.y = 0;
@ -757,7 +757,7 @@ __kernel
ind = s_ind;
for (int i = lidx; i < s_counter; i += get_num_groups(0))
for (int i = lidx; i < s_counter; i += get_local_size(0))
{
st2[ind + i] = s_st[i];
}

845
modules/ocl/src/kernels/pyrlk.cl
View File

File diff suppressed because it is too large Load Diff

764
modules/ocl/src/kernels/pyrlk_no_image.cl Normal file
View File

@ -0,0 +1,764 @@
/*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-2012, Multicoreware, Inc., all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
// Sen Liu, sen@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 oclMaterials 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.
//
//M*/
#define BUFFER 256
void reduce3(float val1, float val2, float val3, __local float *smem1, __local float *smem2, __local float *smem3, int tid)
{
smem1[tid] = val1;
smem2[tid] = val2;
smem3[tid] = val3;
barrier(CLK_LOCAL_MEM_FENCE);
#if BUFFER > 128
if (tid < 128)
{
smem1[tid] = val1 += smem1[tid + 128];
smem2[tid] = val2 += smem2[tid + 128];
smem3[tid] = val3 += smem3[tid + 128];
}
barrier(CLK_LOCAL_MEM_FENCE);
#endif
#if BUFFER > 64
if (tid < 64)
{
smem1[tid] = val1 += smem1[tid + 64];
smem2[tid] = val2 += smem2[tid + 64];
smem3[tid] = val3 += smem3[tid + 64];
}
barrier(CLK_LOCAL_MEM_FENCE);
#endif
if (tid < 32)
{
smem1[tid] = val1 += smem1[tid + 32];
smem2[tid] = val2 += smem2[tid + 32];
smem3[tid] = val3 += smem3[tid + 32];
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 16)
{
smem1[tid] = val1 += smem1[tid + 16];
smem2[tid] = val2 += smem2[tid + 16];
smem3[tid] = val3 += smem3[tid + 16];
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 8)
{
volatile __local float *vmem1 = smem1;
volatile __local float *vmem2 = smem2;
volatile __local float *vmem3 = smem3;
vmem1[tid] = val1 += vmem1[tid + 8];
vmem2[tid] = val2 += vmem2[tid + 8];
vmem3[tid] = val3 += vmem3[tid + 8];
vmem1[tid] = val1 += vmem1[tid + 4];
vmem2[tid] = val2 += vmem2[tid + 4];
vmem3[tid] = val3 += vmem3[tid + 4];
vmem1[tid] = val1 += vmem1[tid + 2];
vmem2[tid] = val2 += vmem2[tid + 2];
vmem3[tid] = val3 += vmem3[tid + 2];
vmem1[tid] = val1 += vmem1[tid + 1];
vmem2[tid] = val2 += vmem2[tid + 1];
vmem3[tid] = val3 += vmem3[tid + 1];
}
}
void reduce2(float val1, float val2, __local float *smem1, __local float *smem2, int tid)
{
smem1[tid] = val1;
smem2[tid] = val2;
barrier(CLK_LOCAL_MEM_FENCE);
#if BUFFER > 128
if (tid < 128)
{
smem1[tid] = val1 += smem1[tid + 128];
smem2[tid] = val2 += smem2[tid + 128];
}
barrier(CLK_LOCAL_MEM_FENCE);
#endif
#if BUFFER > 64
if (tid < 64)
{
smem1[tid] = val1 += smem1[tid + 64];
smem2[tid] = val2 += smem2[tid + 64];
}
barrier(CLK_LOCAL_MEM_FENCE);
#endif
if (tid < 32)
{
smem1[tid] = val1 += smem1[tid + 32];
smem2[tid] = val2 += smem2[tid + 32];
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 16)
{
smem1[tid] = val1 += smem1[tid + 16];
smem2[tid] = val2 += smem2[tid + 16];
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 8)
{
volatile __local float *vmem1 = smem1;
volatile __local float *vmem2 = smem2;
vmem1[tid] = val1 += vmem1[tid + 8];
vmem2[tid] = val2 += vmem2[tid + 8];
vmem1[tid] = val1 += vmem1[tid + 4];
vmem2[tid] = val2 += vmem2[tid + 4];
vmem1[tid] = val1 += vmem1[tid + 2];
vmem2[tid] = val2 += vmem2[tid + 2];
vmem1[tid] = val1 += vmem1[tid + 1];
vmem2[tid] = val2 += vmem2[tid + 1];
}
}
void reduce1(float val1, __local float *smem1, int tid)
{
smem1[tid] = val1;
barrier(CLK_LOCAL_MEM_FENCE);
#if BUFFER > 128
if (tid < 128)
{
smem1[tid] = val1 += smem1[tid + 128];
}
barrier(CLK_LOCAL_MEM_FENCE);
#endif
#if BUFFER > 64
if (tid < 64)
{
smem1[tid] = val1 += smem1[tid + 64];
}
barrier(CLK_LOCAL_MEM_FENCE);
#endif
if (tid < 32)
{
smem1[tid] = val1 += smem1[tid + 32];
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 16)
{
volatile __local float *vmem1 = smem1;
vmem1[tid] = val1 += vmem1[tid + 16];
}
barrier(CLK_LOCAL_MEM_FENCE);
if (tid < 8)
{
volatile __local float *vmem1 = smem1;
vmem1[tid] = val1 += vmem1[tid + 8];
vmem1[tid] = val1 += vmem1[tid + 4];
vmem1[tid] = val1 += vmem1[tid + 2];
vmem1[tid] = val1 += vmem1[tid + 1];
}
}
#define SCALE (1.0f / (1 << 20))
#define THRESHOLD 0.01f
#define DIMENSION 21
float readImage2Df_C1(__global const float *image, const float x, const float y, const int rows, const int cols, const int elemCntPerRow)
{
float2 coor = (float2)(x, y);
int i0 = clamp((int)floor(coor.x), 0, cols - 1);
int j0 = clamp((int)floor(coor.y), 0, rows - 1);
int i1 = clamp((int)floor(coor.x) + 1, 0, cols - 1);
int j1 = clamp((int)floor(coor.y) + 1, 0, rows - 1);
float a = coor.x - floor(coor.x);
float b = coor.y - floor(coor.y);
return (1 - a) * (1 - b) * image[mad24(j0, elemCntPerRow, i0)]
+ a * (1 - b) * image[mad24(j0, elemCntPerRow, i1)]
+ (1 - a) * b * image[mad24(j1, elemCntPerRow, i0)]
+ a * b * image[mad24(j1, elemCntPerRow, i1)];
}
__kernel void lkSparse_C1_D5(__global const float *I, __global const float *J,
__global const float2 *prevPts, int prevPtsStep, __global float2 *nextPts, int nextPtsStep, __global uchar *status, __global float *err,
const int level, const int rows, const int cols, const int elemCntPerRow,
int PATCH_X, int PATCH_Y, int cn, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr)
{
__local float smem1[BUFFER];
__local float smem2[BUFFER];
__local float smem3[BUFFER];
float2 c_halfWin = (float2)((c_winSize_x - 1) >> 1, (c_winSize_y - 1) >> 1);
const int tid = mad24(get_local_id(1), get_local_size(0), get_local_id(0));
float2 prevPt = prevPts[get_group_id(0)] * (1.0f / (1 << level));
if (prevPt.x < 0 || prevPt.x >= cols || prevPt.y < 0 || prevPt.y >= rows)
{
if (tid == 0 && level == 0)
{
status[get_group_id(0)] = 0;
}
return;
}
prevPt -= c_halfWin;
// extract the patch from the first image, compute covariation matrix of derivatives
float A11 = 0;
float A12 = 0;
float A22 = 0;
float I_patch[1][3];
float dIdx_patch[1][3];
float dIdy_patch[1][3];
for (int yBase = get_local_id(1), i = 0; yBase < c_winSize_y; yBase += get_local_size(1), ++i)
{
for (int xBase = get_local_id(0), j = 0; xBase < c_winSize_x; xBase += get_local_size(0), ++j)
{
float x = (prevPt.x + xBase);
float y = (prevPt.y + yBase);
I_patch[i][j] = readImage2Df_C1(I, x, y, rows, cols, elemCntPerRow);
float dIdx = 3.0f * readImage2Df_C1(I, x + 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C1(I, x + 1, y, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C1(I, x + 1, y + 1, rows, cols, elemCntPerRow) -
(3.0f * readImage2Df_C1(I, x - 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C1(I, x - 1, y, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C1(I, x - 1, y + 1, rows, cols, elemCntPerRow));
float dIdy = 3.0f * readImage2Df_C1(I, x - 1, y + 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C1(I, x, y + 1, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C1(I, x + 1, y + 1, rows, cols, elemCntPerRow) -
(3.0f * readImage2Df_C1(I, x - 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C1(I, x, y - 1, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C1(I, x + 1, y - 1, rows, cols, elemCntPerRow));
dIdx_patch[i][j] = dIdx;
dIdy_patch[i][j] = dIdy;
A11 += dIdx * dIdx;
A12 += dIdx * dIdy;
A22 += dIdy * dIdy;
}
}
reduce3(A11, A12, A22, smem1, smem2, smem3, tid);
barrier(CLK_LOCAL_MEM_FENCE);
A11 = smem1[0];
A12 = smem2[0];
A22 = smem3[0];
float D = A11 * A22 - A12 * A12;
if (D < 1.192092896e-07f)
{
if (tid == 0 && level == 0)
{
status[get_group_id(0)] = 0;
}
return;
}
D = 1.f / D;
A11 *= D;
A12 *= D;
A22 *= D;
float2 nextPt = nextPts[get_group_id(0)];
nextPt = nextPt * 2.0f - c_halfWin;
for (int k = 0; k < c_iters; ++k)
{
if (nextPt.x < -c_halfWin.x || nextPt.x >= cols || nextPt.y < -c_halfWin.y || nextPt.y >= rows)
{
if (tid == 0 && level == 0)
{
status[get_group_id(0)] = 0;
}
return;
}
float b1 = 0;
float b2 = 0;
for (int y = get_local_id(1), i = 0; y < c_winSize_y; y += get_local_size(1), ++i)
{
for (int x = get_local_id(0), j = 0; x < c_winSize_x; x += get_local_size(0), ++j)
{
float diff = (readImage2Df_C1(J, nextPt.x + x, nextPt.y + y, rows, cols, elemCntPerRow) - I_patch[i][j]) * 32.0f;
b1 += diff * dIdx_patch[i][j];
b2 += diff * dIdy_patch[i][j];
}
}
reduce2(b1, b2, smem1, smem2, tid);
barrier(CLK_LOCAL_MEM_FENCE);
b1 = smem1[0];
b2 = smem2[0];
float2 delta;
delta.x = A12 * b2 - A22 * b1;
delta.y = A12 * b1 - A11 * b2;
nextPt += delta;
//if (fabs(delta.x) < THRESHOLD && fabs(delta.y) < THRESHOLD)
// break;
}
float errval = 0.0f;
if (calcErr)
{
for (int y = get_local_id(1), i = 0; y < c_winSize_y; y += get_local_size(1), ++i)
{
for (int x = get_local_id(0), j = 0; x < c_winSize_x; x += get_local_size(0), ++j)
{
float diff = readImage2Df_C1(J, nextPt.x + x, nextPt.y + y, rows, cols, elemCntPerRow) - I_patch[i][j];
errval += fabs(diff);
}
}
reduce1(errval, smem1, tid);
}
if (tid == 0)
{
nextPt += c_halfWin;
nextPts[get_group_id(0)] = nextPt;
if (calcErr)
{
err[get_group_id(0)] = smem1[0] / (c_winSize_x * c_winSize_y);
}
}
}
float4 readImage2Df_C4(__global const float4 *image, const float x, const float y, const int rows, const int cols, const int elemCntPerRow)
{
float2 coor = (float2)(x, y);
int i0 = clamp((int)floor(coor.x), 0, cols - 1);
int j0 = clamp((int)floor(coor.y), 0, rows - 1);
int i1 = clamp((int)floor(coor.x) + 1, 0, cols - 1);
int j1 = clamp((int)floor(coor.y) + 1, 0, rows - 1);
float a = coor.x - floor(coor.x);
float b = coor.y - floor(coor.y);
return (1 - a) * (1 - b) * image[mad24(j0, elemCntPerRow, i0)]
+ a * (1 - b) * image[mad24(j0, elemCntPerRow, i1)]
+ (1 - a) * b * image[mad24(j1, elemCntPerRow, i0)]
+ a * b * image[mad24(j1, elemCntPerRow, i1)];
}
__kernel void lkSparse_C4_D5(__global const float *I, __global const float *J,
__global const float2 *prevPts, int prevPtsStep, __global float2 *nextPts, int nextPtsStep, __global uchar *status, __global float *err,
const int level, const int rows, const int cols, const int elemCntPerRow,
int PATCH_X, int PATCH_Y, int cn, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr)
{
__local float smem1[BUFFER];
__local float smem2[BUFFER];
__local float smem3[BUFFER];
float2 c_halfWin = (float2)((c_winSize_x - 1) >> 1, (c_winSize_y - 1) >> 1);
const int tid = mad24(get_local_id(1), get_local_size(0), get_local_id(0));
float2 prevPt = prevPts[get_group_id(0)] * (1.0f / (1 << level));
if (prevPt.x < 0 || prevPt.x >= cols || prevPt.y < 0 || prevPt.y >= rows)
{
if (tid == 0 && level == 0)
{
status[get_group_id(0)] = 0;
}
return;
}
prevPt -= c_halfWin;
// extract the patch from the first image, compute covariation matrix of derivatives
float A11 = 0;
float A12 = 0;
float A22 = 0;
float4 I_patch[1][3];
float4 dIdx_patch[1][3];
float4 dIdy_patch[1][3];
__global float4 *ptrI = (__global float4 *)I;
for (int yBase = get_local_id(1), i = 0; yBase < c_winSize_y; yBase += get_local_size(1), ++i)
{
for (int xBase = get_local_id(0), j = 0; xBase < c_winSize_x; xBase += get_local_size(0), ++j)
{
float x = (prevPt.x + xBase);
float y = (prevPt.y + yBase);
I_patch[i][j] = readImage2Df_C4(ptrI, x, y, rows, cols, elemCntPerRow);
float4 dIdx = 3.0f * readImage2Df_C4(ptrI, x + 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C4(ptrI, x + 1, y, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C4(ptrI, x + 1, y + 1, rows, cols, elemCntPerRow) -
(3.0f * readImage2Df_C4(ptrI, x - 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C4(ptrI, x - 1, y, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C4(ptrI, x - 1, y + 1, rows, cols, elemCntPerRow));
float4 dIdy = 3.0f * readImage2Df_C4(ptrI, x - 1, y + 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C4(ptrI, x, y + 1, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C4(ptrI, x + 1, y + 1, rows, cols, elemCntPerRow) -
(3.0f * readImage2Df_C4(ptrI, x - 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C4(ptrI, x, y - 1, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C4(ptrI, x + 1, y - 1, rows, cols, elemCntPerRow));
dIdx_patch[i][j] = dIdx;
dIdy_patch[i][j] = dIdy;
A11 += (dIdx * dIdx).x + (dIdx * dIdx).y + (dIdx * dIdx).z;
A12 += (dIdx * dIdy).x + (dIdx * dIdy).y + (dIdx * dIdy).z;
A22 += (dIdy * dIdy).x + (dIdy * dIdy).y + (dIdy * dIdy).z;
}
}
reduce3(A11, A12, A22, smem1, smem2, smem3, tid);
barrier(CLK_LOCAL_MEM_FENCE);
A11 = smem1[0];
A12 = smem2[0];
A22 = smem3[0];
float D = A11 * A22 - A12 * A12;
//pD[get_group_id(0)] = D;
if (D < 1.192092896e-07f)
{
if (tid == 0 && level == 0)
{
status[get_group_id(0)] = 0;
}
return;
}
D = 1.f / D;
A11 *= D;
A12 *= D;
A22 *= D;
float2 nextPt = nextPts[get_group_id(0)];
nextPt = nextPt * 2.0f - c_halfWin;
__global float4 *ptrJ = (__global float4 *)J;
for (int k = 0; k < c_iters; ++k)
{
if (nextPt.x < -c_halfWin.x || nextPt.x >= cols || nextPt.y < -c_halfWin.y || nextPt.y >= rows)
{
if (tid == 0 && level == 0)
{
status[get_group_id(0)] = 0;
}
return;
}
float b1 = 0;
float b2 = 0;
for (int y = get_local_id(1), i = 0; y < c_winSize_y; y += get_local_size(1), ++i)
{
for (int x = get_local_id(0), j = 0; x < c_winSize_x; x += get_local_size(0), ++j)
{
float4 diff = (readImage2Df_C4(ptrJ, nextPt.x + x, nextPt.y + y, rows, cols, elemCntPerRow) - I_patch[i][j]) * 32.0f;
b1 += (diff * dIdx_patch[i][j]).x + (diff * dIdx_patch[i][j]).y + (diff * dIdx_patch[i][j]).z;
b2 += (diff * dIdy_patch[i][j]).x + (diff * dIdy_patch[i][j]).y + (diff * dIdy_patch[i][j]).z;
}
}
reduce2(b1, b2, smem1, smem2, tid);
barrier(CLK_LOCAL_MEM_FENCE);
b1 = smem1[0];
b2 = smem2[0];
float2 delta;
delta.x = A12 * b2 - A22 * b1;
delta.y = A12 * b1 - A11 * b2;
nextPt += delta;
//if (fabs(delta.x) < THRESHOLD && fabs(delta.y) < THRESHOLD)
// break;
}
float errval = 0.0f;
if (calcErr)
{
for (int y = get_local_id(1), i = 0; y < c_winSize_y; y += get_local_size(1), ++i)
{
for (int x = get_local_id(0), j = 0; x < c_winSize_x; x += get_local_size(0), ++j)
{
float4 diff = readImage2Df_C4(ptrJ, nextPt.x + x, nextPt.y + y, rows, cols, elemCntPerRow) - I_patch[i][j];
errval += fabs(diff.x) + fabs(diff.y) + fabs(diff.z);
}
}
reduce1(errval, smem1, tid);
}
if (tid == 0)
{
nextPt += c_halfWin;
nextPts[get_group_id(0)] = nextPt;
if (calcErr)
{
err[get_group_id(0)] = smem1[0] / (3 * c_winSize_x * c_winSize_y);
}
}
}
int readImage2Di_C1(__global const int *image, float2 coor, int2 size, const int elemCntPerRow)
{
int i = clamp((int)floor(coor.x), 0, size.x - 1);
int j = clamp((int)floor(coor.y), 0, size.y - 1);
return image[mad24(j, elemCntPerRow, i)];
}
__kernel void lkDense_C1_D0(__global const int *I, __global const int *J, __global float *u, int uStep, __global float *v, int vStep, __global const float *prevU, int prevUStep, __global const float *prevV, int prevVStep,
const int rows, const int cols, /*__global float* err, int errStep, int cn,*/
const int elemCntPerRow, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr)
{
int c_halfWin_x = (c_winSize_x - 1) / 2;
int c_halfWin_y = (c_winSize_y - 1) / 2;
const int patchWidth = get_local_size(0) + 2 * c_halfWin_x;
const int patchHeight = get_local_size(1) + 2 * c_halfWin_y;
__local int smem[8192];
__local int *I_patch = smem;
__local int *dIdx_patch = I_patch + patchWidth * patchHeight;
__local int *dIdy_patch = dIdx_patch + patchWidth * patchHeight;
const int xBase = get_group_id(0) * get_local_size(0);
const int yBase = get_group_id(1) * get_local_size(1);
int2 size = (int2)(cols, rows);
for (int i = get_local_id(1); i < patchHeight; i += get_local_size(1))
{
for (int j = get_local_id(0); j < patchWidth; j += get_local_size(0))
{
float x = xBase - c_halfWin_x + j + 0.5f;
float y = yBase - c_halfWin_y + i + 0.5f;
I_patch[i * patchWidth + j] = readImage2Di_C1(I, (float2)(x, y), size, elemCntPerRow);
// Sharr Deriv
dIdx_patch[i * patchWidth + j] = 3 * readImage2Di_C1(I, (float2)(x + 1, y - 1), size, elemCntPerRow) + 10 * readImage2Di_C1(I, (float2)(x + 1, y), size, elemCntPerRow) + 3 * readImage2Di_C1(I, (float2)(x + 1, y + 1), size, elemCntPerRow) -
(3 * readImage2Di_C1(I, (float2)(x - 1, y - 1), size, elemCntPerRow) + 10 * readImage2Di_C1(I, (float2)(x - 1, y), size, elemCntPerRow) + 3 * readImage2Di_C1(I, (float2)(x - 1, y + 1), size, elemCntPerRow));
dIdy_patch[i * patchWidth + j] = 3 * readImage2Di_C1(I, (float2)(x - 1, y + 1), size, elemCntPerRow) + 10 * readImage2Di_C1(I, (float2)(x, y + 1), size, elemCntPerRow) + 3 * readImage2Di_C1(I, (float2)(x + 1, y + 1), size, elemCntPerRow) -
(3 * readImage2Di_C1(I, (float2)(x - 1, y - 1), size, elemCntPerRow) + 10 * readImage2Di_C1(I, (float2)(x, y - 1), size, elemCntPerRow) + 3 * readImage2Di_C1(I, (float2)(x + 1, y - 1), size, elemCntPerRow));
}
}
barrier(CLK_LOCAL_MEM_FENCE);
// extract the patch from the first image, compute covariation matrix of derivatives
const int x = get_global_id(0);
const int y = get_global_id(1);
if (x >= cols || y >= rows)
{
return;
}
int A11i = 0;
int A12i = 0;
int A22i = 0;
for (int i = 0; i < c_winSize_y; ++i)
{
for (int j = 0; j < c_winSize_x; ++j)
{
int dIdx = dIdx_patch[(get_local_id(1) + i) * patchWidth + (get_local_id(0) + j)];
int dIdy = dIdy_patch[(get_local_id(1) + i) * patchWidth + (get_local_id(0) + j)];
A11i += dIdx * dIdx;
A12i += dIdx * dIdy;
A22i += dIdy * dIdy;
}
}
float A11 = A11i;
float A12 = A12i;
float A22 = A22i;
float D = A11 * A22 - A12 * A12;
//if (calcErr && GET_MIN_EIGENVALS)
// (err + y * errStep)[x] = minEig;
if (D < 1.192092896e-07f)
{
//if (calcErr)
// err(y, x) = 3.402823466e+38f;
return;
}
D = 1.f / D;
A11 *= D;
A12 *= D;
A22 *= D;
float2 nextPt;
nextPt.x = x + prevU[y / 2 * prevUStep / 4 + x / 2] * 2.0f;
nextPt.y = y + prevV[y / 2 * prevVStep / 4 + x / 2] * 2.0f;
for (int k = 0; k < c_iters; ++k)
{
if (nextPt.x < 0 || nextPt.x >= cols || nextPt.y < 0 || nextPt.y >= rows)
{
//if (calcErr)
// err(y, x) = 3.402823466e+38f;
return;
}
int b1 = 0;
int b2 = 0;
for (int i = 0; i < c_winSize_y; ++i)
{
for (int j = 0; j < c_winSize_x; ++j)
{
int iI = I_patch[(get_local_id(1) + i) * patchWidth + get_local_id(0) + j];
int iJ = readImage2Di_C1(J, (float2)(nextPt.x - c_halfWin_x + j + 0.5f, nextPt.y - c_halfWin_y + i + 0.5f), size, elemCntPerRow);
int diff = (iJ - iI) * 32;
int dIdx = dIdx_patch[(get_local_id(1) + i) * patchWidth + (get_local_id(0) + j)];
int dIdy = dIdy_patch[(get_local_id(1) + i) * patchWidth + (get_local_id(0) + j)];
b1 += diff * dIdx;
b2 += diff * dIdy;
}
}
float2 delta;
delta.x = A12 * b2 - A22 * b1;
delta.y = A12 * b1 - A11 * b2;
nextPt.x += delta.x;
nextPt.y += delta.y;
if (fabs(delta.x) < 0.01f && fabs(delta.y) < 0.01f)
{
break;
}
}
u[y * uStep / 4 + x] = nextPt.x - x;
v[y * vStep / 4 + x] = nextPt.y - y;
if (calcErr)
{
int errval = 0;
for (int i = 0; i < c_winSize_y; ++i)
{
for (int j = 0; j < c_winSize_x; ++j)
{
int iI = I_patch[(get_local_id(1) + i) * patchWidth + get_local_id(0) + j];
int iJ = readImage2Di_C1(J, (float2)(nextPt.x - c_halfWin_x + j + 0.5f, nextPt.y - c_halfWin_y + i + 0.5f), size, elemCntPerRow);
errval += abs(iJ - iI);
}
}
//err[y * errStep / 4 + x] = static_cast<float>(errval) / (c_winSize_x * c_winSize_y);
}
}

2
modules/ocl/src/match_template.cpp
View File

@ -99,7 +99,7 @@ namespace cv
// Evaluates optimal template's area threshold. If
// template's area is less than the threshold, we use naive match
// template version, otherwise FFT-based (if available)
int getTemplateThreshold(int method, int depth)
static int getTemplateThreshold(int method, int depth)
{
switch (method)
{

14
modules/ocl/src/matrix_operations.cpp
View File

@ -126,7 +126,7 @@ namespace cv
////////////////////////////////////////////////////////////////////////
// convert_C3C4
void convert_C3C4(const cl_mem &src, oclMat &dst)
static void convert_C3C4(const cl_mem &src, oclMat &dst)
{
int dstStep_in_pixel = dst.step1() / dst.oclchannels();
int pixel_end = dst.wholecols * dst.wholerows - 1;
@ -174,7 +174,7 @@ void convert_C3C4(const cl_mem &src, oclMat &dst)
}
////////////////////////////////////////////////////////////////////////
// convert_C4C3
void convert_C4C3(const oclMat &src, cl_mem &dst)
static void convert_C4C3(const oclMat &src, cl_mem &dst)
{
int srcStep_in_pixel = src.step1() / src.oclchannels();
int pixel_end = src.wholecols * src.wholerows - 1;
@ -336,7 +336,7 @@ inline int divUp(int total, int grain)
///////////////////////////////////////////////////////////////////////////
////////////////////////////////// CopyTo /////////////////////////////////
///////////////////////////////////////////////////////////////////////////
void copy_to_with_mask(const oclMat &src, oclMat &dst, const oclMat &mask, string kernelName)
static void copy_to_with_mask(const oclMat &src, oclMat &dst, const oclMat &mask, string kernelName)
{
CV_DbgAssert( dst.rows == mask.rows && dst.cols == mask.cols &&
src.rows == dst.rows && src.cols == dst.cols
@ -401,7 +401,7 @@ void cv::ocl::oclMat::copyTo( oclMat &mat, const oclMat &mask) const
///////////////////////////////////////////////////////////////////////////
//////////////////////////////// ConvertTo ////////////////////////////////
///////////////////////////////////////////////////////////////////////////
void convert_run(const oclMat &src, oclMat &dst, double alpha, double beta)
static void convert_run(const oclMat &src, oclMat &dst, double alpha, double beta)
{
string kernelName = "convert_to_S";
stringstream idxStr;
@ -472,7 +472,7 @@ oclMat &cv::ocl::oclMat::operator = (const Scalar &s)
setTo(s);
return *this;
}
void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kernelName)
static void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kernelName)
{
vector<pair<size_t , const void *> > args;
@ -642,7 +642,7 @@ void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kern
default:
CV_Error(CV_StsUnsupportedFormat, "unknown depth");
}
#if CL_VERSION_1_2
#ifdef CL_VERSION_1_2
if(dst.offset == 0 && dst.cols == dst.wholecols)
{
clEnqueueFillBuffer(dst.clCxt->impl->clCmdQueue, (cl_mem)dst.data, args[0].second, args[0].first, 0, dst.step * dst.rows, 0, NULL, NULL);
@ -668,7 +668,7 @@ void set_to_withoutmask_run(const oclMat &dst, const Scalar &scalar, string kern
#endif
}
void set_to_withmask_run(const oclMat &dst, const Scalar &scalar, const oclMat &mask, string kernelName)
static void set_to_withmask_run(const oclMat &dst, const Scalar &scalar, const oclMat &mask, string kernelName)
{
CV_DbgAssert( dst.rows == mask.rows && dst.cols == mask.cols);
vector<pair<size_t , const void *> > args;

2
modules/ocl/src/mcwutil.cpp
View File

@ -62,7 +62,7 @@ namespace cv
}
// provide additional methods for the user to interact with the command queue after a task is fired
void openCLExecuteKernel_2(Context *clCxt , const char **source, string kernelName, size_t globalThreads[3],
static void openCLExecuteKernel_2(Context *clCxt , const char **source, string kernelName, size_t globalThreads[3],
size_t localThreads[3], vector< pair<size_t, const void *> > &args, int channels,
int depth, char *build_options, FLUSH_MODE finish_mode)
{

2
modules/ocl/src/precomp.hpp
View File

@ -46,7 +46,7 @@
#ifndef __OPENCV_PRECOMP_H__
#define __OPENCV_PRECOMP_H__
#if _MSC_VER >= 1200
#if defined _MSC_VER && _MSC_VER >= 1200
#pragma warning( disable: 4267 4324 4244 4251 4710 4711 4514 4996 )
#endif

2
modules/ocl/src/pyrdown.cpp
View File

@ -66,7 +66,7 @@ namespace cv
//////////////////////////////////////////////////////////////////////////////
/////////////////////// add subtract multiply divide /////////////////////////
//////////////////////////////////////////////////////////////////////////////
void pyrdown_run(const oclMat &src, const oclMat &dst)
static void pyrdown_run(const oclMat &src, const oclMat &dst)
{
CV_Assert(src.type() == dst.type());

415
modules/ocl/src/pyrlk.cpp
View File

@ -48,23 +48,24 @@ using namespace cv::ocl;
#if !defined (HAVE_OPENCL)
void cv::ocl::PyrLKOpticalFlow::sparse(const oclMat &, const oclMat &, const oclMat &, oclMat &, oclMat &, oclMat *) { }
void cv::ocl::PyrLKOpticalFlow::sparse(const oclMat &, const oclMat &, const oclMat &, oclMat &, oclMat &, oclMat &) { }
void cv::ocl::PyrLKOpticalFlow::dense(const oclMat &, const oclMat &, oclMat &, oclMat &, oclMat *) { }
#else /* !defined (HAVE_OPENCL) */
namespace cv
{
namespace ocl
{
///////////////////////////OpenCL kernel strings///////////////////////////
extern const char *pyrlk;
extern const char *operator_setTo;
extern const char *operator_convertTo;
extern const char *operator_copyToM;
extern const char *arithm_mul;
extern const char *pyr_down;
}
namespace ocl
{
///////////////////////////OpenCL kernel strings///////////////////////////
extern const char *pyrlk;
extern const char *pyrlk_no_image;
extern const char *operator_setTo;
extern const char *operator_convertTo;
extern const char *operator_copyToM;
extern const char *arithm_mul;
extern const char *pyr_down;
}
}
struct dim3
@ -84,26 +85,26 @@ struct int2
namespace
{
void calcPatchSize(cv::Size winSize, int cn, dim3 &block, dim3 &patch, bool isDeviceArch11)
void calcPatchSize(cv::Size winSize, int cn, dim3 &block, dim3 &patch, bool isDeviceArch11)
{
winSize.width *= cn;
if (winSize.width > 32 && winSize.width > 2 * winSize.height)
{
winSize.width *= cn;
if (winSize.width > 32 && winSize.width > 2 * winSize.height)
{
block.x = isDeviceArch11 ? 16 : 32;
block.y = 8;
}
else
{
block.x = 16;
block.y = isDeviceArch11 ? 8 : 16;
}
patch.x = (winSize.width + block.x - 1) / block.x;
patch.y = (winSize.height + block.y - 1) / block.y;
block.z = patch.z = 1;
block.x = isDeviceArch11 ? 16 : 32;
block.y = 8;
}
else
{
block.x = 16;
block.y = isDeviceArch11 ? 8 : 16;
}
patch.x = (winSize.width + block.x - 1) / block.x;
patch.y = (winSize.height + block.y - 1) / block.y;
block.z = patch.z = 1;
}
}
inline int divUp(int total, int grain)
@ -114,7 +115,7 @@ inline int divUp(int total, int grain)
///////////////////////////////////////////////////////////////////////////
//////////////////////////////// ConvertTo ////////////////////////////////
///////////////////////////////////////////////////////////////////////////
void convert_run_cus(const oclMat &src, oclMat &dst, double alpha, double beta)
static void convert_run_cus(const oclMat &src, oclMat &dst, double alpha, double beta)
{
string kernelName = "convert_to_S";
stringstream idxStr;
@ -185,7 +186,7 @@ void convertTo( const oclMat &src, oclMat &dst, int rtype, double alpha, double
// setTo(s);
// return *this;
//}
void set_to_withoutmask_run_cus(const oclMat &dst, const Scalar &scalar, string kernelName)
static void set_to_withoutmask_run_cus(const oclMat &dst, const Scalar &scalar, string kernelName)
{
vector<pair<size_t , const void *> > args;
@ -355,7 +356,7 @@ void set_to_withoutmask_run_cus(const oclMat &dst, const Scalar &scalar, string
default:
CV_Error(CV_StsUnsupportedFormat, "unknown depth");
}
#if CL_VERSION_1_2
#ifdef CL_VERSION_1_2
if(dst.offset == 0 && dst.cols == dst.wholecols)
{
clEnqueueFillBuffer(dst.clCxt->impl->clCmdQueue, (cl_mem)dst.data, args[0].second, args[0].first, 0, dst.step * dst.rows, 0, NULL, NULL);
@ -381,7 +382,7 @@ void set_to_withoutmask_run_cus(const oclMat &dst, const Scalar &scalar, string
#endif
}
oclMat &setTo(oclMat &src, const Scalar &scalar)
static oclMat &setTo(oclMat &src, const Scalar &scalar)
{
CV_Assert( src.depth() >= 0 && src.depth() <= 6 );
CV_DbgAssert( !src.empty());
@ -401,48 +402,48 @@ oclMat &setTo(oclMat &src, const Scalar &scalar)
///////////////////////////////////////////////////////////////////////////
////////////////////////////////// CopyTo /////////////////////////////////
///////////////////////////////////////////////////////////////////////////
void copy_to_with_mask_cus(const oclMat &src, oclMat &dst, const oclMat &mask, string kernelName)
{
CV_DbgAssert( dst.rows == mask.rows && dst.cols == mask.cols &&
src.rows == dst.rows && src.cols == dst.cols
&& mask.type() == CV_8UC1);
// static void copy_to_with_mask_cus(const oclMat &src, oclMat &dst, const oclMat &mask, string kernelName)
// {
// CV_DbgAssert( dst.rows == mask.rows && dst.cols == mask.cols &&
// src.rows == dst.rows && src.cols == dst.cols
// && mask.type() == CV_8UC1);
vector<pair<size_t , const void *> > args;
// vector<pair<size_t , const void *> > args;
std::string string_types[4][7] = {{"uchar", "char", "ushort", "short", "int", "float", "double"},
{"uchar2", "char2", "ushort2", "short2", "int2", "float2", "double2"},
{"uchar3", "char3", "ushort3", "short3", "int3", "float3", "double3"},
{"uchar4", "char4", "ushort4", "short4", "int4", "float4", "double4"}
};
char compile_option[32];
sprintf(compile_option, "-D GENTYPE=%s", string_types[dst.oclchannels() - 1][dst.depth()].c_str());
size_t localThreads[3] = {16, 16, 1};
size_t globalThreads[3];
// std::string string_types[4][7] = {{"uchar", "char", "ushort", "short", "int", "float", "double"},
// {"uchar2", "char2", "ushort2", "short2", "int2", "float2", "double2"},
// {"uchar3", "char3", "ushort3", "short3", "int3", "float3", "double3"},
// {"uchar4", "char4", "ushort4", "short4", "int4", "float4", "double4"}
// };
// char compile_option[32];
// sprintf(compile_option, "-D GENTYPE=%s", string_types[dst.oclchannels() - 1][dst.depth()].c_str());
// size_t localThreads[3] = {16, 16, 1};
// size_t globalThreads[3];
globalThreads[0] = divUp(dst.cols, localThreads[0]) * localThreads[0];
globalThreads[1] = divUp(dst.rows, localThreads[1]) * localThreads[1];
globalThreads[2] = 1;
// globalThreads[0] = divUp(dst.cols, localThreads[0]) * localThreads[0];
// globalThreads[1] = divUp(dst.rows, localThreads[1]) * localThreads[1];
// globalThreads[2] = 1;
int dststep_in_pixel = dst.step / dst.elemSize(), dstoffset_in_pixel = dst.offset / dst.elemSize();
int srcstep_in_pixel = src.step / src.elemSize(), srcoffset_in_pixel = src.offset / src.elemSize();
// int dststep_in_pixel = dst.step / dst.elemSize(), dstoffset_in_pixel = dst.offset / dst.elemSize();
// int srcstep_in_pixel = src.step / src.elemSize(), srcoffset_in_pixel = src.offset / src.elemSize();
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data ));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data ));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&mask.data ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&srcstep_in_pixel ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&srcoffset_in_pixel ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dststep_in_pixel ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dstoffset_in_pixel ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&mask.step ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&mask.offset ));
// args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data ));
// args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data ));
// args.push_back( make_pair( sizeof(cl_mem) , (void *)&mask.data ));
// args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols ));
// args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows ));
// args.push_back( make_pair( sizeof(cl_int) , (void *)&srcstep_in_pixel ));
// args.push_back( make_pair( sizeof(cl_int) , (void *)&srcoffset_in_pixel ));
// args.push_back( make_pair( sizeof(cl_int) , (void *)&dststep_in_pixel ));
// args.push_back( make_pair( sizeof(cl_int) , (void *)&dstoffset_in_pixel ));
// args.push_back( make_pair( sizeof(cl_int) , (void *)&mask.step ));
// args.push_back( make_pair( sizeof(cl_int) , (void *)&mask.offset ));
openCLExecuteKernel2(dst.clCxt , &operator_copyToM, kernelName, globalThreads,
localThreads, args, -1, -1, compile_option, CLFLUSH);
}
// openCLExecuteKernel2(dst.clCxt , &operator_copyToM, kernelName, globalThreads,
// localThreads, args, -1, -1, compile_option, CLFLUSH);
// }
void copyTo(const oclMat &src, oclMat &m )
static void copyTo(const oclMat &src, oclMat &m )
{
CV_DbgAssert(!src.empty());
m.create(src.size(), src.type());
@ -450,20 +451,20 @@ void copyTo(const oclMat &src, oclMat &m )
src.data, src.step, src.cols * src.elemSize(), src.rows, src.offset);
}
void copyTo(const oclMat &src, oclMat &mat, const oclMat &mask)
{
if (mask.empty())
{
copyTo(src, mat);
}
else
{
mat.create(src.size(), src.type());
copy_to_with_mask_cus(src, mat, mask, "copy_to_with_mask");
}
}
// static void copyTo(const oclMat &src, oclMat &mat, const oclMat &mask)
// {
// if (mask.empty())
// {
// copyTo(src, mat);
// }
// else
// {
// mat.create(src.size(), src.type());
// copy_to_with_mask_cus(src, mat, mask, "copy_to_with_mask");
// }
// }
void arithmetic_run(const oclMat &src1, oclMat &dst, string kernelName, const char **kernelString, void *_scalar)
static void arithmetic_run(const oclMat &src1, oclMat &dst, string kernelName, const char **kernelString, void *_scalar)
{
if(src1.clCxt -> impl -> double_support == 0 && src1.type() == CV_64F)
{
@ -528,12 +529,12 @@ void arithmetic_run(const oclMat &src1, oclMat &dst, string kernelName, const ch
openCLExecuteKernel2(clCxt, kernelString, kernelName, globalThreads, localThreads, args, -1, src1.depth(), CLFLUSH);
}
void multiply_cus(const oclMat &src1, oclMat &dst, float scalar)
static void multiply_cus(const oclMat &src1, oclMat &dst, float scalar)
{
arithmetic_run(src1, dst, "arithm_muls", &pyrlk, (void *)(&scalar));
arithmetic_run(src1, dst, "arithm_muls", &arithm_mul, (void *)(&scalar));
}
void pyrdown_run_cus(const oclMat &src, const oclMat &dst)
static void pyrdown_run_cus(const oclMat &src, const oclMat &dst)
{
CV_Assert(src.type() == dst.type());
@ -558,7 +559,7 @@ void pyrdown_run_cus(const oclMat &src, const oclMat &dst)
openCLExecuteKernel2(clCxt, &pyr_down, kernelName, globalThreads, localThreads, args, src.oclchannels(), src.depth(), CLFLUSH);
}
void pyrDown_cus(const oclMat &src, oclMat &dst)
static void pyrDown_cus(const oclMat &src, oclMat &dst)
{
CV_Assert(src.depth() <= CV_32F && src.channels() <= 4);
@ -581,26 +582,26 @@ void pyrDown_cus(const oclMat &src, oclMat &dst)
//
//void callF(const oclMat& src, oclMat& dst, MultiplyScalar op, int mask)
//{
// Mat srcTemp;
// Mat dstTemp;
// src.download(srcTemp);
// dst.download(dstTemp);
// Mat srcTemp;
// Mat dstTemp;
// src.download(srcTemp);
// dst.download(dstTemp);
//
// int i;
// int j;
// int k;
// for(i = 0; i < srcTemp.rows; i++)
// {
// for(j = 0; j < srcTemp.cols; j++)
// {
// for(k = 0; k < srcTemp.channels(); k++)
// {
// ((float*)dstTemp.data)[srcTemp.channels() * (i * srcTemp.rows + j) + k] = (float)op(((float*)srcTemp.data)[srcTemp.channels() * (i * srcTemp.rows + j) + k]);
// }
// }
// }
// int i;
// int j;
// int k;
// for(i = 0; i < srcTemp.rows; i++)
// {
// for(j = 0; j < srcTemp.cols; j++)
// {
// for(k = 0; k < srcTemp.channels(); k++)
// {
// ((float*)dstTemp.data)[srcTemp.channels() * (i * srcTemp.rows + j) + k] = (float)op(((float*)srcTemp.data)[srcTemp.channels() * (i * srcTemp.rows + j) + k]);
// }
// }
// }
//
// dst = dstTemp;
// dst = dstTemp;
//}
//
//static inline bool isAligned(const unsigned char* ptr, size_t size)
@ -622,57 +623,57 @@ void pyrDown_cus(const oclMat &src, oclMat &dst)
// return;
// }
//
// Mat srcTemp;
// Mat dstTemp;
// src.download(srcTemp);
// dst.download(dstTemp);
// Mat srcTemp;
// Mat dstTemp;
// src.download(srcTemp);
// dst.download(dstTemp);
//
// int x_shifted;
// int x_shifted;
//
// int i;
// int j;
// for(i = 0; i < srcTemp.rows; i++)
// {
// const double* srcRow = (const double*)srcTemp.data + i * srcTemp.rows;
// int i;
// int j;
// for(i = 0; i < srcTemp.rows; i++)
// {
// const double* srcRow = (const double*)srcTemp.data + i * srcTemp.rows;
// double* dstRow = (double*)dstTemp.data + i * dstTemp.rows;;
//
// for(j = 0; j < srcTemp.cols; j++)
// {
// x_shifted = j * 4;
// for(j = 0; j < srcTemp.cols; j++)
// {
// x_shifted = j * 4;
//
// if(x_shifted + 4 - 1 < srcTemp.cols)
// {
// dstRow[x_shifted ] = op(srcRow[x_shifted ]);
// dstRow[x_shifted + 1] = op(srcRow[x_shifted + 1]);
// dstRow[x_shifted + 2] = op(srcRow[x_shifted + 2]);
// dstRow[x_shifted + 3] = op(srcRow[x_shifted + 3]);
// }
// else
// {
// for (int real_x = x_shifted; real_x < srcTemp.cols; ++real_x)
// {
// ((float*)dstTemp.data)[i * srcTemp.rows + real_x] = op(((float*)srcTemp.data)[i * srcTemp.rows + real_x]);
// }
// }
// }
// }
// if(x_shifted + 4 - 1 < srcTemp.cols)
// {
// dstRow[x_shifted ] = op(srcRow[x_shifted ]);
// dstRow[x_shifted + 1] = op(srcRow[x_shifted + 1]);
// dstRow[x_shifted + 2] = op(srcRow[x_shifted + 2]);
// dstRow[x_shifted + 3] = op(srcRow[x_shifted + 3]);
// }
// else
// {
// for (int real_x = x_shifted; real_x < srcTemp.cols; ++real_x)
// {
// ((float*)dstTemp.data)[i * srcTemp.rows + real_x] = op(((float*)srcTemp.data)[i * srcTemp.rows + real_x]);
// }
// }
// }
// }
//}
//
//void multiply(const oclMat& src1, double val, oclMat& dst, double scale = 1.0f);
//void multiply(const oclMat& src1, double val, oclMat& dst, double scale)
//{
// MultiplyScalar op(val, scale);
// //if(src1.channels() == 1 && dst.channels() == 1)
// //{
// // callT(src1, dst, op, 0);
// //}
// //else
// //{
// callF(src1, dst, op, 0);
// //}
// //if(src1.channels() == 1 && dst.channels() == 1)
// //{
// // callT(src1, dst, op, 0);
// //}
// //else
// //{
// callF(src1, dst, op, 0);
// //}
//}
cl_mem bindTexture(const oclMat &mat, int depth, int channels)
static cl_mem bindTexture(const oclMat &mat, int depth, int channels)
{
cl_mem texture;
cl_image_format format;
@ -697,7 +698,7 @@ cl_mem bindTexture(const oclMat &mat, int depth, int channels)
{
format.image_channel_order = CL_RGBA;
}
#if CL_VERSION_1_2
#ifdef CL_VERSION_1_2
cl_image_desc desc;
desc.image_type = CL_MEM_OBJECT_IMAGE2D;
desc.image_width = mat.step / mat.elemSize();
@ -729,52 +730,75 @@ cl_mem bindTexture(const oclMat &mat, int depth, int channels)
return texture;
}
void releaseTexture(cl_mem texture)
static void releaseTexture(cl_mem texture)
{
openCLFree(texture);
}
void lkSparse_run(oclMat &I, oclMat &J,
const oclMat &prevPts, oclMat &nextPts, oclMat &status, oclMat *err, bool GET_MIN_EIGENVALS, int ptcount,
static void lkSparse_run(oclMat &I, oclMat &J,
const oclMat &prevPts, oclMat &nextPts, oclMat &status, oclMat& err, bool /*GET_MIN_EIGENVALS*/, int ptcount,
int level, /*dim3 block, */dim3 patch, Size winSize, int iters)
{
Context *clCxt = I.clCxt;
char platform[256] = {0};
cl_platform_id pid;
clGetDeviceInfo(*clCxt->impl->devices, CL_DEVICE_PLATFORM, sizeof(pid), &pid, NULL);
clGetPlatformInfo(pid, CL_PLATFORM_NAME, 256, platform, NULL);
std::string namestr = platform;
bool isImageSupported = true;
if(namestr.find("NVIDIA")!=string::npos || namestr.find("Intel")!=string::npos)
isImageSupported = false;
int elemCntPerRow = I.step / I.elemSize();
string kernelName = "lkSparse";
size_t localThreads[3] = { 8, 32, 1 };
size_t globalThreads[3] = { 8 * ptcount, 32, 1};
size_t localThreads[3] = { 8, isImageSupported?8:32, 1 };
size_t globalThreads[3] = { 8 * ptcount, isImageSupported?8:32, 1};
int cn = I.oclchannels();
bool calcErr;
if (err)
char calcErr;
if (level == 0)
{
calcErr = true;
calcErr = 1;
}
else
{
calcErr = false;
calcErr = 0;
}
calcErr = true;
cl_mem ITex = bindTexture(I, I.depth(), cn);
cl_mem JTex = bindTexture(J, J.depth(), cn);
vector<pair<size_t , const void *> > args;
cl_mem ITex;
cl_mem JTex;
if (isImageSupported)
{
ITex = bindTexture(I, I.depth(), cn);
JTex = bindTexture(J, J.depth(), cn);
}
else
{
ITex = (cl_mem)I.data;
JTex = (cl_mem)J.data;
}
args.push_back( make_pair( sizeof(cl_mem), (void *)&ITex ));
args.push_back( make_pair( sizeof(cl_mem), (void *)&JTex ));
//cl_mem clmD = clCreateBuffer(clCxt, CL_MEM_READ_WRITE, ptcount * sizeof(float), NULL, NULL);
args.push_back( make_pair( sizeof(cl_mem), (void *)&prevPts.data ));
args.push_back( make_pair( sizeof(cl_int), (void *)&prevPts.step ));
args.push_back( make_pair( sizeof(cl_mem), (void *)&nextPts.data ));
args.push_back( make_pair( sizeof(cl_int), (void *)&nextPts.step ));
args.push_back( make_pair( sizeof(cl_mem), (void *)&status.data ));
//args.push_back( make_pair( sizeof(cl_mem), (void *)&(err->data) ));
args.push_back( make_pair( sizeof(cl_mem), (void *)&err.data ));
args.push_back( make_pair( sizeof(cl_int), (void *)&level ));
args.push_back( make_pair( sizeof(cl_int), (void *)&I.rows ));
args.push_back( make_pair( sizeof(cl_int), (void *)&I.cols ));
if (!isImageSupported)
{
args.push_back( make_pair( sizeof(cl_int), (void *)&elemCntPerRow ) );
}
args.push_back( make_pair( sizeof(cl_int), (void *)&patch.x ));
args.push_back( make_pair( sizeof(cl_int), (void *)&patch.y ));
args.push_back( make_pair( sizeof(cl_int), (void *)&cn ));
@ -782,27 +806,29 @@ void lkSparse_run(oclMat &I, oclMat &J,
args.push_back( make_pair( sizeof(cl_int), (void *)&winSize.height ));
args.push_back( make_pair( sizeof(cl_int), (void *)&iters ));
args.push_back( make_pair( sizeof(cl_char), (void *)&calcErr ));
args.push_back( make_pair( sizeof(cl_char), (void *)&GET_MIN_EIGENVALS ));
//args.push_back( make_pair( sizeof(cl_char), (void *)&GET_MIN_EIGENVALS ));
openCLExecuteKernel2(clCxt, &pyrlk, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth(), CLFLUSH);
if (isImageSupported)
{
openCLExecuteKernel2(clCxt, &pyrlk, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth(), CLFLUSH);
releaseTexture(ITex);
releaseTexture(JTex);
releaseTexture(ITex);
releaseTexture(JTex);
}
else
{
//printf("Warning: The image2d_t is not supported by the device. Using alternative method!\n");
openCLExecuteKernel2(clCxt, &pyrlk_no_image, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth(), CLFLUSH);
}
}
void cv::ocl::PyrLKOpticalFlow::sparse(const oclMat &prevImg, const oclMat &nextImg, const oclMat &prevPts, oclMat &nextPts, oclMat &status, oclMat *err)
{
if (prevImg.clCxt->impl->devName.find("Intel(R) HD Graphics") != string::npos)
{
cout << " Intel HD GPU device unsupported " << endl;
return;
}
if (prevPts.empty())
{
nextPts.release();
status.release();
if (err) err->release();
//if (err) err->release();
return;
}
@ -836,8 +862,15 @@ void cv::ocl::PyrLKOpticalFlow::sparse(const oclMat &prevImg, const oclMat &next
//status.setTo(Scalar::all(1));
setTo(status, Scalar::all(1));
//if (err)
// ensureSizeIsEnough(1, prevPts.cols, CV_32FC1, *err);
bool errMat = false;
if (!err)
{
err = new oclMat(1, prevPts.cols, CV_32FC1);
errMat = true;
}
else
ensureSizeIsEnough(1, prevPts.cols, CV_32FC1, *err);
//ensureSizeIsEnough(1, prevPts.cols, CV_32FC1, err);
// build the image pyramids.
@ -872,17 +905,22 @@ void cv::ocl::PyrLKOpticalFlow::sparse(const oclMat &prevImg, const oclMat &next
for (int level = maxLevel; level >= 0; level--)
{
lkSparse_run(prevPyr_[level], nextPyr_[level],
prevPts, nextPts, status, level == 0 && err ? err : 0, getMinEigenVals, prevPts.cols,
prevPts, nextPts, status, *err, getMinEigenVals, prevPts.cols,
level, /*block, */patch, winSize, iters);
}
clFinish(prevImg.clCxt->impl->clCmdQueue);
if(errMat)
delete err;
}
void lkDense_run(oclMat &I, oclMat &J, oclMat &u, oclMat &v,
static void lkDense_run(oclMat &I, oclMat &J, oclMat &u, oclMat &v,
oclMat &prevU, oclMat &prevV, oclMat *err, Size winSize, int iters)
{
Context *clCxt = I.clCxt;
bool isImageSupported = clCxt->impl->devName.find("Intel(R) HD Graphics") == string::npos;
int elemCntPerRow = I.step / I.elemSize();
string kernelName = "lkDense";
@ -901,8 +939,19 @@ void lkDense_run(oclMat &I, oclMat &J, oclMat &u, oclMat &v,
calcErr = false;
}
cl_mem ITex = bindTexture(I, I.depth(), cn);
cl_mem JTex = bindTexture(J, J.depth(), cn);
cl_mem ITex;
cl_mem JTex;
if (isImageSupported)
{
ITex = bindTexture(I, I.depth(), cn);
JTex = bindTexture(J, J.depth(), cn);
}
else
{
ITex = (cl_mem)I.data;
JTex = (cl_mem)J.data;
}
//int2 halfWin = {(winSize.width - 1) / 2, (winSize.height - 1) / 2};
//const int patchWidth = 16 + 2 * halfWin.x;
@ -926,15 +975,27 @@ void lkDense_run(oclMat &I, oclMat &J, oclMat &u, oclMat &v,
args.push_back( make_pair( sizeof(cl_int), (void *)&I.cols ));
//args.push_back( make_pair( sizeof(cl_mem), (void *)&(*err).data ));
//args.push_back( make_pair( sizeof(cl_int), (void *)&(*err).step ));
if (!isImageSupported)
{
args.push_back( make_pair( sizeof(cl_int), (void *)&elemCntPerRow ) );
}
args.push_back( make_pair( sizeof(cl_int), (void *)&winSize.width ));
args.push_back( make_pair( sizeof(cl_int), (void *)&winSize.height ));
args.push_back( make_pair( sizeof(cl_int), (void *)&iters ));
args.push_back( make_pair( sizeof(cl_char), (void *)&calcErr ));
openCLExecuteKernel2(clCxt, &pyrlk, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth(), CLFLUSH);
if (isImageSupported)
{
openCLExecuteKernel2(clCxt, &pyrlk, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth(), CLFLUSH);
releaseTexture(ITex);
releaseTexture(JTex);
releaseTexture(ITex);
releaseTexture(JTex);
}
else
{
//printf("Warning: The image2d_t is not supported by the device. Using alternative method!\n");
openCLExecuteKernel2(clCxt, &pyrlk_no_image, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth(), CLFLUSH);
}
}
void cv::ocl::PyrLKOpticalFlow::dense(const oclMat &prevImg, const oclMat &nextImg, oclMat &u, oclMat &v, oclMat *err)

164
modules/ocl/src/split_merge.cpp
View File

@ -111,52 +111,52 @@ namespace cv
////////////////////////////////////////////////////////////////////////////
////////////////////merge//////////////////////////////////////////////////
////////////////////////////////////////////////////////////////////////////
void merge_vector_run_no_roi(const oclMat *mat_src, size_t n, oclMat &mat_dst)
{
Context *clCxt = mat_dst.clCxt;
int channels = mat_dst.oclchannels();
int depth = mat_dst.depth();
// static void merge_vector_run_no_roi(const oclMat *mat_src, size_t n, oclMat &mat_dst)
// {
// Context *clCxt = mat_dst.clCxt;
// int channels = mat_dst.oclchannels();
// int depth = mat_dst.depth();
string kernelName = "merge_vector";
// string kernelName = "merge_vector";
int indexes[4][7] = {{0, 0, 0, 0, 0, 0, 0},
{4, 4, 2, 2, 1, 1, 1},
{4, 4, 2, 2 , 1, 1, 1},
{4, 4, 2, 2, 1, 1, 1}
};
// int indexes[4][7] = {{0, 0, 0, 0, 0, 0, 0},
// {4, 4, 2, 2, 1, 1, 1},
// {4, 4, 2, 2 , 1, 1, 1},
// {4, 4, 2, 2, 1, 1, 1}
// };
size_t index = indexes[channels - 1][mat_dst.depth()];
int cols = divUp(mat_dst.cols, index);
size_t localThreads[3] = { 64, 4, 1 };
size_t globalThreads[3] = { divUp(cols, localThreads[0]) *localThreads[0],
divUp(mat_dst.rows, localThreads[1]) *localThreads[1],
1
};
// size_t index = indexes[channels - 1][mat_dst.depth()];
// int cols = divUp(mat_dst.cols, index);
// size_t localThreads[3] = { 64, 4, 1 };
// size_t globalThreads[3] = { divUp(cols, localThreads[0]) *localThreads[0],
// divUp(mat_dst.rows, localThreads[1]) *localThreads[1],
// 1
// };
vector<pair<size_t , const void *> > args;
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst.rows));
args.push_back( make_pair( sizeof(cl_int), (void *)&cols));
args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_dst.data));
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst.step));
args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_src[0].data));
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src[0].step));
args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_src[1].data));
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src[1].step));
if(n >= 3)
{
args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_src[2].data));
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src[2].step));
}
if(n >= 4)
{
args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_src[3].data));
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src[3].step));
}
// vector<pair<size_t , const void *> > args;
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst.rows));
// args.push_back( make_pair( sizeof(cl_int), (void *)&cols));
// args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_dst.data));
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst.step));
// args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_src[0].data));
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src[0].step));
// args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_src[1].data));
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src[1].step));
// if(n >= 3)
// {
// args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_src[2].data));
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src[2].step));
// }
// if(n >= 4)
// {
// args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_src[3].data));
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src[3].step));
// }
openCLExecuteKernel(clCxt, &merge_mat, kernelName, globalThreads, localThreads, args, channels, depth);
}
// openCLExecuteKernel(clCxt, &merge_mat, kernelName, globalThreads, localThreads, args, channels, depth);
// }
void merge_vector_run(const oclMat *mat_src, size_t n, oclMat &mat_dst)
static void merge_vector_run(const oclMat *mat_src, size_t n, oclMat &mat_dst)
{
if(mat_dst.clCxt -> impl -> double_support == 0 && mat_dst.type() == CV_64F)
{
@ -228,7 +228,7 @@ namespace cv
openCLExecuteKernel(clCxt, &merge_mat, kernelName, globalThreads, localThreads, args, channels, depth);
}
void merge(const oclMat *mat_src, size_t n, oclMat &mat_dst)
static void merge(const oclMat *mat_src, size_t n, oclMat &mat_dst)
{
CV_Assert(mat_src);
CV_Assert(n > 0);
@ -260,51 +260,51 @@ namespace cv
////////////////////////////////////////////////////////////////////////////////////////////////////
//////////////////////////////////////split/////////////////////////////////////////////////////////////
//////////////////////////////////////////////////////////////////////////////////////////////////
void split_vector_run_no_roi(const oclMat &mat_src, oclMat *mat_dst)
{
Context *clCxt = mat_src.clCxt;
int channels = mat_src.oclchannels();
int depth = mat_src.depth();
// static void split_vector_run_no_roi(const oclMat &mat_src, oclMat *mat_dst)
// {
// Context *clCxt = mat_src.clCxt;
// int channels = mat_src.oclchannels();
// int depth = mat_src.depth();
string kernelName = "split_vector";
// string kernelName = "split_vector";
int indexes[4][7] = {{0, 0, 0, 0, 0, 0, 0},
{8, 8, 8, 8, 4, 4, 2},
{8, 8, 8, 8 , 4, 4, 4},
{4, 4, 2, 2, 1, 1, 1}
};
// int indexes[4][7] = {{0, 0, 0, 0, 0, 0, 0},
// {8, 8, 8, 8, 4, 4, 2},
// {8, 8, 8, 8 , 4, 4, 4},
// {4, 4, 2, 2, 1, 1, 1}
// };
size_t index = indexes[channels - 1][mat_dst[0].depth()];
int cols = divUp(mat_src.cols, index);
size_t localThreads[3] = { 64, 4, 1 };
size_t globalThreads[3] = { divUp(cols, localThreads[0]) *localThreads[0],
divUp(mat_src.rows, localThreads[1]) *localThreads[1],
1
};
// size_t index = indexes[channels - 1][mat_dst[0].depth()];
// int cols = divUp(mat_src.cols, index);
// size_t localThreads[3] = { 64, 4, 1 };
// size_t globalThreads[3] = { divUp(cols, localThreads[0]) *localThreads[0],
// divUp(mat_src.rows, localThreads[1]) *localThreads[1],
// 1
// };
vector<pair<size_t , const void *> > args;
args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_src.data));
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src.step));
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src.rows));
args.push_back( make_pair( sizeof(cl_int), (void *)&cols));
args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_dst[0].data));
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst[0].step));
args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_dst[1].data));
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst[1].step));
if(channels >= 3)
{
args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_dst[2].data));
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst[2].step));
}
if(channels >= 4)
{
args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_dst[3].data));
args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst[3].step));
}
// vector<pair<size_t , const void *> > args;
// args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_src.data));
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src.step));
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_src.rows));
// args.push_back( make_pair( sizeof(cl_int), (void *)&cols));
// args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_dst[0].data));
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst[0].step));
// args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_dst[1].data));
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst[1].step));
// if(channels >= 3)
// {
// args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_dst[2].data));
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst[2].step));
// }
// if(channels >= 4)
// {
// args.push_back( make_pair( sizeof(cl_mem), (void *)&mat_dst[3].data));
// args.push_back( make_pair( sizeof(cl_int), (void *)&mat_dst[3].step));
// }
openCLExecuteKernel(clCxt, &split_mat, kernelName, globalThreads, localThreads, args, channels, depth);
}
void split_vector_run(const oclMat &mat_src, oclMat *mat_dst)
// openCLExecuteKernel(clCxt, &split_mat, kernelName, globalThreads, localThreads, args, channels, depth);
// }
static void split_vector_run(const oclMat &mat_src, oclMat *mat_dst)
{
if(mat_src.clCxt -> impl -> double_support == 0 && mat_src.type() == CV_64F)
@ -374,7 +374,7 @@ namespace cv
openCLExecuteKernel(clCxt, &split_mat, kernelName, globalThreads, localThreads, args, channels, depth);
}
void split(const oclMat &mat_src, oclMat *mat_dst)
static void split(const oclMat &mat_src, oclMat *mat_dst)
{
CV_Assert(mat_dst);

2
modules/ocl/src/surf.cpp
View File

@ -536,7 +536,7 @@ void SURF_OCL_Invoker::bindImgTex(const oclMat &img, cl_mem &texture)
openCLFree(texture);
}
#if CL_VERSION_1_2
#ifdef CL_VERSION_1_2
cl_image_desc desc;
desc.image_type = CL_MEM_OBJECT_IMAGE2D;
desc.image_width = img.step / img.elemSize();

9
modules/ocl/test/precomp.hpp
View File

@ -38,6 +38,15 @@
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#ifdef __GNUC__
# pragma GCC diagnostic ignored "-Wmissing-declarations"
# if defined __clang__ || defined __APPLE__
# pragma GCC diagnostic ignored "-Wmissing-prototypes"
# pragma GCC diagnostic ignored "-Wextra"
# endif
#endif
#ifndef __OPENCV_TEST_PRECOMP_HPP__
#define __OPENCV_TEST_PRECOMP_HPP__

10
modules/ocl/test/test_arithm.cpp
View File

@ -1110,8 +1110,8 @@ TEST_P(Phase, Mat)
for(int j = 0; j < LOOP_TIMES; j++)
{
random_roi();
cv::phase(mat1_roi, mat2_roi, dst_roi, angelInDegrees);
cv::ocl::phase(gmat1, gmat2, gdst, angelInDegrees);
cv::phase(mat1_roi, mat2_roi, dst_roi, angelInDegrees ? true : false);
cv::ocl::phase(gmat1, gmat2, gdst, angelInDegrees ? true : false);
cv::Mat cpu_dst;
gdst_whole.download(cpu_dst);
@ -1449,8 +1449,8 @@ TEST_P(MagnitudeSqr, Mat)
for(int j = 0; j < LOOP_TIMES; j++)
{
// random_roi();
int64 start, end;
start = cv::getTickCount();
// int64 start, end;
// start = cv::getTickCount();
for(int i = 0; i < mat1.rows; ++i)
for(int j = 0; j < mat1.cols; ++j)
{
@ -1465,7 +1465,7 @@ TEST_P(MagnitudeSqr, Mat)
// ((float *)(dst.data))[i*dst.step/4 +j]= val1 * val1 +val2 * val2;
}
end = cv::getTickCount();
// end = cv::getTickCount();

2
modules/ocl/test/test_blend.cpp
View File

@ -74,7 +74,7 @@ TEST_P(Blend, Accuracy)
else
blendLinearGold<float>(img1, img2, weights1, weights2, result_gold);
EXPECT_MAT_NEAR(result_gold, result, CV_MAT_DEPTH(type) == CV_8U ? 1 : 1e-5f, NULL)
EXPECT_MAT_NEAR(result_gold, result, CV_MAT_DEPTH(type) == CV_8U ? 1.f : 1e-5f, 0);
}
INSTANTIATE_TEST_CASE_P(GPU_ImgProc, Blend, Combine(

193
modules/ocl/test/test_color.cpp Normal file
View File

@ -0,0 +1,193 @@
/*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-2012, Multicoreware, Inc., all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
// Peng Xiao, pengxiao@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 oclMaterials 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.
//
//M*/
#include "precomp.hpp"
#ifdef HAVE_OPENCL
//#define MAT_DEBUG
#ifdef MAT_DEBUG
#define MAT_DIFF(mat, mat2)\
{\
for(int i = 0; i < mat.rows; i ++)\
{\
for(int j = 0; j < mat.cols; j ++)\
{\
cv::Vec4b s = mat.at<cv::Vec4b>(i, j);\
cv::Vec4b s2 = mat2.at<cv::Vec4b>(i, j);\
if(s != s2) printf("*");\
else printf(".");\
}\
puts("\n");\
}\
}
#else
#define MAT_DIFF(mat, mat2)
#endif
namespace
{
///////////////////////////////////////////////////////////////////////////////////////////////////////
// cvtColor
PARAM_TEST_CASE(CvtColor, cv::Size, MatDepth)
{
cv::Size size;
int depth;
bool useRoi;
cv::Mat img;
virtual void SetUp()
{
size = GET_PARAM(0);
depth = GET_PARAM(1);
img = randomMat(size, CV_MAKE_TYPE(depth, 3), 0.0, depth == CV_32F ? 1.0 : 255.0);
}
};
#define CVTCODE(name) cv::COLOR_ ## name
#define TEST_P_CVTCOLOR(name) TEST_P(CvtColor, name)\
{\
cv::Mat src = img;\
cv::ocl::oclMat ocl_img, dst;\
ocl_img.upload(img);\
cv::ocl::cvtColor(ocl_img, dst, CVTCODE(name));\
cv::Mat dst_gold;\
cv::cvtColor(src, dst_gold, CVTCODE(name));\
cv::Mat dst_mat;\
dst.download(dst_mat);\
EXPECT_MAT_NEAR(dst_gold, dst_mat, 1e-5, "");\
}
//add new ones here using macro
TEST_P_CVTCOLOR(RGB2GRAY)
TEST_P_CVTCOLOR(BGR2GRAY)
TEST_P_CVTCOLOR(RGBA2GRAY)
TEST_P_CVTCOLOR(BGRA2GRAY)
TEST_P_CVTCOLOR(RGB2YUV)
TEST_P_CVTCOLOR(BGR2YUV)
TEST_P_CVTCOLOR(YUV2RGB)
TEST_P_CVTCOLOR(YUV2BGR)
TEST_P_CVTCOLOR(RGB2YCrCb)
TEST_P_CVTCOLOR(BGR2YCrCb)
PARAM_TEST_CASE(CvtColor_Gray2RGB, cv::Size, MatDepth, int)
{
cv::Size size;
int code;
int depth;
cv::Mat img;
virtual void SetUp()
{
size = GET_PARAM(0);
depth = GET_PARAM(1);
code = GET_PARAM(2);
img = randomMat(size, CV_MAKETYPE(depth, 1), 0.0, depth == CV_32F ? 1.0 : 255.0);
}
};
TEST_P(CvtColor_Gray2RGB, Accuracy)
{
cv::Mat src = img;
cv::ocl::oclMat ocl_img, dst;
ocl_img.upload(src);
cv::ocl::cvtColor(ocl_img, dst, code);
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, code);
cv::Mat dst_mat;
dst.download(dst_mat);
EXPECT_MAT_NEAR(dst_gold, dst_mat, 1e-5, "");
}
PARAM_TEST_CASE(CvtColor_YUV420, cv::Size, int)
{
cv::Size size;
int code;
cv::Mat img;
virtual void SetUp()
{
size = GET_PARAM(0);
code = GET_PARAM(1);
img = randomMat(size, CV_8UC1, 0.0, 255.0);
}
};
TEST_P(CvtColor_YUV420, Accuracy)
{
cv::Mat src = img;
cv::ocl::oclMat ocl_img, dst;
ocl_img.upload(src);
cv::ocl::cvtColor(ocl_img, dst, code);
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, code);
cv::Mat dst_mat;
dst.download(dst_mat);
MAT_DIFF(dst_mat, dst_gold);
EXPECT_MAT_NEAR(dst_gold, dst_mat, 1e-5, "");
}
INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor, testing::Combine(
DIFFERENT_SIZES,
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F))
));
INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor_YUV420, testing::Combine(
testing::Values(cv::Size(128, 45), cv::Size(46, 132), cv::Size(1024, 1023)),
testing::Values((int)CV_YUV2RGBA_NV12, (int)CV_YUV2BGRA_NV12, (int)CV_YUV2RGB_NV12, (int)CV_YUV2BGR_NV12)
));
INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor_Gray2RGB, testing::Combine(
DIFFERENT_SIZES,
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)),
testing::Values((int)CV_GRAY2BGR, (int)CV_GRAY2BGRA, (int)CV_GRAY2RGB, (int)CV_GRAY2RGBA)
));
}
#endif

18
modules/ocl/test/test_haar.cpp
View File

@ -109,15 +109,15 @@ TEST_F(Haar, FaceDetect)
//double t = 0;
vector<Rect> faces, oclfaces;
const static Scalar colors[] = { CV_RGB(0, 0, 255),
CV_RGB(0, 128, 255),
CV_RGB(0, 255, 255),
CV_RGB(0, 255, 0),
CV_RGB(255, 128, 0),
CV_RGB(255, 255, 0),
CV_RGB(255, 0, 0),
CV_RGB(255, 0, 255)
} ;
// const static Scalar colors[] = { CV_RGB(0, 0, 255),
// CV_RGB(0, 128, 255),
// CV_RGB(0, 255, 255),
// CV_RGB(0, 255, 0),
// CV_RGB(255, 128, 0),
// CV_RGB(255, 255, 0),
// CV_RGB(255, 0, 0),
// CV_RGB(255, 0, 255)
// } ;
Mat gray, smallImg(cvRound (img.rows / scale), cvRound(img.cols / scale), CV_8UC1 );
MemStorage storage(cvCreateMemStorage(0));

7
modules/ocl/test/test_imgproc.cpp
View File

@ -498,11 +498,11 @@ TEST_P(bilateralFilter, Mat)
}
else
{
for(int i = 0; i < sizeof(bordertype) / sizeof(int); i++)
for(size_t i = 0; i < sizeof(bordertype) / sizeof(int); i++)
for(int j = 0; j < LOOP_TIMES; j++)
{
random_roi();
if(((bordertype[i] != cv::BORDER_CONSTANT) && (bordertype[i] != cv::BORDER_REPLICATE)) && (mat1_roi.cols <= radius) || (mat1_roi.cols <= radius) || (mat1_roi.rows <= radius) || (mat1_roi.rows <= radius))
if(((bordertype[i] != cv::BORDER_CONSTANT) && (bordertype[i] != cv::BORDER_REPLICATE) && (mat1_roi.cols <= radius)) || (mat1_roi.cols <= radius) || (mat1_roi.rows <= radius) || (mat1_roi.rows <= radius))
{
continue;
}
@ -563,7 +563,7 @@ TEST_P(CopyMakeBorder, Mat)
}
else
{
for(int i = 0; i < sizeof(bordertype) / sizeof(int); i++)
for(size_t i = 0; i < sizeof(bordertype) / sizeof(int); i++)
for(int j = 0; j < LOOP_TIMES; j++)
{
random_roi();
@ -911,7 +911,6 @@ PARAM_TEST_CASE(Remap, MatType, MatType, MatType, int, int)
cv::RNG &rng = TS::ptr()->get_rng();
cv::Size srcSize = cv::Size(MWIDTH, MHEIGHT);
cv::Size dstSize = cv::Size(MWIDTH, MHEIGHT);
cv::Size map1Size = cv::Size(MWIDTH, MHEIGHT);
double min = 5, max = 16;

4
modules/ocl/test/test_match_template.cpp
View File

@ -100,7 +100,7 @@ TEST_P(MatchTemplate8U, Accuracy)
EXPECT_MAT_NEAR(dst_gold, mat_dst, templ_size.area() * 1e-1, sss);
#if PERF_TEST
#ifdef PERF_TEST
{
P_TEST_FULL( {}, {cv::ocl::matchTemplate(ocl_image, ocl_templ, dst, method);}, {});
P_TEST_FULL( {}, {cv::matchTemplate(image, templ, dst_gold, method);}, {});
@ -145,7 +145,7 @@ TEST_P(MatchTemplate32F, Accuracy)
EXPECT_MAT_NEAR(dst_gold, mat_dst, templ_size.area() * 1e-1, sss);
#if PERF_TEST
#ifdef PERF_TEST
{
std::cout << "Method: " << TEMPLATE_METHOD_NAMES[method] << std::endl;
std::cout << "Image Size: (" << size.width << ", " << size.height << ")" << std::endl;

10
modules/ocl/test/test_pyrlk.cpp
View File

@ -118,9 +118,9 @@ TEST_P(Sparse, Mat)
cv::Mat status_mat(1, d_status.cols, CV_8UC1, (void *)&status[0]);
d_status.download(status_mat);
//std::vector<float> err(d_err.cols);
//cv::Mat err_mat(1, d_err.cols, CV_32FC1, (void*)&err[0]);
//d_err.download(err_mat);
std::vector<float> err(d_err.cols);
cv::Mat err_mat(1, d_err.cols, CV_32FC1, (void*)&err[0]);
d_err.download(err_mat);
std::vector<cv::Point2f> nextPts_gold;
std::vector<unsigned char> status_gold;
@ -153,9 +153,9 @@ TEST_P(Sparse, Mat)
}
}
double bad_ratio = static_cast<double>(mistmatch) / (nextPts.size() * 2);
double bad_ratio = static_cast<double>(mistmatch) / (nextPts.size());
ASSERT_LE(bad_ratio, 0.05f);
ASSERT_LE(bad_ratio, 0.02f);
}

12
modules/ocl/test/utility.hpp
View File

@ -76,12 +76,12 @@ double checkSimilarity(const cv::Mat &m1, const cv::Mat &m2);
EXPECT_LE(checkNorm(cv::Mat(mat)), eps) \
}
//#define EXPECT_MAT_NEAR(mat1, mat2, eps) \
//{ \
// ASSERT_EQ(mat1.type(), mat2.type()); \
// ASSERT_EQ(mat1.size(), mat2.size()); \
// EXPECT_LE(checkNorm(cv::Mat(mat1), cv::Mat(mat2)), eps); \
//}
/*#define EXPECT_MAT_NEAR(mat1, mat2, eps) \
{ \
ASSERT_EQ(mat1.type(), mat2.type()); \
ASSERT_EQ(mat1.size(), mat2.size()); \
EXPECT_LE(checkNorm(cv::Mat(mat1), cv::Mat(mat2)), eps); \
}*/
#define EXPECT_MAT_NEAR(mat1, mat2, eps,s) \
{ \

4
modules/python/CMakeLists.txt
View File

@ -80,6 +80,10 @@ set_target_properties(${the_module} PROPERTIES
OUTPUT_NAME cv2
SUFFIX ${CVPY_SUFFIX})
if(ENABLE_SOLUTION_FOLDERS)
set_target_properties(${the_module} PROPERTIES FOLDER "bindings")
endif()
if(CMAKE_COMPILER_IS_GNUCXX AND NOT ENABLE_NOISY_WARNINGS)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-unused-function")
endif()

6
modules/python/src2/gen2.py
View File

@ -243,9 +243,9 @@ class ClassInfo(object):
if decl:
self.bases = decl[1].split()[1:]
if len(self.bases) > 1:
print "Warning: class %s has more than 1 base class (not supported by Python C extensions)" % (self.name,)
print "Bases: ", " ".join(self.bases)
print "Only the first base class will be used"
print "Note: Class %s has more than 1 base class (not supported by Python C extensions)" % (self.name,)
print " Bases: ", " ".join(self.bases)
print " Only the first base class will be used"
self.bases = [self.bases[0].strip(",")]
#return sys.exit(-1)
if self.bases and self.bases[0].startswith("cv::"):

9
modules/ts/include/opencv2/ts/ts.hpp
View File

@ -547,6 +547,15 @@ struct CV_EXPORTS DefaultRngAuto
}
namespace cvtest
{
// test images generation functions
CV_EXPORTS void fillGradient(Mat& img, int delta = 5);
CV_EXPORTS void smoothBorder(Mat& img, const Scalar& color, int delta = 3);
} //namespace cvtest
// fills c with zeros
CV_EXPORTS void cvTsZero( CvMat* c, const CvMat* mask=0 );

31
modules/ts/misc/perf_tests_timing.py
View File

@ -56,33 +56,39 @@ if __name__ == "__main__":
test_cases[name] = [None] * setsCount
test_cases[name][i] = case
testsuits = [] # testsuit name, time, flag for failed tests
testsuits = [] # testsuit name, time, num, flag for failed tests
overall_time = 0
prevGroupName = None
suit_time = 0
has_failed = False
suit_num = 0
fails_num = 0
for name in sorted(test_cases.iterkeys(), key=alphanum_keyselector):
cases = test_cases[name]
groupName = next(c for c in cases if c).shortName()
if groupName != prevGroupName:
if prevGroupName != None:
suit_time = suit_time/60 #from seconds to minutes
testsuits.append({'name': prevGroupName, 'time': suit_time, \
'failed': has_failed})
has_failed = False
'num': suit_num, 'failed': fails_num})
overall_time += suit_time
suit_time = 0
suit_num = 0
fails_num = 0
prevGroupName = groupName
for i in range(setsCount):
case = cases[i]
if not case is None:
suit_num += 1
if case.get('status') == 'run':
suit_time += case.get('time')
if case.get('status') == 'failed':
has_failed = True
fails_num += 1
testsuits.append({'name': prevGroupName, 'time': suit_time, \
'failed': has_failed})
'num': suit_num, 'failed': fails_num})
if len(testsuits)==0:
print 'No testsuits found'
@ -91,17 +97,19 @@ if __name__ == "__main__":
tbl = table()
# header
tbl.newColumn('name', 'Name of testsuit', align = 'left', cssclass = 'col_name')
tbl.newColumn('time', 'Time (ms)', align = 'left', cssclass = 'col_name')
tbl.newColumn('failed', 'Failed tests', align = 'center', cssclass = 'col_name')
tbl.newColumn('name', 'Testsuit', align = 'left', cssclass = 'col_name')
tbl.newColumn('time', 'Time (min)', align = 'center', cssclass = 'col_name')
tbl.newColumn('num', 'Num of tests', align = 'center', cssclass = 'col_name')
tbl.newColumn('failed', 'Failed', align = 'center', cssclass = 'col_name')
# rows
for suit in sorted(testsuits, key = lambda suit: suit['time'], reverse = True):
tbl.newRow()
tbl.newCell('name', suit['name'])
tbl.newCell('time', formatValue(suit['time'], '', ''), suit['time'])
if (suit['failed']):
tbl.newCell('failed', 'Yes')
tbl.newCell('num', suit['num'])
if (suit['failed'] != 0):
tbl.newCell('failed', suit['failed'])
else:
tbl.newCell('failed', ' ')
@ -116,5 +124,6 @@ if __name__ == "__main__":
splitter = 15 * '*'
print '\n%s\n %s\n%s\n' % (splitter, module_name, splitter)
print 'Overall time: %.2f min\n' % overall_time
tbl.consolePrintTable(sys.stdout)
print 4 * '\n'

14
modules/ts/misc/run.py
View File

@ -69,6 +69,8 @@ parse_patterns = (
{'name': "ndk_path", 'default': None, 'pattern': re.compile("^(?:ANDROID_NDK|ANDROID_STANDALONE_TOOLCHAIN)?:PATH=(.*)$")},
{'name': "android_abi", 'default': None, 'pattern': re.compile("^ANDROID_ABI:STRING=(.*)$")},
{'name': "android_executable", 'default': None, 'pattern': re.compile("^ANDROID_EXECUTABLE:FILEPATH=(.*android.*)$")},
{'name': "ant_executable", 'default': None, 'pattern': re.compile("^ANT_EXECUTABLE:FILEPATH=(.*ant.*)$")},
{'name': "java_test_binary_dir", 'default': None, 'pattern': re.compile("^opencv_test_java_BINARY_DIR:STATIC=(.*)$")},
{'name': "is_x64", 'default': "OFF", 'pattern': re.compile("^CUDA_64_BIT_DEVICE_CODE:BOOL=(ON)$")},#ugly(
{'name': "cmake_generator", 'default': None, 'pattern': re.compile("^CMAKE_GENERATOR:INTERNAL=(.+)$")},
{'name': "cxx_compiler", 'default': None, 'pattern': re.compile("^CMAKE_CXX_COMPILER:FILEPATH=(.+)$")},
@ -431,6 +433,8 @@ class TestSuite(object):
if self.tests_dir and os.path.isdir(self.tests_dir):
files = glob.glob(os.path.join(self.tests_dir, self.nameprefix + "*"))
files = [f for f in files if self.isTest(f)]
if self.ant_executable and self.java_test_binary_dir:
files.append("java")
return files
return []
@ -740,6 +744,16 @@ class TestSuite(object):
if os.path.isfile(hostlogpath):
return hostlogpath
return None
elif path == "java":
cmd = [self.ant_executable, "-DjavaLibraryPath=" + self.tests_dir, "buildAndTest"]
print >> _stderr, "Run command:", " ".join(cmd)
try:
Popen(cmd, stdout=_stdout, stderr=_stderr, cwd = self.java_test_binary_dir + "/.build").wait()
except OSError:
pass
return None
else:
cmd = [exe]
if self.options.help:

Some files were not shown because too many files have changed in this diff Show More